Skip to content

Commit

Permalink
keccak: avoid to use twice cuda_default_throughput
Browse files Browse the repository at this point in the history 
and drop useless gpu hash alloc...
  • Loading branch information
@tpruvot
tpruvot committed Jan 4, 2018
1 parent 11a512f commit 73dd6aa
Show file tree
Hide file tree
Showing 2 changed files with 11 additions and 23 deletions.
20 changes: 6 additions & 14 deletions Algo256/cuda_keccak256_sm3.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,6 @@ static const uint64_t host_keccak_round_constants[24] = {
0x0000000080000001ull, 0x8000000080008008ull
};

static uint32_t *d_nounce[MAX_GPUS];
static uint32_t *d_KNonce[MAX_GPUS];

__constant__ uint32_t pTarget[8];
Expand Down Expand Up @@ -170,7 +169,7 @@ static void keccak_blockv30(uint64_t *s, const uint64_t *keccak_round_constants)
#endif

__global__ __launch_bounds__(128,5)
void keccak256_sm3_gpu_hash_80(uint32_t threads, uint32_t startNounce, void *outputHash, uint32_t *resNounce)
void keccak256_sm3_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t *resNounce)
{
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
Expand Down Expand Up @@ -210,25 +209,20 @@ void keccak256_sm3_gpu_hash_80(uint32_t threads, uint32_t startNounce, void *out
}

__host__
uint32_t keccak256_sm3_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_outputHash, int order)
void keccak256_sm3_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *resNonces, int order)
{
uint32_t result = UINT32_MAX;
cudaMemset(d_KNonce[thr_id], 0xff, sizeof(uint32_t));
cudaMemset(d_KNonce[thr_id], 0xff, 2*sizeof(uint32_t));
const uint32_t threadsperblock = 128;

dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);

size_t shared_size = 0;

keccak256_sm3_gpu_hash_80<<<grid, block, shared_size>>>(threads, startNounce, d_outputHash, d_KNonce[thr_id]);
keccak256_sm3_gpu_hash_80<<<grid, block, shared_size>>>(threads, startNounce, d_KNonce[thr_id]);

MyStreamSynchronize(NULL, order, thr_id);
cudaMemcpy(d_nounce[thr_id], d_KNonce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost);
cudaMemcpy(resNonces, d_KNonce[thr_id], 2*sizeof(uint32_t), cudaMemcpyDeviceToHost);
cudaThreadSynchronize();
result = *d_nounce[thr_id];

return result;
}

#if 0
Expand Down Expand Up @@ -299,13 +293,11 @@ void keccak256_sm3_init(int thr_id, uint32_t threads)
{
CUDA_SAFE_CALL(cudaMemcpyToSymbol(keccak_round_constants, host_keccak_round_constants,
sizeof(host_keccak_round_constants), 0, cudaMemcpyHostToDevice));
CUDA_SAFE_CALL(cudaMalloc(&d_KNonce[thr_id], sizeof(uint32_t)));
CUDA_SAFE_CALL(cudaMallocHost(&d_nounce[thr_id], 1*sizeof(uint32_t)));
CUDA_SAFE_CALL(cudaMalloc(&d_KNonce[thr_id], 2*sizeof(uint32_t)));
}

__host__
void keccak256_sm3_free(int thr_id)
{
cudaFree(d_KNonce[thr_id]);
cudaFreeHost(d_nounce[thr_id]);
}
14 changes: 5 additions & 9 deletions Algo256/keccak256.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -14,8 +14,6 @@ extern "C"

#include "cuda_helper.h"

static uint32_t *d_hash[MAX_GPUS];

// SM5+ cuda
extern void keccak256_cpu_init(int thr_id);
extern void keccak256_cpu_free(int thr_id);
Expand All @@ -27,7 +25,7 @@ extern void keccak256_setOutput(int thr_id);
extern void keccak256_sm3_init(int thr_id, uint32_t threads);
extern void keccak256_sm3_free(int thr_id);
extern void keccak256_sm3_setBlock_80(void *pdata, const void *ptarget);
extern uint32_t keccak256_sm3_hash_80(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_hash, int order);
extern uint32_t keccak256_sm3_hash_80(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t* resNonces, int order);

// CPU Hash
extern "C" void keccak256_hash(void *state, const void *input)
Expand All @@ -52,13 +50,13 @@ extern "C" int scanhash_keccak256(int thr_id, struct work* work, uint32_t max_no
uint32_t *ptarget = work->target;
const uint32_t first_nonce = pdata[19];
const int dev_id = device_map[thr_id];
uint32_t throughput = cuda_default_throughput(thr_id, 1U << 21); // 256*256*8*4
uint32_t throughput;
uint32_t intensity = 23;
if(!use_compat_kernels[thr_id]) {
uint32_t intensity = 23;
if (strstr(device_name[dev_id], "GTX 1070")) intensity = 25;
if (strstr(device_name[dev_id], "GTX 1080")) intensity = 26;
throughput = cuda_default_throughput(thr_id, 1U << intensity);
}
throughput = cuda_default_throughput(thr_id, 1U << intensity);
if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);

if (opt_benchmark)
Expand All @@ -80,7 +78,6 @@ extern "C" int scanhash_keccak256(int thr_id, struct work* work, uint32_t max_no
keccak256_cpu_init(thr_id);
} else {
// really useful ?
CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], throughput * 64));
keccak256_sm3_init(thr_id, throughput);
}

Expand All @@ -107,7 +104,7 @@ extern "C" int scanhash_keccak256(int thr_id, struct work* work, uint32_t max_no
*hashes_done = pdata[19] - first_nonce + throughput;

if(use_compat_kernels[thr_id])
work->nonces[0] = keccak256_sm3_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
keccak256_sm3_hash_80(thr_id, throughput, pdata[19], work->nonces, order++);
else {
keccak256_cpu_hash_80(thr_id, throughput, pdata[19], work->nonces, highTarget);
}
Expand Down Expand Up @@ -170,7 +167,6 @@ extern "C" void free_keccak256(int thr_id)
if(!use_compat_kernels[thr_id])
keccak256_cpu_free(thr_id);
else {
cudaFree(d_hash[thr_id]);
keccak256_sm3_free(thr_id);
}

Expand Down

0 comments on commit 73dd6aa

Please sign in to comment.