From 634bea21f5e45de05c197b9b180aef4d26129f01 Mon Sep 17 00:00:00 2001 From: Tanguy Pruvot Date: Wed, 17 Jun 2015 03:39:07 +0200 Subject: [PATCH] luffa/cube: unroll 1 really required on the 9xx --- x11/cuda_x11_luffa512_Cubehash.cu | 44 +++++++++++++++---------------- 1 file changed, 21 insertions(+), 23 deletions(-) diff --git a/x11/cuda_x11_luffa512_Cubehash.cu b/x11/cuda_x11_luffa512_Cubehash.cu index 2fc868c..1e7b914 100644 --- a/x11/cuda_x11_luffa512_Cubehash.cu +++ b/x11/cuda_x11_luffa512_Cubehash.cu @@ -1,16 +1,9 @@ /* - * Merged LUFFA512 64 + CUBE512 64 + * Merged LUFFA512 64 + CUBE512 64 - from sp */ #include "cuda_helper.h" -typedef unsigned char BitSequence; - -typedef struct { - uint32_t buffer[8]; /* Buffer to be hashed */ - uint32_t chainv[40]; /* Chaining values */ -} hashState; - #define MULT0(a) {\ tmp = a[7]; \ a[7] = a[6]; \ @@ -204,7 +197,7 @@ static void rnd512(uint32_t *statebuffer, uint32_t *statechainv) chainv[i] = statechainv[i]; } - #pragma unroll + #pragma unroll 1 for(i=0;i<8;i++) { STEP(c_CNS[(2*i)],c_CNS[(2*i)+1]); } @@ -217,7 +210,7 @@ static void rnd512(uint32_t *statebuffer, uint32_t *statechainv) TWEAK(chainv[4],chainv[5],chainv[6],chainv[7],1); - #pragma unroll + #pragma unroll 1 for(i=0;i<8;i++) { STEP(c_CNS[(2*i)+16],c_CNS[(2*i)+16+1]); } @@ -230,7 +223,7 @@ static void rnd512(uint32_t *statebuffer, uint32_t *statechainv) TWEAK(chainv[4],chainv[5],chainv[6],chainv[7],2); - #pragma unroll + #pragma unroll 1 for(i=0;i<8;i++) { STEP(c_CNS[(2*i)+32],c_CNS[(2*i)+32+1]); } @@ -243,7 +236,7 @@ static void rnd512(uint32_t *statebuffer, uint32_t *statechainv) TWEAK(chainv[4],chainv[5],chainv[6],chainv[7],3); - #pragma unroll + #pragma unroll 1 for(i=0;i<8;i++) { STEP(c_CNS[(2*i)+48],c_CNS[(2*i)+48+1]); } @@ -256,7 +249,7 @@ static void rnd512(uint32_t *statebuffer, uint32_t *statechainv) TWEAK(chainv[4],chainv[5],chainv[6],chainv[7],4); - #pragma unroll + #pragma unroll 1 for(i=0;i<8;i++) { STEP(c_CNS[(2*i)+64],c_CNS[(2*i)+64+1]); } @@ -287,7 +280,7 @@ static void rnd512_first(uint32_t state[40], uint32_t buffer[8]) for (i = 0; i<8; i++) chainv[i] = state[i]; - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i)], c_CNS[(2 * i) + 1]); } @@ -300,7 +293,7 @@ static void rnd512_first(uint32_t state[40], uint32_t buffer[8]) TWEAK(chainv[4], chainv[5], chainv[6], chainv[7], 1); - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i) + 16], c_CNS[(2 * i) + 16 + 1]); } @@ -313,7 +306,7 @@ static void rnd512_first(uint32_t state[40], uint32_t buffer[8]) TWEAK(chainv[4], chainv[5], chainv[6], chainv[7], 2); - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i) + 32], c_CNS[(2 * i) + 32 + 1]); } @@ -326,7 +319,7 @@ static void rnd512_first(uint32_t state[40], uint32_t buffer[8]) TWEAK(chainv[4], chainv[5], chainv[6], chainv[7], 3); - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i) + 48], c_CNS[(2 * i) + 48 + 1]); } @@ -339,7 +332,7 @@ static void rnd512_first(uint32_t state[40], uint32_t buffer[8]) TWEAK(chainv[4], chainv[5], chainv[6], chainv[7], 4); - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i) + 64], c_CNS[(2 * i) + 64 + 1]); } @@ -419,7 +412,7 @@ static void rnd512_nullhash(uint32_t *state) for (i = 0; i<8; i++) chainv[i] = state[i]; - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i)], c_CNS[(2 * i) + 1]); } @@ -432,7 +425,7 @@ static void rnd512_nullhash(uint32_t *state) TWEAK(chainv[4], chainv[5], chainv[6], chainv[7], 1); - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i) + 16], c_CNS[(2 * i) + 16 + 1]); } @@ -445,7 +438,7 @@ static void rnd512_nullhash(uint32_t *state) TWEAK(chainv[4], chainv[5], chainv[6], chainv[7], 2); - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i) + 32], c_CNS[(2 * i) + 32 + 1]); } @@ -458,7 +451,7 @@ static void rnd512_nullhash(uint32_t *state) TWEAK(chainv[4], chainv[5], chainv[6], chainv[7], 3); - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i) + 48], c_CNS[(2 * i) + 48 + 1]); } @@ -471,7 +464,7 @@ static void rnd512_nullhash(uint32_t *state) TWEAK(chainv[4], chainv[5], chainv[6], chainv[7], 4); - #pragma unroll + #pragma unroll 1 for (i = 0; i<8; i++) { STEP(c_CNS[(2 * i) + 64], c_CNS[(2 * i) + 64 + 1]); } @@ -736,6 +729,11 @@ static void finalization512(uint32_t *statebuffer, uint32_t *statechainv, uint32 } __global__ +#if __CUDA_ARCH__ > 500 +__launch_bounds__(256, 4) +#else +__launch_bounds__(256, 3) +#endif void x11_luffaCubehash512_gpu_hash_64(uint32_t threads, uint32_t *g_hash) { const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);