Add Sha3

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);

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<uint8_t, uint64_t>
   {
   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<uint8_t, uint64_t>(KECCAK_STATE_SIZE, KECCAK_STATE_SIZE, rate, 0) {}
+    Keccak(unsigned int rate, unsigned int padding_const)
+        : Hasher<uint8_t, uint64_t>(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
 

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<uint8_t, uint64_t>(
       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<uint8_t, uint64_t>(
+      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<uint8_t, uint64_t>(
+      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<uint8_t, uint64_t>(
       leaves, digests, height, input_block_len, keccak, keccak, tree_config);
   }

icicle/src/hash/keccak/keccak.cu

@@ -180,8 +180,13 @@ namespace keccak {
   }
 
   template <const int R>
-  __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<KECCAK_256_RATE><<<number_of_gpu_blocks, number_of_threads, 0, ctx.stream>>>(
-        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<KECCAK_512_RATE><<<number_of_gpu_blocks, number_of_threads, 0, ctx.stream>>>(
-        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]");

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(

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<u8> + ?Sized),
+    input_block_size: u32,
+    number_of_blocks: u32,
+    output: &mut (impl HostOrDeviceSlice<u8> + ?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<u8> + ?Sized),
+    input_block_size: u32,
+    number_of_blocks: u32,
+    output: &mut (impl HostOrDeviceSlice<u8> + ?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<u8> + ?Sized),
     digests: &mut (impl HostOrDeviceSlice<u64> + ?Sized),
@@ -123,3 +199,41 @@ pub fn build_keccak512_merkle_tree(
         .wrap()
     }
 }
+
+pub fn build_sha3_256_merkle_tree(
+    leaves: &(impl HostOrDeviceSlice<u8> + ?Sized),
+    digests: &mut (impl HostOrDeviceSlice<u64> + ?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<u8> + ?Sized),
+    digests: &mut (impl HostOrDeviceSlice<u64> + ?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()
+    }
+}

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);