Aviad/blake2s #576
Aviad/blake2s #576
Conversation
| */ | ||
|
|
||
| #pragma once | ||
| typedef unsigned char BYTE; |
There was a problem hiding this comment.
please move these inside blake2s namespace
icicle/src/merkle-tree/merkle.cu
Outdated
| THROW_ICICLE_ERR( | ||
| IcicleError_t::InvalidArgument, | ||
| "Hash max preimage length does not match merkle tree arity multiplied by digest elements"); | ||
| // if (compression.preimage_max_length < tree_config.arity * tree_config.digest_elements) |
There was a problem hiding this comment.
You can just delete those at this point
icicle/src/hash/blake2s/blake2s.cu
Outdated
| __device__ __forceinline__ void cuda_blake2s_init_state(cuda_blake2s_ctx_t* ctx) | ||
| { | ||
| memcpy(ctx->state, ctx->chain, BLAKE2S_CHAIN_LENGTH); | ||
| // ctx->state[8] = ctx->t0; |
There was a problem hiding this comment.
Why are these commented?
There was a problem hiding this comment.
it's from an old version. deleted
icicle/src/hash/blake2s/blake2s.cu
Outdated
| return a; | ||
| } | ||
|
|
||
| __device__ uint32_t cuda_blake2s_ROTR32(uint32_t a, uint8_t b) { return (a >> b) | (a << (32 - b)); } |
There was a problem hiding this comment.
Maybe worth to add __inline__
icicle/src/hash/blake2s/blake2s.cu
Outdated
| cudaMalloc(&cuda_outdata, BLAKE2S_BLOCK_SIZE * n_batch); | ||
| assert(keylen <= 32); | ||
|
|
||
| // CUDA_BLAKE2S_CTX ctx; |
There was a problem hiding this comment.
These should be removed
| WORD block = (n_batch + thread - 1) / thread; | ||
| kernel_blake2s_hash<<<block, thread>>>(cuda_indata, inlen, cuda_outdata, n_batch, BLAKE2S_BLOCK_SIZE); | ||
| cudaMemcpy(out, cuda_outdata, BLAKE2S_BLOCK_SIZE * n_batch, cudaMemcpyDeviceToHost); | ||
| cudaDeviceSynchronize(); |
There was a problem hiding this comment.
Current implementation does not support async. All of our other primitives do. So maybe worth adding HashConfig as an input and changing all the functions to their async alternatives
There was a problem hiding this comment.
mcm_cuda_blake2s_hash_batch is not a part of the API, it's only for unit tests.
run_hash_many_kernel() is used in the same way as our other implementations
icicle/src/hash/blake2s/blake2s.cu
Outdated
| kernel_blake2s_hash<<<block, thread>>>(cuda_indata, inlen, cuda_outdata, n_batch, BLAKE2S_BLOCK_SIZE); | ||
| cudaMemcpy(out, cuda_outdata, BLAKE2S_BLOCK_SIZE * n_batch, cudaMemcpyDeviceToHost); | ||
| cudaDeviceSynchronize(); | ||
| cudaError_t error = cudaGetLastError(); |
There was a problem hiding this comment.
Please use our error-management functions (you can find an example in any of our primitives)
| extern "C" cudaError_t blake2s_cuda( | ||
| BYTE* input, BYTE* output, WORD number_of_blocks, WORD input_block_size, WORD output_block_size, HashConfig& config) | ||
| { |
There was a problem hiding this comment.
I noticed that the parameter order in blake2s_cuda is a bit different from keccak256_cuda (and other keccak functions).
keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, HashConfig& config)To keep things consistent, so my suggestion smt like this:
blake2s_cuda(BYTE* input, WORD input_block_size, WORD number_of_blocks, WORD output_block_size, BYTE* output, HashConfig& config)There was a problem hiding this comment.
you are right. I'm not sure if it is worth changing it since this is usually not exposed to the user and it will require updates to the rust signatures as well
## Changes This PR updates the existing Blake2s implementation by integrating a C++ API and a Rust wrapper API.
This reverts commit 166c615.
* stwo blake2s mmcs * cleanup * add rust wrappers * fix old blake2s mode
Describe the changes
adds Blake2s cuda capability.