lukechilds
/
ethminer
mirror of https://github.com/lukechilds/ethminer.git
1
0
Fork 0
Code Issues Projects Releases Wiki Activity
Browse Source

trying a faster keccak.

cl-refactor
Jan Willem Penterman 9 years ago
parent
8759cfe0e3
commit
109aaa783c
4 changed files with 452 additions and 48 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 22
      libethash-cuda/dagger_shared.cuh
  2. 19
      libethash-cuda/dagger_shuffled.cuh
  3. 1
      libethash-cuda/ethash_cuda_miner_kernel.h
  4. 458
      libethash-cuda/keccak.cuh

22
libethash-cuda/dagger_shared.cuh
View File

@ -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);

19
libethash-cuda/dagger_shuffled.cuh
View File

@ -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]);
}

1
libethash-cuda/ethash_cuda_miner_kernel.h
View File

@ -57,6 +57,7 @@ void run_ethash_search(
uint64_t start_nonce
);
#define CUDA_SAFE_CALL(call) \
do { \
cudaError_t err = call; \

458
libethash-cuda/keccak.cuh
View File

@ -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]);
}
Write Preview
Loading…
Cancel
Save