diff --git a/libethash-cuda/dagger_shared.cuh b/libethash-cuda/dagger_shared.cuh index 185b1f5e5..cd67b31d1 100644 --- a/libethash-cuda/dagger_shared.cuh +++ b/libethash-cuda/dagger_shared.cuh @@ -21,16 +21,8 @@ __device__ hash64_t init_hash(uint64_t nonce) state[2] = d_header.uint64s[2]; state[3] = d_header.uint64s[3]; state[4] = nonce; - state[5] = 0x0000000000000001; - state[6] = 0; - state[7] = 0; - state[8] = 0x8000000000000000; - for (uint32_t i = 9; i < 25; i++) - { - state[i] = 0; - } - keccak_f1600_block((uint2 *)state, 8); + keccak_f1600_init((uint2 *)state); copy(init.uint64s, state, 8); return init; } @@ -83,18 +75,8 @@ __device__ hash32_t final_hash(hash64_t const* init, hash32_t const* mix) // keccak_256(keccak_512(header..nonce) .. mix); copy(state, init->uint64s, 8); copy(state + 8, mix->uint64s, 4); - state[12] = 0x0000000000000001; - for (uint32_t i = 13; i < 16; i++) - { - state[i] = 0; - } - state[16] = 0x8000000000000000; - for (uint32_t i = 17; i < 25; i++) - { - state[i] = 0; - } - keccak_f1600_block((uint2 *)state, 1); + keccak_f1600_final((uint2 *)state); // copy out copy(hash.uint64s, state, 4); diff --git a/libethash-cuda/dagger_shuffled.cuh b/libethash-cuda/dagger_shuffled.cuh index 90b78bb49..641db5173 100644 --- a/libethash-cuda/dagger_shuffled.cuh +++ b/libethash-cuda/dagger_shuffled.cuh @@ -9,19 +9,14 @@ __device__ uint64_t compute_hash_shuffle( { // sha3_512(header .. nonce) uint2 state[25]; - + state[0] = d_header.uint2s[0]; state[1] = d_header.uint2s[1]; state[2] = d_header.uint2s[2]; state[3] = d_header.uint2s[3]; state[4] = vectorize(nonce); - state[5] = vectorize(0x0000000000000001ULL); - for (uint32_t i = 6; i < 25; i++) - { - state[i] = make_uint2(0, 0); - } - state[8] = vectorize(0x8000000000000000ULL); - keccak_f1600_block(state,8); + + keccak_f1600_init(state); // Threads work together in this phase in groups of 8. const int thread_id = threadIdx.x & (THREADS_PER_HASH - 1); @@ -94,13 +89,7 @@ __device__ uint64_t compute_hash_shuffle( } // keccak_256(keccak_512(header..nonce) .. mix); - state[12] = vectorize(0x0000000000000001ULL); - for (uint32_t i = 13; i < 25; i++) - { - state[i] = vectorize(0ULL); - } - state[16] = vectorize(0x8000000000000000); - keccak_f1600_block(state, 1); + keccak_f1600_final(state); return devectorize(state[0]); } \ No newline at end of file diff --git a/libethash-cuda/ethash_cuda_miner_kernel.h b/libethash-cuda/ethash_cuda_miner_kernel.h index 54d3f2530..e47a22a75 100644 --- a/libethash-cuda/ethash_cuda_miner_kernel.h +++ b/libethash-cuda/ethash_cuda_miner_kernel.h @@ -57,6 +57,7 @@ void run_ethash_search( uint64_t start_nonce ); + #define CUDA_SAFE_CALL(call) \ do { \ cudaError_t err = call; \ diff --git a/libethash-cuda/keccak.cuh b/libethash-cuda/keccak.cuh index ddee7ef99..260dd0eda 100644 --- a/libethash-cuda/keccak.cuh +++ b/libethash-cuda/keccak.cuh @@ -52,12 +52,132 @@ uint2 chi(const uint2 a, const uint2 b, const uint2 c) { #endif -__device__ __forceinline__ void keccak_f1600_block(uint2* s, uint32_t out_size) +__device__ __forceinline__ void keccak_f1600_init(uint2* s) { uint2 t[5], u, v; + s[5] = vectorize(0x0000000000000001ULL); + for (uint32_t i = 6; i < 25; i++) + { + s[i] = make_uint2(0, 0); + } + s[8] = vectorize(0x8000000000000000ULL); + + /* theta: c = a[0,i] ^ a[1,i] ^ .. a[4,i] */ + t[0] = s[0] ^ s[5]; + t[1] = s[1] ^ s[11]; + t[2] = s[2]; + t[3] = s[3] ^ s[8]; + t[4] = s[4]; + + /* theta: d[i] = c[i+4] ^ rotl(c[i+1],1) */ + /* theta: a[0,i], a[1,i], .. a[4,i] ^= d[i] */ + + u = ROL2(t[1], 1); + s[0] = xor3(s[0], t[4], u); + s[5] = xor3(s[5], t[4], u); + s[10] = xor3(s[10], t[4], u); + s[15] = xor3(s[15], t[4], u); + s[20] = xor3(s[20], t[4], u); + + u = ROL2(t[2], 1); + s[1] = xor3(s[1], t[0], u); + s[6] = xor3(s[6], t[0], u); + s[11] = xor3(s[11], t[0], u); + s[16] = xor3(s[16], t[0], u); + s[21] = xor3(s[21], t[0], u); + + u = ROL2(t[3], 1); + s[2] = xor3(s[2], t[1], u); + s[7] = xor3(s[7], t[1], u); + s[12] = xor3(s[12], t[1], u); + s[17] = xor3(s[17], t[1], u); + s[22] = xor3(s[22], t[1], u); + + u = ROL2(t[4], 1); + s[3] = xor3(s[3], t[2], u); + s[8] = xor3(s[8], t[2], u); + s[13] = xor3(s[13], t[2], u); + s[18] = xor3(s[18], t[2], u); + s[23] = xor3(s[23], t[2], u); + + + u = ROL2(t[0], 1); + s[4] = xor3(s[4], t[3], u); + s[9] = xor3(s[9], t[3], u); + s[14] = xor3(s[14], t[3], u); + s[19] = xor3(s[19], t[3], u); + s[24] = xor3(s[24], t[3], u); + + /* rho pi: b[..] = rotl(a[..], ..) */ + u = s[1]; - for (int i = 0; i < 24; i++) + s[1] = ROL2(s[6], 44); + s[6] = ROL2(s[9], 20); + s[9] = ROL2(s[22], 61); + s[22] = ROL2(s[14], 39); + s[14] = ROL2(s[20], 18); + s[20] = ROL2(s[2], 62); + s[2] = ROL2(s[12], 43); + s[12] = ROL2(s[13], 25); + s[13] = ROL2(s[19], 8); + s[19] = ROL2(s[23], 56); + s[23] = ROL2(s[15], 41); + s[15] = ROL2(s[4], 27); + s[4] = ROL2(s[24], 14); + s[24] = ROL2(s[21], 2); + s[21] = ROL2(s[8], 55); + s[8] = ROL2(s[16], 45); + s[16] = ROL2(s[5], 36); + s[5] = ROL2(s[3], 28); + s[3] = ROL2(s[18], 21); + s[18] = ROL2(s[17], 15); + s[17] = ROL2(s[11], 10); + s[11] = ROL2(s[7], 6); + s[7] = ROL2(s[10], 3); + s[10] = ROL2(u, 1); + + /* chi: a[i,j] ^= ~b[i,j+1] & b[i,j+2] */ + + u = s[0]; v = s[1]; + s[0] = chi(s[0], s[1], s[2]); + s[1] = chi(s[1], s[2], s[3]); + s[2] = chi(s[2], s[3], s[4]); + s[3] = chi(s[3], s[4], u); + s[4] = chi(s[4], u, v); + + u = s[5]; v = s[6]; + s[5] = chi(s[5], s[6], s[7]); + s[6] = chi(s[6], s[7], s[8]); + s[7] = chi(s[7], s[8], s[9]); + s[8] = chi(s[8], s[9], u); + s[9] = chi(s[9], u, v); + + u = s[10]; v = s[11]; + s[10] = chi(s[10], s[11], s[12]); + s[11] = chi(s[11], s[12], s[13]); + s[12] = chi(s[12], s[13], s[14]); + s[13] = chi(s[13], s[14], u); + s[14] = chi(s[14], u, v); + + u = s[15]; v = s[16]; + s[15] = chi(s[15], s[16], s[17]); + s[16] = chi(s[16], s[17], s[18]); + s[17] = chi(s[17], s[18], s[19]); + s[18] = chi(s[18], s[19], u); + s[19] = chi(s[19], u, v); + + u = s[20]; v = s[21]; + s[20] = chi(s[20], s[21], s[22]); + s[21] = chi(s[21], s[22], s[23]); + s[22] = chi(s[22], s[23], s[24]); + s[23] = chi(s[23], s[24], u); + s[24] = chi(s[24], u, v); + + /* iota: a[0,0] ^= round constant */ + s[0] ^= vectorize(keccak_round_constants[0]); + + for (int i = 1; i < 23; i++) { /* theta: c = a[0,i] ^ a[1,i] ^ .. a[4,i] */ t[0] = xor5(s[0] , s[5] , s[10] , s[15] , s[20]); @@ -133,34 +253,325 @@ __device__ __forceinline__ void keccak_f1600_block(uint2* s, uint32_t out_size) s[7] = ROL2(s[10], 3); s[10] = ROL2(u, 1); - // squeeze this in here /* chi: a[i,j] ^= ~b[i,j+1] & b[i,j+2] */ - u = s[0]; + + u = s[0]; v = s[1]; s[0] = chi(s[0], s[1], s[2]); + s[1] = chi(s[1], s[2], s[3]); + s[2] = chi(s[2], s[3], s[4]); + s[3] = chi(s[3], s[4], u); + s[4] = chi(s[4], u, v); + + u = s[5]; v = s[6]; + s[5] = chi(s[5], s[6], s[7]); + s[6] = chi(s[6], s[7], s[8]); + s[7] = chi(s[7], s[8], s[9]); + s[8] = chi(s[8], s[9], u); + s[9] = chi(s[9], u, v); + + u = s[10]; v = s[11]; + s[10] = chi(s[10], s[11], s[12]); + s[11] = chi(s[11], s[12], s[13]); + s[12] = chi(s[12], s[13], s[14]); + s[13] = chi(s[13], s[14], u); + s[14] = chi(s[14], u, v); + + u = s[15]; v = s[16]; + s[15] = chi(s[15], s[16], s[17]); + s[16] = chi(s[16], s[17], s[18]); + s[17] = chi(s[17], s[18], s[19]); + s[18] = chi(s[18], s[19], u); + s[19] = chi(s[19], u, v); + + u = s[20]; v = s[21]; + s[20] = chi(s[20], s[21], s[22]); + s[21] = chi(s[21], s[22], s[23]); + s[22] = chi(s[22], s[23], s[24]); + s[23] = chi(s[23], s[24], u); + s[24] = chi(s[24], u, v); /* iota: a[0,0] ^= round constant */ s[0] ^= vectorize(keccak_round_constants[i]); - if (i == 23 && out_size == 1) return; + } + + /* theta: c = a[0,i] ^ a[1,i] ^ .. a[4,i] */ + t[0] = xor5(s[0], s[5], s[10], s[15], s[20]); + t[1] = xor5(s[1], s[6], s[11], s[16], s[21]); + t[2] = xor5(s[2], s[7], s[12], s[17], s[22]); + t[3] = xor5(s[3], s[8], s[13], s[18], s[23]); + t[4] = xor5(s[4], s[9], s[14], s[19], s[24]); + + /* theta: d[i] = c[i+4] ^ rotl(c[i+1],1) */ + /* theta: a[0,i], a[1,i], .. a[4,i] ^= d[i] */ + + u = ROL2(t[1], 1); + s[0] = xor3(s[0], t[4], u); + s[10] = xor3(s[10], t[4], u); + + u = ROL2(t[2], 1); + s[6] = xor3(s[6], t[0], u); + s[16] = xor3(s[16], t[0], u); + + + u = ROL2(t[3], 1); + s[12] = xor3(s[12], t[1], u); + s[22] = xor3(s[22], t[1], u); + + u = ROL2(t[4], 1); + s[3] = xor3(s[3], t[2], u); + s[18] = xor3(s[18], t[2], u); + + u = ROL2(t[0], 1); + s[9] = xor3(s[9], t[3], u); + s[24] = xor3(s[24], t[3], u); + + /* rho pi: b[..] = rotl(a[..], ..) */ + u = s[1]; + + s[1] = ROL2(s[6], 44); + s[6] = ROL2(s[9], 20); + s[9] = ROL2(s[22], 61); + s[2] = ROL2(s[12], 43); + s[4] = ROL2(s[24], 14); + s[8] = ROL2(s[16], 45); + s[5] = ROL2(s[3], 28); + s[3] = ROL2(s[18], 21); + s[7] = ROL2(s[10], 3); + + /* chi: a[i,j] ^= ~b[i,j+1] & b[i,j+2] */ + + u = s[0]; v = s[1]; + s[0] = chi(s[0], s[1], s[2]); + s[1] = chi(s[1], s[2], s[3]); + s[2] = chi(s[2], s[3], s[4]); + s[3] = chi(s[3], s[4], u); + s[4] = chi(s[4], u, v); + + u = s[5]; v = s[6]; + s[5] = chi(s[5], s[6], s[7]); + s[6] = chi(s[6], s[7], s[8]); + s[7] = chi(s[7], s[8], s[9]); + + /* iota: a[0,0] ^= round constant */ + s[0] ^= vectorize(keccak_round_constants[23]); +} + +__device__ __forceinline__ void keccak_f1600_final(uint2* s) +{ + uint2 t[5], u, v; + + s[12] = vectorize(0x0000000000000001ULL); + for (uint32_t i = 13; i < 25; i++) + { + s[i] = vectorize(0ULL); + } + s[16] = vectorize(0x8000000000000000); + + /* theta: c = a[0,i] ^ a[1,i] ^ .. a[4,i] */ + t[0] = xor3(s[0], s[5], s[10]); + t[1] = xor3(s[1], s[6], s[11]) ^ s[16]; + t[2] = xor3(s[2], s[7], s[12]); + t[3] = s[3] ^ s[8]; + t[4] = s[4] ^ s[9]; + + /* theta: d[i] = c[i+4] ^ rotl(c[i+1],1) */ + /* theta: a[0,i], a[1,i], .. a[4,i] ^= d[i] */ + + u = ROL2(t[1], 1); + s[0] = xor3(s[0], t[4], u); + s[5] = xor3(s[5], t[4], u); + s[10] = xor3(s[10], t[4], u); + s[15] = xor3(s[15], t[4], u); + s[20] = xor3(s[20], t[4], u); + + u = ROL2(t[2], 1); + s[1] = xor3(s[1], t[0], u); + s[6] = xor3(s[6], t[0], u); + s[11] = xor3(s[11], t[0], u); + s[16] = xor3(s[16], t[0], u); + s[21] = xor3(s[21], t[0], u); + + u = ROL2(t[3], 1); + s[2] = xor3(s[2], t[1], u); + s[7] = xor3(s[7], t[1], u); + s[12] = xor3(s[12], t[1], u); + s[17] = xor3(s[17], t[1], u); + s[22] = xor3(s[22], t[1], u); + + u = ROL2(t[4], 1); + s[3] = xor3(s[3], t[2], u); + s[8] = xor3(s[8], t[2], u); + s[13] = xor3(s[13], t[2], u); + s[18] = xor3(s[18], t[2], u); + s[23] = xor3(s[23], t[2], u); + + + u = ROL2(t[0], 1); + s[4] = xor3(s[4], t[3], u); + s[9] = xor3(s[9], t[3], u); + s[14] = xor3(s[14], t[3], u); + s[19] = xor3(s[19], t[3], u); + s[24] = xor3(s[24], t[3], u); + + /* rho pi: b[..] = rotl(a[..], ..) */ + u = s[1]; + + s[1] = ROL2(s[6], 44); + s[6] = ROL2(s[9], 20); + s[9] = ROL2(s[22], 61); + s[22] = ROL2(s[14], 39); + s[14] = ROL2(s[20], 18); + s[20] = ROL2(s[2], 62); + s[2] = ROL2(s[12], 43); + s[12] = ROL2(s[13], 25); + s[13] = ROL2(s[19], 8); + s[19] = ROL2(s[23], 56); + s[23] = ROL2(s[15], 41); + s[15] = ROL2(s[4], 27); + s[4] = ROL2(s[24], 14); + s[24] = ROL2(s[21], 2); + s[21] = ROL2(s[8], 55); + s[8] = ROL2(s[16], 45); + s[16] = ROL2(s[5], 36); + s[5] = ROL2(s[3], 28); + s[3] = ROL2(s[18], 21); + s[18] = ROL2(s[17], 15); + s[17] = ROL2(s[11], 10); + s[11] = ROL2(s[7], 6); + s[7] = ROL2(s[10], 3); + s[10] = ROL2(u, 1); + + /* chi: a[i,j] ^= ~b[i,j+1] & b[i,j+2] */ + u = s[0]; v = s[1]; + s[0] = chi(s[0], s[1], s[2]); + s[1] = chi(s[1], s[2], s[3]); + s[2] = chi(s[2], s[3], s[4]); + s[3] = chi(s[3], s[4], u); + s[4] = chi(s[4], u, v); + + u = s[5]; v = s[6]; + s[5] = chi(s[5], s[6], s[7]); + s[6] = chi(s[6], s[7], s[8]); + s[7] = chi(s[7], s[8], s[9]); + s[8] = chi(s[8], s[9], u); + s[9] = chi(s[9], u, v); + + u = s[10]; v = s[11]; + s[10] = chi(s[10], s[11], s[12]); + s[11] = chi(s[11], s[12], s[13]); + s[12] = chi(s[12], s[13], s[14]); + s[13] = chi(s[13], s[14], u); + s[14] = chi(s[14], u, v); - v = s[1]; - // continue chi + u = s[15]; v = s[16]; + s[15] = chi(s[15], s[16], s[17]); + s[16] = chi(s[16], s[17], s[18]); + s[17] = chi(s[17], s[18], s[19]); + s[18] = chi(s[18], s[19], u); + s[19] = chi(s[19], u, v); + + u = s[20]; v = s[21]; + s[20] = chi(s[20], s[21], s[22]); + s[21] = chi(s[21], s[22], s[23]); + s[22] = chi(s[22], s[23], s[24]); + s[23] = chi(s[23], s[24], u); + s[24] = chi(s[24], u, v); + + /* iota: a[0,0] ^= round constant */ + s[0] ^= vectorize(keccak_round_constants[0]); + + for (int i = 1; i < 23; i++) + { + /* theta: c = a[0,i] ^ a[1,i] ^ .. a[4,i] */ + t[0] = xor5(s[0], s[5], s[10], s[15], s[20]); + t[1] = xor5(s[1], s[6], s[11], s[16], s[21]); + t[2] = xor5(s[2], s[7], s[12], s[17], s[22]); + t[3] = xor5(s[3], s[8], s[13], s[18], s[23]); + t[4] = xor5(s[4], s[9], s[14], s[19], s[24]); + + /* theta: d[i] = c[i+4] ^ rotl(c[i+1],1) */ + /* theta: a[0,i], a[1,i], .. a[4,i] ^= d[i] */ + + u = ROL2(t[1], 1); + s[0] = xor3(s[0], t[4], u); + s[5] = xor3(s[5], t[4], u); + s[10] = xor3(s[10], t[4], u); + s[15] = xor3(s[15], t[4], u); + s[20] = xor3(s[20], t[4], u); + + u = ROL2(t[2], 1); + s[1] = xor3(s[1], t[0], u); + s[6] = xor3(s[6], t[0], u); + s[11] = xor3(s[11], t[0], u); + s[16] = xor3(s[16], t[0], u); + s[21] = xor3(s[21], t[0], u); + + u = ROL2(t[3], 1); + s[2] = xor3(s[2], t[1], u); + s[7] = xor3(s[7], t[1], u); + s[12] = xor3(s[12], t[1], u); + s[17] = xor3(s[17], t[1], u); + s[22] = xor3(s[22], t[1], u); + + u = ROL2(t[4], 1); + s[3] = xor3(s[3], t[2], u); + s[8] = xor3(s[8], t[2], u); + s[13] = xor3(s[13], t[2], u); + s[18] = xor3(s[18], t[2], u); + s[23] = xor3(s[23], t[2], u); + + + u = ROL2(t[0], 1); + s[4] = xor3(s[4], t[3], u); + s[9] = xor3(s[9], t[3], u); + s[14] = xor3(s[14], t[3], u); + s[19] = xor3(s[19], t[3], u); + s[24] = xor3(s[24], t[3], u); + + /* rho pi: b[..] = rotl(a[..], ..) */ + u = s[1]; + + s[1] = ROL2(s[6], 44); + s[6] = ROL2(s[9], 20); + s[9] = ROL2(s[22], 61); + s[22] = ROL2(s[14], 39); + s[14] = ROL2(s[20], 18); + s[20] = ROL2(s[2], 62); + s[2] = ROL2(s[12], 43); + s[12] = ROL2(s[13], 25); + s[13] = ROL2(s[19], 8); + s[19] = ROL2(s[23], 56); + s[23] = ROL2(s[15], 41); + s[15] = ROL2(s[4], 27); + s[4] = ROL2(s[24], 14); + s[24] = ROL2(s[21], 2); + s[21] = ROL2(s[8], 55); + s[8] = ROL2(s[16], 45); + s[16] = ROL2(s[5], 36); + s[5] = ROL2(s[3], 28); + s[3] = ROL2(s[18], 21); + s[18] = ROL2(s[17], 15); + s[17] = ROL2(s[11], 10); + s[11] = ROL2(s[7], 6); + s[7] = ROL2(s[10], 3); + s[10] = ROL2(u, 1); + + /* chi: a[i,j] ^= ~b[i,j+1] & b[i,j+2] */ + u = s[0]; v = s[1]; + s[0] = chi(s[0], s[1], s[2]); s[1] = chi(s[1], s[2], s[3]); s[2] = chi(s[2], s[3], s[4]); s[3] = chi(s[3], s[4], u); s[4] = chi(s[4], u, v); - u = s[5]; v = s[6]; - + u = s[5]; v = s[6]; s[5] = chi(s[5], s[6], s[7]); s[6] = chi(s[6], s[7], s[8]); s[7] = chi(s[7], s[8], s[9]); - - if (i == 23) return; - s[8] = chi(s[8], s[9], u); s[9] = chi(s[9], u, v); - u = s[10]; v = s[11]; + u = s[10]; v = s[11]; s[10] = chi(s[10], s[11], s[12]); s[11] = chi(s[11], s[12], s[13]); s[12] = chi(s[12], s[13], s[14]); @@ -180,5 +591,26 @@ __device__ __forceinline__ void keccak_f1600_block(uint2* s, uint32_t out_size) s[22] = chi(s[22], s[23], s[24]); s[23] = chi(s[23], s[24], u); s[24] = chi(s[24], u, v); + + /* iota: a[0,0] ^= round constant */ + s[0] ^= vectorize(keccak_round_constants[i]); } + + t[0] = xor5(s[0], s[5], s[10], s[15], s[20]); + t[1] = xor5(s[1], s[6], s[11], s[16], s[21]); + t[2] = xor5(s[2], s[7], s[12], s[17], s[22]); + t[3] = xor5(s[3], s[8], s[13], s[18], s[23]); + t[4] = xor5(s[4], s[9], s[14], s[19], s[24]); + + s[0] = xor3(s[0], t[4], ROL2(t[1], 1)); + s[6] = xor3(s[6], t[0], ROL2(t[2], 1)); + s[12] = xor3(s[12], t[1], ROL2(t[3], 1)); + + s[1] = ROL2(s[6], 44); + s[2] = ROL2(s[12], 43); + + s[0] = chi(s[0], s[1], s[2]); + + /* iota: a[0,0] ^= round constant */ + s[0] ^= vectorize(keccak_round_constants[23]); } \ No newline at end of file