|
|
@ -29,12 +29,12 @@ __constant__ uint32_t c_PaddedMessage2[18]; // 44 bytes of remaining message (No |
|
|
|
*p = (uint32_t)((v)); *(p+1) = (uint32_t)((v) >> 32); |
|
|
|
*p = (uint32_t)((v)); *(p+1) = (uint32_t)((v) >> 32); |
|
|
|
|
|
|
|
|
|
|
|
static __device__ void mycpy72(uint32_t *d, const uint32_t *s) { |
|
|
|
static __device__ void mycpy72(uint32_t *d, const uint32_t *s) { |
|
|
|
#pragma unroll 18 |
|
|
|
#pragma unroll 18 |
|
|
|
for (int k=0; k < 18; ++k) d[k] = s[k]; |
|
|
|
for (int k=0; k < 18; ++k) d[k] = s[k]; |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
|
static __device__ void mycpy32(uint32_t *d, const uint32_t *s) { |
|
|
|
static __device__ void mycpy32(uint32_t *d, const uint32_t *s) { |
|
|
|
#pragma unroll 8 |
|
|
|
#pragma unroll 8 |
|
|
|
for (int k=0; k < 8; ++k) d[k] = s[k]; |
|
|
|
for (int k=0; k < 8; ++k) d[k] = s[k]; |
|
|
|
} |
|
|
|
} |
|
|
|
|
|
|
|
|
|
|
@ -45,7 +45,7 @@ typedef struct keccak_hash_state_t { |
|
|
|
|
|
|
|
|
|
|
|
__device__ void statecopy(uint64_t *d, uint64_t *s) |
|
|
|
__device__ void statecopy(uint64_t *d, uint64_t *s) |
|
|
|
{ |
|
|
|
{ |
|
|
|
#pragma unroll 25 |
|
|
|
#pragma unroll 25 |
|
|
|
for (int i=0; i < 25; ++i) |
|
|
|
for (int i=0; i < 25; ++i) |
|
|
|
d[i] = s[i]; |
|
|
|
d[i] = s[i]; |
|
|
|
} |
|
|
|
} |
|
|
@ -74,7 +74,6 @@ keccak_block(uint64_t *s, const uint32_t *in, const uint64_t *keccak_round_const |
|
|
|
uint64_t t[5], u[5], v, w; |
|
|
|
uint64_t t[5], u[5], v, w; |
|
|
|
|
|
|
|
|
|
|
|
/* absorb input */ |
|
|
|
/* absorb input */ |
|
|
|
#pragma unroll 9 |
|
|
|
|
|
|
|
for (i = 0; i < 9 /* 72/8 */; i++, in += 2) |
|
|
|
for (i = 0; i < 9 /* 72/8 */; i++, in += 2) |
|
|
|
s[i] ^= U32TO64_LE(in); |
|
|
|
s[i] ^= U32TO64_LE(in); |
|
|
|
|
|
|
|
|
|
|
|