AtHeartEngineer/icicle
1
0
Fork 0
mirror of https://github.com/pseXperiments/icicle.git synced 2026-01-08 23:17:54 -05:00
Code Issues Packages Projects Releases Wiki Activity

moved initialization to device

Browse Source
This commit is contained in:
aviadingo
2024-08-05 12:49:31 +03:00
parent 4229b9d26c
commit 6ce4a6cf1e
2 changed files with 8 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
icicle/src/hash/blake2s/Makefile
View File

@@ -3,6 +3,7 @@ test_blake2s: test_blake2s.cu blake2s.cu
test_blake2s_hasher: test_blake2s_hasher.cu blake2s.cu
nvcc -o test_blake2s_hasher -I. -I../../../include test_blake2s_hasher.cu -g
./test_blake2s_hasher
test_blake2s_tree: test_tree.cu blake2s.cu ../../merkle-tree/merkle.cu
nvcc -DMERKLE_DEBUG -o test_blake2s_tree -I../../../include test_tree.cu

21
icicle/src/hash/blake2s/blake2s.cu
View File

@@ -225,11 +225,15 @@ __global__ void kernel_blake2s_hash(const BYTE *indata, WORD inlen, BYTE *outdat
if (thread >= n_batch) {
return;
}
BYTE key[32] = ""; // Null key
WORD keylen = 0;
CUDA_BLAKE2S_CTX blake_ctx;
const BYTE *in = indata + thread * inlen;
BYTE *out = outdata + thread * BLAKE2S_BLOCK_SIZE;
CUDA_BLAKE2S_CTX ctx = c_CTX;
cuda_blake2s_update(&ctx, in, inlen);
cuda_blake2s_final(&ctx, out);
cuda_blake2s_init(&blake_ctx, key, keylen, (BLAKE2S_BLOCK_SIZE << 3));
cuda_blake2s_update(&blake_ctx, in, inlen);
cuda_blake2s_final(&blake_ctx, out);
}
extern "C" {
@@ -269,22 +273,11 @@ cudaError_t Blake2s::run_hash_many_kernel(
WORD output_len,
const device_context::DeviceContext& ctx) const
{
BYTE key[32] = ""; // Null key
WORD keylen = strlen((char *)key);
const WORD BLAKE2S_BLOCK_SIZE = (output_len >> 3);
CUDA_BLAKE2S_CTX blake_ctx;
cpu_blake2s_init(&blake_ctx, key, keylen, output_len);
cudaMemcpyToSymbol(c_CTX, &blake_ctx, sizeof(CUDA_BLAKE2S_CTX), 0, cudaMemcpyHostToDevice);
WORD thread = 256;
WORD block = (number_of_states + thread - 1) / thread;
kernel_blake2s_hash<<<block, thread, 0, ctx.stream>>>(input, input_len, output, number_of_states, BLAKE2S_BLOCK_SIZE);
// cudaDeviceSynchronize();
// cudaError_t error = cudaGetLastError();
// if (error != cudaSuccess) {
// printf("Error cuda blake2s hash: %s \n", cudaGetErrorString(error));
// }
CHK_IF_RETURN(cudaPeekAtLastError());
return CHK_LAST();