From 54b898cb788a461302eb6318664991bd28b75146 Mon Sep 17 00:00:00 2001 From: orignal Date: Mon, 8 May 2017 11:39:40 -0400 Subject: [PATCH] eliminate extra copy --- gost/cuda_gosthash.cu | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/gost/cuda_gosthash.cu b/gost/cuda_gosthash.cu index a799325..e3e761c 100644 --- a/gost/cuda_gosthash.cu +++ b/gost/cuda_gosthash.cu @@ -684,20 +684,22 @@ static void GOST_F(uint64_t* state) __device__ -static void GOST_E12(uint64_t* K, uint64_t *state) +static void GOST_E12(uint64_t* const K, uint64_t *state) { + uint64_t state1[8], K1[8]; + GOST_Copy512(K1, K); //#pragma unroll 1 for(int i=0; i<12; i++) { - GOST_F(state); - GOST_Xor512(K, K, CC[i]); - GOST_F(K); - GOST_Xor512(state, state, K); + GOST_FS(state, state1); + GOST_Xor512(K, K1, CC[i]); + GOST_FS(K, K1); + GOST_Xor512(state, state1, K1); } } __device__ -void GOST_E(uint64_t* K, uint64_t* const m, uint64_t *state /* out only */) +void GOST_E(uint64_t* const K, uint64_t* const m, uint64_t *state /* out only */) { GOST_Xor512(state, m, K); // state = m ^ K GOST_E12(K, state);