From 25e0ab280b443a398a8273eb9937de40fa8c851b Mon Sep 17 00:00:00 2001 From: orignal Date: Wed, 26 Apr 2017 13:41:52 -0400 Subject: [PATCH] correct GOST_Add512 implementation --- gost/cuda_gosthash.cu | 47 +++++++------------------------------------ 1 file changed, 7 insertions(+), 40 deletions(-) diff --git a/gost/cuda_gosthash.cu b/gost/cuda_gosthash.cu index caf92a8..b6712f3 100644 --- a/gost/cuda_gosthash.cu +++ b/gost/cuda_gosthash.cu @@ -570,7 +570,13 @@ void ADD_ASM_512_32(uint32_t* x, uint32_t* const a, uint32_t* const b) __device__ __forceinline__ void GOST_Add512(void *x, void * const a, void * const b) { - ADD_ASM_512_32((uint32_t*)x, (uint32_t*)a, (uint32_t*)b); + //ADD_ASM_512_32((uint32_t*)x, (uint32_t*)a, (uint32_t*)b); + uint16_t t = 0; + for(int i = 63; i >= 0; i--) + { + t = ((uint8_t *)a)[i] + ((uint8_t *)b)[i] + (t >> 8); + ((uint8_t *)x)[i] = t & 0xFF; + } } __device__ __forceinline__ @@ -973,20 +979,6 @@ void GOST_E(uint64_t* K, uint64_t* const m, uint64_t *state /* out only */) GOST_E12(K, state); } -__device__ -void GOST_g_0(uint64_t* h, uint64_t* const M) -{ - uint64_t K0[8]; - // GOST_F(0); - K0[0] = K0[1] = K0[2] = K0[3] = 0x74a5d4ce2efc83b3; - K0[4] = K0[5] = K0[6] = K0[7] = 0x74a5d4ce2efc83b3; - - uint64_t t[8]; - GOST_Xor512(t, M, K0); // t = M ^ K0 - GOST_E12(K0, t); - GOST_Xor512(h, t, M); -} - __device__ void GOST_g_N(uint64_t* h, uint64_t* const M, uint64_t* const N) { @@ -1003,31 +995,6 @@ void GOST_g_N(uint64_t* h, uint64_t* const M, uint64_t* const N) GOST_Xor512_3(h, t, M); // h = h ^ t ^ M } -__device__ -void GOST_g_1(uint64_t* h, uint64_t* N) -{ - uint64_t K[8]; - GOST_FS(h, K); // K = F(h) - GOST_Xor512(h, h, N); - - uint64_t t[8]; - GOST_E(K, N, t); - GOST_Xor512(h, h, t); -} - -__device__ -void GOST_g_F(uint64_t* out, uint64_t* h, uint64_t* M) -{ - uint64_t K[8]; - GOST_FS(h, K); // K = F(h) - GOST_Add512(M, M, out); - - uint64_t t[8]; - GOST_E(K, M, t); - GOST_Xor512(t, t, M); - - GOST_Xor512(out, t, h); -} __device__ __forceinline__ void GOST_hash_X(uint64_t *hash, uchar * const message, uint64_t len)