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

Add Sha3 (#560)

Browse Source 
## Describe the changes

This PR...

## Linked Issues

Resolves #
This commit is contained in:
ChickenLover
2024-07-28 15:31:28 +07:00
committed by GitHub
parent 8411ed1451
commit 0cb0b49be9
7 changed files with 213 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
docs/docs/icicle/primitives/keccak.md
View File

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

34
icicle/include/hash/keccak/keccak.cuh
View File

@@ -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

50
icicle/src/hash/keccak/extern.cu
View File

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

15
icicle/src/hash/keccak/keccak.cu
View File

@@ -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]");

10
icicle/src/merkle-tree/merkle.cu
View File

@@ -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(

114
wrappers/rust/icicle-hash/src/keccak/mod.rs
View File

@@ -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()
}
}

2
wrappers/rust/icicle-hash/src/keccak/tests.rs
View File

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