Aviad/blake2s #576
Open
Aviad/blake2s #576
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
ChickenLover
reviewed
Sep 6, 2024
| */ | ||
|
|
||
| #pragma once | ||
| typedef unsigned char BYTE; |
There was a problem hiding this comment.
please move these inside blake2s namespace
| 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
| __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?
| 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__
| 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
| 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)
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Describe the changes
adds Blake2s cuda capability.