diff --git a/docs/docs/icicle/primitives/keccak.md b/docs/docs/icicle/primitives/keccak.md index 37ea9904d..e81b8f86c 100644 --- a/docs/docs/icicle/primitives/keccak.md +++ b/docs/docs/icicle/primitives/keccak.md @@ -12,6 +12,10 @@ At its core, Keccak consists of a permutation function operating on a state arra - **Chi:** This step applies a nonlinear mixing operation to each lane of the state array. - **Iota:** This step introduces a round constant to the state array. +## Keccak vs Sha3 + +There exists a [confusion](https://www.cybertest.com/blog/keccak-vs-sha3) between what is called `Keccak` and `Sha3`. In ICICLE we support both. `Keccak256` relates to the old hash function used in Ethereum, and `Sha3-256` relates to the modern hash function. + ## Using Keccak ICICLE Keccak supports batch hashing, which can be utilized for constructing a merkle tree or running multiple hashes in parallel. @@ -35,7 +39,7 @@ let input_block_len = 136; let number_of_hashes = 1024; let preimages = vec![1u8; number_of_hashes * input_block_len]; -let mut digests = vec![0u8; number_of_hashes * 64]; +let mut digests = vec![0u8; number_of_hashes * 32]; let preimages_slice = HostSlice::from_slice(&preimages); let digests_slice = HostSlice::from_mut_slice(&mut digests); diff --git a/icicle/include/hash/keccak/keccak.cuh b/icicle/include/hash/keccak/keccak.cuh index d95c22a40..01bed87ea 100644 --- a/icicle/include/hash/keccak/keccak.cuh +++ b/icicle/include/hash/keccak/keccak.cuh @@ -22,9 +22,14 @@ namespace keccak { // Number of state elements in u64 const int KECCAK_STATE_SIZE = 25; + const int KECCAK_PADDING_CONST = 1; + const int SHA3_PADDING_CONST = 6; + class Keccak : public Hasher { public: + const int PADDING_CONST; + cudaError_t run_hash_many_kernel( const uint8_t* input, uint64_t* output, @@ -33,7 +38,34 @@ namespace keccak { unsigned int output_len, const device_context::DeviceContext& ctx) const override; - Keccak(unsigned int rate) : Hasher(KECCAK_STATE_SIZE, KECCAK_STATE_SIZE, rate, 0) {} + Keccak(unsigned int rate, unsigned int padding_const) + : Hasher(KECCAK_STATE_SIZE, KECCAK_STATE_SIZE, rate, 0), PADDING_CONST(padding_const) + { + } + }; + + class Keccak256 : public Keccak + { + public: + Keccak256() : Keccak(KECCAK_256_RATE, KECCAK_PADDING_CONST) {} + }; + + class Keccak512 : public Keccak + { + public: + Keccak512() : Keccak(KECCAK_512_RATE, KECCAK_PADDING_CONST) {} + }; + + class Sha3_256 : public Keccak + { + public: + Sha3_256() : Keccak(KECCAK_256_RATE, SHA3_PADDING_CONST) {} + }; + + class Sha3_512 : public Keccak + { + public: + Sha3_512() : Keccak(KECCAK_512_RATE, SHA3_PADDING_CONST) {} }; } // namespace keccak diff --git a/icicle/src/hash/keccak/extern.cu b/icicle/src/hash/keccak/extern.cu index b1e6d6aaa..519a0e14f 100644 --- a/icicle/src/hash/keccak/extern.cu +++ b/icicle/src/hash/keccak/extern.cu @@ -11,15 +11,29 @@ namespace keccak { extern "C" cudaError_t keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, HashConfig& config) { - return Keccak(KECCAK_256_RATE) - .hash_many(input, (uint64_t*)output, number_of_blocks, input_block_size, KECCAK_256_DIGEST, config); + return Keccak256().hash_many( + input, (uint64_t*)output, number_of_blocks, input_block_size, KECCAK_256_DIGEST, config); } extern "C" cudaError_t keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, HashConfig& config) { - return Keccak(KECCAK_512_RATE) - .hash_many(input, (uint64_t*)output, number_of_blocks, input_block_size, KECCAK_512_DIGEST, config); + return Keccak512().hash_many( + input, (uint64_t*)output, number_of_blocks, input_block_size, KECCAK_512_DIGEST, config); + } + + extern "C" cudaError_t + sha3_256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, HashConfig& config) + { + return Sha3_256().hash_many( + input, (uint64_t*)output, number_of_blocks, input_block_size, KECCAK_256_DIGEST, config); + } + + extern "C" cudaError_t + sha3_512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, HashConfig& config) + { + return Sha3_512().hash_many( + input, (uint64_t*)output, number_of_blocks, input_block_size, KECCAK_512_DIGEST, config); } extern "C" cudaError_t build_keccak256_merkle_tree_cuda( @@ -29,7 +43,7 @@ namespace keccak { unsigned int input_block_len, const merkle_tree::TreeBuilderConfig& tree_config) { - Keccak keccak(KECCAK_256_RATE); + Keccak256 keccak; return merkle_tree::build_merkle_tree( leaves, digests, height, input_block_len, keccak, keccak, tree_config); } @@ -41,7 +55,31 @@ namespace keccak { unsigned int input_block_len, const merkle_tree::TreeBuilderConfig& tree_config) { - Keccak keccak(KECCAK_512_RATE); + Keccak512 keccak; + return merkle_tree::build_merkle_tree( + leaves, digests, height, input_block_len, keccak, keccak, tree_config); + } + + extern "C" cudaError_t build_sha3_256_merkle_tree_cuda( + const uint8_t* leaves, + uint64_t* digests, + unsigned int height, + unsigned int input_block_len, + const merkle_tree::TreeBuilderConfig& tree_config) + { + Sha3_256 keccak; + return merkle_tree::build_merkle_tree( + leaves, digests, height, input_block_len, keccak, keccak, tree_config); + } + + extern "C" cudaError_t build_sha3_512_merkle_tree_cuda( + const uint8_t* leaves, + uint64_t* digests, + unsigned int height, + unsigned int input_block_len, + const merkle_tree::TreeBuilderConfig& tree_config) + { + Sha3_512 keccak; return merkle_tree::build_merkle_tree( leaves, digests, height, input_block_len, keccak, keccak, tree_config); } diff --git a/icicle/src/hash/keccak/keccak.cu b/icicle/src/hash/keccak/keccak.cu index e805bcf63..542da24ae 100644 --- a/icicle/src/hash/keccak/keccak.cu +++ b/icicle/src/hash/keccak/keccak.cu @@ -180,8 +180,13 @@ namespace keccak { } template - __global__ void - keccak_hash_blocks(const uint8_t* input, int input_block_size, int output_len, int number_of_blocks, uint64_t* output) + __global__ void keccak_hash_blocks( + const uint8_t* input, + int input_block_size, + int output_len, + int number_of_blocks, + uint64_t* output, + int padding_const) { int sid = (blockIdx.x * blockDim.x) + threadIdx.x; if (sid >= number_of_blocks) { return; } @@ -209,7 +214,7 @@ namespace keccak { } // pad 10*1 - last_block[input_len] = 1; + last_block[input_len] = padding_const; for (int i = 0; i < R - input_len - 1; i++) { last_block[input_len + i + 1] = 0; } @@ -240,11 +245,11 @@ namespace keccak { switch (rate) { case KECCAK_256_RATE: keccak_hash_blocks<<>>( - input, input_len, output_len, number_of_states, output); + input, input_len, output_len, number_of_states, output, PADDING_CONST); break; case KECCAK_512_RATE: keccak_hash_blocks<<>>( - input, input_len, output_len, number_of_states, output); + input, input_len, output_len, number_of_states, output, PADDING_CONST); break; default: THROW_ICICLE_ERR(IcicleError_t::InvalidArgument, "KeccakHash: #rate must be one of [136, 72]"); diff --git a/icicle/src/merkle-tree/merkle.cu b/icicle/src/merkle-tree/merkle.cu index 53f3b8f6f..2fe171634 100644 --- a/icicle/src/merkle-tree/merkle.cu +++ b/icicle/src/merkle-tree/merkle.cu @@ -129,8 +129,9 @@ namespace merkle_tree { while (number_of_states > 0) { CHK_IF_RETURN(compression.run_hash_many_kernel( - (L*)prev_layer, next_layer, number_of_states, tree_config.digest_elements * tree_config.arity, - tree_config.digest_elements, hash_config.ctx)); + (L*)prev_layer, next_layer, number_of_states, + tree_config.digest_elements * tree_config.arity * (sizeof(D) / sizeof(L)), tree_config.digest_elements, + hash_config.ctx)); if (!keep_rows || subtree_height < keep_rows) { D* digests_with_offset = @@ -298,8 +299,9 @@ namespace merkle_tree { size_t segment_offset = start_segment_offset; while (number_of_states > 0) { CHK_IF_RETURN(compression.run_hash_many_kernel( - (L*)prev_layer, next_layer, number_of_states, tree_config.digest_elements * tree_config.arity, - tree_config.digest_elements, tree_config.ctx)); + (L*)prev_layer, next_layer, number_of_states, + tree_config.digest_elements * tree_config.arity * (sizeof(D) / sizeof(L)), tree_config.digest_elements, + tree_config.ctx)); if (!tree_config.keep_rows || cap_height < tree_config.keep_rows + (int)caps_mode) { D* digests_with_offset = digests + segment_offset; CHK_IF_RETURN(cudaMemcpyAsync( diff --git a/wrappers/rust/icicle-hash/src/keccak/mod.rs b/wrappers/rust/icicle-hash/src/keccak/mod.rs index 88fc5a2c7..df2cbfb73 100644 --- a/wrappers/rust/icicle-hash/src/keccak/mod.rs +++ b/wrappers/rust/icicle-hash/src/keccak/mod.rs @@ -25,6 +25,22 @@ extern "C" { config: &HashConfig, ) -> CudaError; + pub(crate) fn sha3_256_cuda( + input: *const u8, + input_block_size: u32, + number_of_blocks: u32, + output: *mut u8, + config: &HashConfig, + ) -> CudaError; + + pub(crate) fn sha3_512_cuda( + input: *const u8, + input_block_size: u32, + number_of_blocks: u32, + output: *mut u8, + config: &HashConfig, + ) -> CudaError; + pub(crate) fn build_keccak256_merkle_tree_cuda( leaves: *const u8, digests: *mut u64, @@ -40,6 +56,22 @@ extern "C" { input_block_len: u32, config: &TreeBuilderConfig, ) -> CudaError; + + pub(crate) fn build_sha3_256_merkle_tree_cuda( + leaves: *const u8, + digests: *mut u64, + height: u32, + input_block_len: u32, + config: &TreeBuilderConfig, + ) -> CudaError; + + pub(crate) fn build_sha3_512_merkle_tree_cuda( + leaves: *const u8, + digests: *mut u64, + height: u32, + input_block_len: u32, + config: &TreeBuilderConfig, + ) -> CudaError; } pub fn keccak256( @@ -86,6 +118,50 @@ pub fn keccak512( } } +pub fn sha3_256( + input: &(impl HostOrDeviceSlice + ?Sized), + input_block_size: u32, + number_of_blocks: u32, + output: &mut (impl HostOrDeviceSlice + ?Sized), + config: &HashConfig, +) -> IcicleResult<()> { + let mut local_cfg = config.clone(); + local_cfg.are_inputs_on_device = input.is_on_device(); + local_cfg.are_outputs_on_device = output.is_on_device(); + unsafe { + sha3_256_cuda( + input.as_ptr(), + input_block_size, + number_of_blocks, + output.as_mut_ptr(), + &local_cfg, + ) + .wrap() + } +} + +pub fn sha3_512( + input: &(impl HostOrDeviceSlice + ?Sized), + input_block_size: u32, + number_of_blocks: u32, + output: &mut (impl HostOrDeviceSlice + ?Sized), + config: &HashConfig, +) -> IcicleResult<()> { + let mut local_cfg = config.clone(); + local_cfg.are_inputs_on_device = input.is_on_device(); + local_cfg.are_outputs_on_device = output.is_on_device(); + unsafe { + sha3_512_cuda( + input.as_ptr(), + input_block_size, + number_of_blocks, + output.as_mut_ptr(), + &local_cfg, + ) + .wrap() + } +} + pub fn build_keccak256_merkle_tree( leaves: &(impl HostOrDeviceSlice + ?Sized), digests: &mut (impl HostOrDeviceSlice + ?Sized), @@ -123,3 +199,41 @@ pub fn build_keccak512_merkle_tree( .wrap() } } + +pub fn build_sha3_256_merkle_tree( + leaves: &(impl HostOrDeviceSlice + ?Sized), + digests: &mut (impl HostOrDeviceSlice + ?Sized), + height: usize, + input_block_len: usize, + config: &TreeBuilderConfig, +) -> IcicleResult<()> { + unsafe { + build_sha3_256_merkle_tree_cuda( + leaves.as_ptr(), + digests.as_mut_ptr(), + height as u32, + input_block_len as u32, + config, + ) + .wrap() + } +} + +pub fn build_sha3_512_merkle_tree( + leaves: &(impl HostOrDeviceSlice + ?Sized), + digests: &mut (impl HostOrDeviceSlice + ?Sized), + height: usize, + input_block_len: usize, + config: &TreeBuilderConfig, +) -> IcicleResult<()> { + unsafe { + build_sha3_512_merkle_tree_cuda( + leaves.as_ptr(), + digests.as_mut_ptr(), + height as u32, + input_block_len as u32, + config, + ) + .wrap() + } +} diff --git a/wrappers/rust/icicle-hash/src/keccak/tests.rs b/wrappers/rust/icicle-hash/src/keccak/tests.rs index 5675e9158..0aa8cbc0e 100644 --- a/wrappers/rust/icicle-hash/src/keccak/tests.rs +++ b/wrappers/rust/icicle-hash/src/keccak/tests.rs @@ -15,7 +15,7 @@ pub(crate) mod tests { let number_of_hashes = 1024; let preimages = vec![1u8; number_of_hashes * input_block_len]; - let mut digests = vec![0u8; number_of_hashes * 64]; + let mut digests = vec![0u8; number_of_hashes * 32]; let preimages_slice = HostSlice::from_slice(&preimages); let digests_slice = HostSlice::from_mut_slice(&mut digests);