From d7acec6dcc43783f1b28b79a947ed80305d0ee50 Mon Sep 17 00:00:00 2001 From: Jan Willem Penterman Date: Fri, 25 Sep 2015 18:03:12 +0200 Subject: [PATCH] minor keccak change --- libethash-cuda/CMakeLists.txt | 2 +- libethash-cuda/dagger_shuffled.cuh | 4 +--- libethash-cuda/keccak.cuh | 8 +++----- 3 files changed, 5 insertions(+), 9 deletions(-) diff --git a/libethash-cuda/CMakeLists.txt b/libethash-cuda/CMakeLists.txt index 9fed4faa6..8b697e80d 100644 --- a/libethash-cuda/CMakeLists.txt +++ b/libethash-cuda/CMakeLists.txt @@ -13,7 +13,7 @@ LIST(APPEND CUDA_NVCC_FLAGS_DEBUG -G) if(COMPUTE AND (COMPUTE GREATER 0)) LIST(APPEND CUDA_NVCC_FLAGS -gencode arch=compute_${COMPUTE},code=sm_${COMPUTE}) else(COMPUTE AND (COMPUTE GREATER 0)) - set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_20,code=sm_20;-gencode arch=compute_30,code=sm_30;-gencode arch=compute_35,code=sm_35;-gencode arch=compute_50,code=sm_50;-gencode arch=compute_52,code=sm_52) + set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_35,code=sm_35;-gencode arch=compute_50,code=sm_50) endif(COMPUTE AND (COMPUTE GREATER 0)) diff --git a/libethash-cuda/dagger_shuffled.cuh b/libethash-cuda/dagger_shuffled.cuh index c046e5d44..e815449f8 100644 --- a/libethash-cuda/dagger_shuffled.cuh +++ b/libethash-cuda/dagger_shuffled.cuh @@ -85,7 +85,5 @@ __device__ uint64_t compute_hash_shuffle( } // keccak_256(keccak_512(header..nonce) .. mix); - keccak_f1600_final(state); - - return devectorize(state[0]); + return keccak_f1600_final(state); } \ No newline at end of file diff --git a/libethash-cuda/keccak.cuh b/libethash-cuda/keccak.cuh index d66c512e9..836b39b67 100644 --- a/libethash-cuda/keccak.cuh +++ b/libethash-cuda/keccak.cuh @@ -289,7 +289,6 @@ __device__ __forceinline__ void keccak_f1600_init(uint2* s) 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); @@ -323,8 +322,6 @@ __device__ __forceinline__ void keccak_f1600_init(uint2* s) 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]); @@ -333,7 +330,7 @@ __device__ __forceinline__ void keccak_f1600_init(uint2* s) s[0] ^= vectorize(keccak_round_constants[23]); } -__device__ __forceinline__ void keccak_f1600_final(uint2* s) +__device__ __forceinline__ uint64_t keccak_f1600_final(uint2* s) { uint2 t[5], u, v; @@ -589,5 +586,6 @@ __device__ __forceinline__ void keccak_f1600_final(uint2* s) s[0] = chi(s[0], s[1], s[2]); /* iota: a[0,0] ^= round constant */ - s[0] ^= vectorize(keccak_round_constants[23]); + //s[0] ^= vectorize(keccak_round_constants[23]); + return devectorize(s[0]) ^ keccak_round_constants[23]; } \ No newline at end of file