|
|
|
@ -173,6 +173,7 @@ __kernel void search1(__global hash_t* hashes)
@@ -173,6 +173,7 @@ __kernel void search1(__global hash_t* hashes)
|
|
|
|
|
T6_L[i] = T6[i]; |
|
|
|
|
T7_L[i] = T7[i]; |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE); |
|
|
|
|
|
|
|
|
|
#define T0 T0_L |
|
|
|
@ -185,8 +186,10 @@ __kernel void search1(__global hash_t* hashes)
@@ -185,8 +186,10 @@ __kernel void search1(__global hash_t* hashes)
|
|
|
|
|
#define T7 T7_L |
|
|
|
|
|
|
|
|
|
sph_u64 H[16]; |
|
|
|
|
|
|
|
|
|
for (unsigned int u = 0; u < 15; u ++) |
|
|
|
|
H[u] = 0; |
|
|
|
|
|
|
|
|
|
#if USE_LE |
|
|
|
|
H[15] = ((sph_u64)(512 & 0xFF) << 56) | ((sph_u64)(512 & 0xFF00) << 40); |
|
|
|
|
#else |
|
|
|
@ -202,8 +205,10 @@ __kernel void search1(__global hash_t* hashes)
@@ -202,8 +205,10 @@ __kernel void search1(__global hash_t* hashes)
|
|
|
|
|
m[5] = DEC64E(hash->h8[5]); |
|
|
|
|
m[6] = DEC64E(hash->h8[6]); |
|
|
|
|
m[7] = DEC64E(hash->h8[7]); |
|
|
|
|
|
|
|
|
|
for (unsigned int u = 0; u < 16; u ++) |
|
|
|
|
g[u] = m[u] ^ H[u]; |
|
|
|
|
|
|
|
|
|
m[8] = 0x80; g[8] = m[8] ^ H[8]; |
|
|
|
|
m[9] = 0; g[9] = m[9] ^ H[9]; |
|
|
|
|
m[10] = 0; g[10] = m[10] ^ H[10]; |
|
|
|
@ -212,16 +217,23 @@ __kernel void search1(__global hash_t* hashes)
@@ -212,16 +217,23 @@ __kernel void search1(__global hash_t* hashes)
|
|
|
|
|
m[13] = 0; g[13] = m[13] ^ H[13]; |
|
|
|
|
m[14] = 0; g[14] = m[14] ^ H[14]; |
|
|
|
|
m[15] = 0x100000000000000; g[15] = m[15] ^ H[15]; |
|
|
|
|
|
|
|
|
|
PERM_BIG_P(g); |
|
|
|
|
PERM_BIG_Q(m); |
|
|
|
|
|
|
|
|
|
for (unsigned int u = 0; u < 16; u ++) |
|
|
|
|
H[u] ^= g[u] ^ m[u]; |
|
|
|
|
|
|
|
|
|
sph_u64 xH[16]; |
|
|
|
|
|
|
|
|
|
for (unsigned int u = 0; u < 16; u ++) |
|
|
|
|
xH[u] = H[u]; |
|
|
|
|
|
|
|
|
|
PERM_BIG_P(xH); |
|
|
|
|
|
|
|
|
|
for (unsigned int u = 0; u < 16; u ++) |
|
|
|
|
H[u] ^= xH[u]; |
|
|
|
|
|
|
|
|
|
for (unsigned int u = 0; u < 8; u ++) |
|
|
|
|
hash->h8[u] = DEC64E(H[u + 8]); |
|
|
|
|
|
|
|
|
@ -239,10 +251,10 @@ __kernel void search2(__global hash_t* hashes)
@@ -239,10 +251,10 @@ __kernel void search2(__global hash_t* hashes)
|
|
|
|
|
sph_u64 h4h = C64e(0x0169e60541e34a69), h4l = C64e(0x46b58a8e2e6fe65a), h5h = C64e(0x1047a7d0c1843c24), h5l = C64e(0x3b6e71b12d5ac199), h6h = C64e(0xcf57f6ec9db1f856), h6l = C64e(0xa706887c5716b156), h7h = C64e(0xe3c2fcdfe68517fb), h7l = C64e(0x545a4678cc8cdd4b); |
|
|
|
|
sph_u64 tmp; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
for(int i = 0; i < 2; i++) |
|
|
|
|
{ |
|
|
|
|
if (i == 0) { |
|
|
|
|
if (i == 0) |
|
|
|
|
{ |
|
|
|
|
h0h ^= DEC64E(hash->h8[0]); |
|
|
|
|
h0l ^= DEC64E(hash->h8[1]); |
|
|
|
|
h1h ^= DEC64E(hash->h8[2]); |
|
|
|
@ -251,7 +263,9 @@ __kernel void search2(__global hash_t* hashes)
@@ -251,7 +263,9 @@ __kernel void search2(__global hash_t* hashes)
|
|
|
|
|
h2l ^= DEC64E(hash->h8[5]); |
|
|
|
|
h3h ^= DEC64E(hash->h8[6]); |
|
|
|
|
h3l ^= DEC64E(hash->h8[7]); |
|
|
|
|
} else if(i == 1) { |
|
|
|
|
} |
|
|
|
|
else if(i == 1) |
|
|
|
|
{ |
|
|
|
|
h4h ^= DEC64E(hash->h8[0]); |
|
|
|
|
h4l ^= DEC64E(hash->h8[1]); |
|
|
|
|
h5h ^= DEC64E(hash->h8[2]); |
|
|
|
@ -266,6 +280,7 @@ __kernel void search2(__global hash_t* hashes)
@@ -266,6 +280,7 @@ __kernel void search2(__global hash_t* hashes)
|
|
|
|
|
} |
|
|
|
|
E8; |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
h4h ^= 0x80; |
|
|
|
|
h7l ^= 0x2000000000000; |
|
|
|
|
|
|
|
|
@ -314,6 +329,7 @@ __kernel void search3(__global hash_t* hashes)
@@ -314,6 +329,7 @@ __kernel void search3(__global hash_t* hashes)
|
|
|
|
|
a21 ^= SWAP8(hash->h8[7]); |
|
|
|
|
a31 ^= 0x8000000000000001; |
|
|
|
|
KECCAK_F_1600; |
|
|
|
|
|
|
|
|
|
// Finalize the "lane complement" |
|
|
|
|
a10 = ~a10; |
|
|
|
|
a20 = ~a20; |
|
|
|
@ -340,9 +356,8 @@ __kernel void search4(__global hash_t* hashes, __global uint* output, const ulon
@@ -340,9 +356,8 @@ __kernel void search4(__global hash_t* hashes, __global uint* output, const ulon
|
|
|
|
|
hash_t hash; |
|
|
|
|
__global hash_t *hashp = &(hashes[gid-offset]); |
|
|
|
|
|
|
|
|
|
for (int i = 0; i < 8; i++) { |
|
|
|
|
for (int i = 0; i < 8; i++) |
|
|
|
|
hash.h8[i] = hashes[gid-offset].h8[i]; |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
sph_u64 h0 = SPH_C64(0x4903ADFF749C51CE), h1 = SPH_C64(0x0D95DE399746DF03), h2 = SPH_C64(0x8FD1934127C79BCE), h3 = SPH_C64(0x9A255629FF352CB1), h4 = SPH_C64(0x5DB62599DF6CA7B0), h5 = SPH_C64(0xEABE394CA9D5C3F4), h6 = SPH_C64(0x991112C71A75B523), h7 = SPH_C64(0xAE18A40B660FCC33); |
|
|
|
|
sph_u64 m0, m1, m2, m3, m4, m5, m6, m7; |
|
|
|
@ -356,10 +371,14 @@ __kernel void search4(__global hash_t* hashes, __global uint* output, const ulon
@@ -356,10 +371,14 @@ __kernel void search4(__global hash_t* hashes, __global uint* output, const ulon
|
|
|
|
|
m5 = SWAP8(hash.h8[5]); |
|
|
|
|
m6 = SWAP8(hash.h8[6]); |
|
|
|
|
m7 = SWAP8(hash.h8[7]); |
|
|
|
|
|
|
|
|
|
UBI_BIG(480, 64); |
|
|
|
|
|
|
|
|
|
bcount = 0; |
|
|
|
|
m0 = m1 = m2 = m3 = m4 = m5 = m6 = m7 = 0; |
|
|
|
|
|
|
|
|
|
UBI_BIG(510, 8); |
|
|
|
|
|
|
|
|
|
hash.h8[0] = SWAP8(h0); |
|
|
|
|
hash.h8[1] = SWAP8(h1); |
|
|
|
|
hash.h8[2] = SWAP8(h2); |
|
|
|
@ -372,7 +391,6 @@ __kernel void search4(__global hash_t* hashes, __global uint* output, const ulon
@@ -372,7 +391,6 @@ __kernel void search4(__global hash_t* hashes, __global uint* output, const ulon
|
|
|
|
|
bool result = (SWAP8(hash.h8[3]) <= target); |
|
|
|
|
if (result) |
|
|
|
|
output[atomic_inc(output+0xFF)] = SWAP4(gid); |
|
|
|
|
|
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
#endif // TALKCOIN_MOD_CL |
|
|
|
|