From e2ad621f97d6899cc355632b71b9969f75367ed7 Mon Sep 17 00:00:00 2001 From: nonam3e <71525212+nonam3e@users.noreply.github.com> Date: Wed, 1 May 2024 18:08:33 +0700 Subject: [PATCH] Nonam3e/golang/keccak (#496) ## Describe the changes This PR adds keccak bindings + passes cfg as reference in keccak cuda functions --- .github/workflows/golang.yml | 31 ++++++- icicle/include/api/hash.h | 4 +- icicle/include/hash/keccak/keccak.cuh | 2 +- icicle/src/hash/keccak/keccak.cu | 8 +- icicle/src/hash/keccak/test.cu | 4 +- wrappers/golang/build.sh | 18 ++++ wrappers/golang/hash/keccak/hasher.go | 88 +++++++++++++++++++ wrappers/golang/hash/keccak/include/keccak.h | 20 +++++ wrappers/golang/hash/keccak/main.go | 4 + .../golang/hash/keccak/tests/hasher_test.go | 66 ++++++++++++++ wrappers/rust/icicle-hash/src/keccak/mod.rs | 8 +- 11 files changed, 239 insertions(+), 14 deletions(-) create mode 100644 wrappers/golang/hash/keccak/hasher.go create mode 100644 wrappers/golang/hash/keccak/include/keccak.h create mode 100644 wrappers/golang/hash/keccak/main.go create mode 100644 wrappers/golang/hash/keccak/tests/hasher_test.go diff --git a/.github/workflows/golang.yml b/.github/workflows/golang.yml index 1a89d12d..abc02612 100644 --- a/.github/workflows/golang.yml +++ b/.github/workflows/golang.yml @@ -99,11 +99,40 @@ jobs: path: | icicle/build/lib/libingo_field_${{ matrix.field.name }}.a retention-days: 1 + + build-hashes-linux: + name: Build hashes on Linux + runs-on: [self-hosted, Linux, X64, icicle] + needs: [check-changed-files, check-format] + strategy: + matrix: + hash: + - name: keccak + build_args: + steps: + - name: Checkout Repo + uses: actions/checkout@v4 + - name: Setup go + uses: actions/setup-go@v5 + with: + go-version: '1.20.0' + - name: Build + working-directory: ./wrappers/golang + if: needs.check-changed-files.outputs.golang == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true' + run: ./build.sh -hash=${{ matrix.hash.name }} ${{ matrix.hash.build_args }} # builds a single hash algorithm + - name: Upload ICICLE lib artifacts + uses: actions/upload-artifact@v4 + if: needs.check-changed-files.outputs.golang == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true' + with: + name: icicle-builds-${{ matrix.hash.name }}-${{ github.workflow }}-${{ github.sha }} + path: | + icicle/build/lib/libingo_hash.a + retention-days: 1 test-linux: name: Test on Linux runs-on: [self-hosted, Linux, X64, icicle] - needs: [check-changed-files, build-curves-linux, build-fields-linux] + needs: [check-changed-files, build-curves-linux, build-fields-linux, build-hashes-linux] steps: - name: Checkout Repo uses: actions/checkout@v4 diff --git a/icicle/include/api/hash.h b/icicle/include/api/hash.h index 70649811..a85e1b6c 100644 --- a/icicle/include/api/hash.h +++ b/icicle/include/api/hash.h @@ -8,9 +8,9 @@ #include "hash/keccak/keccak.cuh" extern "C" cudaError_t - keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config); + keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, keccak::KeccakConfig& config); extern "C" cudaError_t - keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config); + keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, keccak::KeccakConfig& config); #endif \ No newline at end of file diff --git a/icicle/include/hash/keccak/keccak.cuh b/icicle/include/hash/keccak/keccak.cuh index 24c856d6..251ace3b 100644 --- a/icicle/include/hash/keccak/keccak.cuh +++ b/icicle/include/hash/keccak/keccak.cuh @@ -50,7 +50,7 @@ namespace keccak { */ template cudaError_t - keccak_hash(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config); + keccak_hash(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig& config); } // namespace keccak #endif \ No newline at end of file diff --git a/icicle/src/hash/keccak/keccak.cu b/icicle/src/hash/keccak/keccak.cu index 572a96cf..8655eed5 100644 --- a/icicle/src/hash/keccak/keccak.cu +++ b/icicle/src/hash/keccak/keccak.cu @@ -224,7 +224,7 @@ namespace keccak { template cudaError_t - keccak_hash(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config) + keccak_hash(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig& config) { CHK_INIT_IF_RETURN(); cudaStream_t& stream = config.ctx.stream; @@ -245,7 +245,7 @@ namespace keccak { CHK_IF_RETURN(cudaMallocAsync(&output_device, number_of_blocks * (D / 8), stream)); } - int number_of_threads = 1024; + int number_of_threads = 512; int number_of_gpu_blocks = (number_of_blocks - 1) / number_of_threads + 1; keccak_hash_blocks<<>>( input_device, input_block_size, number_of_blocks, output_device); @@ -262,13 +262,13 @@ namespace keccak { } extern "C" cudaError_t - keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config) + keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig& config) { return keccak_hash<512, 256>(input, input_block_size, number_of_blocks, output, config); } extern "C" cudaError_t - keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config) + keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig& config) { return keccak_hash<1024, 512>(input, input_block_size, number_of_blocks, output, config); } diff --git a/icicle/src/hash/keccak/test.cu b/icicle/src/hash/keccak/test.cu index 85e09f35..8149dc03 100644 --- a/icicle/src/hash/keccak/test.cu +++ b/icicle/src/hash/keccak/test.cu @@ -1,4 +1,4 @@ -#include "utils/device_context.cuh" +#include "gpu-utils/device_context.cuh" #include "keccak.cu" // #define DEBUG @@ -51,7 +51,7 @@ int main(int argc, char* argv[]) START_TIMER(keccak_timer); KeccakConfig config = default_keccak_config(); - keccak256(in_ptr, input_block_size, number_of_blocks, out_ptr, config); + keccak256_cuda(in_ptr, input_block_size, number_of_blocks, out_ptr, config); END_TIMER(keccak_timer, "Keccak") for (int i = 0; i < number_of_blocks; i++) { diff --git a/wrappers/golang/build.sh b/wrappers/golang/build.sh index b858ac72..caef755e 100755 --- a/wrappers/golang/build.sh +++ b/wrappers/golang/build.sh @@ -7,9 +7,11 @@ DEVMODE=OFF EXT_FIELD=OFF BUILD_CURVES=( ) BUILD_FIELDS=( ) +BUILD_HASHES=( ) SUPPORTED_CURVES=("bn254" "bls12_377" "bls12_381" "bw6_761", "grumpkin") SUPPORTED_FIELDS=("babybear") +SUPPORTED_HASHES=("keccak") if [[ $1 == "-help" ]]; then echo "Build script for building ICICLE cpp libraries" @@ -67,6 +69,15 @@ do -field-ext) EXT_FIELD=ON ;; + -hash*) + hash=$(echo "$arg_lower" | cut -d'=' -f2) + if [[ $hash == "all" ]] + then + BUILD_HASHES=("${SUPPORTED_HASHES[@]}") + else + BUILD_HASHES=( $hash ) + fi + ;; -devmode) DEVMODE=ON ;; @@ -105,3 +116,10 @@ do cmake --build build -j8 && rm build_config.txt done +for HASH in "${BUILD_HASHES[@]}" +do + echo "HASH=${HASH_DEFINED}" > build_config.txt + echo "DEVMODE=${DEVMODE}" >> build_config.txt + cmake -DCMAKE_CUDA_COMPILER=$CUDA_COMPILER_PATH -DBUILD_HASH=$HASH -DDEVMODE=$DEVMODE -DCMAKE_BUILD_TYPE=Release -S . -B build + cmake --build build -j8 && rm build_config.txt +done diff --git a/wrappers/golang/hash/keccak/hasher.go b/wrappers/golang/hash/keccak/hasher.go new file mode 100644 index 00000000..304342a8 --- /dev/null +++ b/wrappers/golang/hash/keccak/hasher.go @@ -0,0 +1,88 @@ +package keccak + +// #cgo CFLAGS: -I./include/ +// #include "keccak.h" +import "C" + +import ( + "fmt" + "unsafe" + + core "github.com/ingonyama-zk/icicle/v2/wrappers/golang/core" + cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime" +) + +type HashSize int + +const ( + Hash256 HashSize = 256 + Hash512 HashSize = 512 +) + +type KeccakConfig struct { + Ctx cr.DeviceContext + areInputsOnDevice bool + areOutputsOnDevice bool + IsAsync bool +} + +func GetDefaultKeccakConfig() KeccakConfig { + ctx, _ := cr.GetDefaultDeviceContext() + return KeccakConfig{ + ctx, + false, + false, + false, + } +} + +func keccakCheck(input core.HostOrDeviceSlice, output core.HostOrDeviceSlice, cfg *KeccakConfig, hashSize HashSize, numberOfBlocks int32) (unsafe.Pointer, unsafe.Pointer, unsafe.Pointer) { + cfg.areInputsOnDevice = input.IsOnDevice() + cfg.areOutputsOnDevice = output.IsOnDevice() + + if input.IsOnDevice() { + input.(core.DeviceSlice).CheckDevice() + } + + if output.IsOnDevice() { + output.(core.DeviceSlice).CheckDevice() + } + + if output.Cap() < int(hashSize)/8*int(numberOfBlocks) { + errorString := fmt.Sprintf( + "Output capacity %d isn't enough for hashSize %d and numberOfBlocks %d", + output.Cap(), + hashSize, + numberOfBlocks, + ) + panic(errorString) + } + + return input.AsUnsafePointer(), output.AsUnsafePointer(), unsafe.Pointer(cfg) +} + +func keccak(input core.HostOrDeviceSlice, inputBlockSize, numberOfBlocks int32, output core.HostOrDeviceSlice, config *KeccakConfig, hashSize HashSize) (ret core.IcicleError) { + inputPointer, outputPointer, cfgPointer := keccakCheck(input, output, config, hashSize, numberOfBlocks) + cInput := (*C.uint8_t)(inputPointer) + cOutput := (*C.uint8_t)(outputPointer) + cInputBlockSize := (C.int)(inputBlockSize) + cNumberOfBlocks := (C.int)(numberOfBlocks) + cConfig := (*C.KeccakConfig)(cfgPointer) + + switch hashSize { + case Hash256: + ret = core.FromCudaError((cr.CudaError)(C.keccak256_cuda(cInput, cInputBlockSize, cNumberOfBlocks, cOutput, cConfig))) + case Hash512: + ret = core.FromCudaError((cr.CudaError)(C.keccak512_cuda(cInput, cInputBlockSize, cNumberOfBlocks, cOutput, cConfig))) + } + + return ret +} + +func Keccak256(input core.HostOrDeviceSlice, inputBlockSize, numberOfBlocks int32, output core.HostOrDeviceSlice, config *KeccakConfig) core.IcicleError { + return keccak(input, inputBlockSize, numberOfBlocks, output, config, Hash256) +} + +func Keccak512(input core.HostOrDeviceSlice, inputBlockSize, numberOfBlocks int32, output core.HostOrDeviceSlice, config *KeccakConfig) core.IcicleError { + return keccak(input, inputBlockSize, numberOfBlocks, output, config, Hash512) +} diff --git a/wrappers/golang/hash/keccak/include/keccak.h b/wrappers/golang/hash/keccak/include/keccak.h new file mode 100644 index 00000000..e2202e76 --- /dev/null +++ b/wrappers/golang/hash/keccak/include/keccak.h @@ -0,0 +1,20 @@ +#include +#include + +#ifndef _KECCAK_HASH_H +#define _KECCAK_HASH_H + +#ifdef __cplusplus +extern "C" { +#endif + +typedef struct KeccakConfig KeccakConfig; + +cudaError_t keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig* config); +cudaError_t keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig* config); + +#ifdef __cplusplus +} +#endif + +#endif diff --git a/wrappers/golang/hash/keccak/main.go b/wrappers/golang/hash/keccak/main.go new file mode 100644 index 00000000..2f342c8e --- /dev/null +++ b/wrappers/golang/hash/keccak/main.go @@ -0,0 +1,4 @@ +package keccak + +// #cgo LDFLAGS: -L${SRCDIR}/../../../../icicle/build/lib -lingo_hash -lstdc++ -lm +import "C" diff --git a/wrappers/golang/hash/keccak/tests/hasher_test.go b/wrappers/golang/hash/keccak/tests/hasher_test.go new file mode 100644 index 00000000..040715be --- /dev/null +++ b/wrappers/golang/hash/keccak/tests/hasher_test.go @@ -0,0 +1,66 @@ +package tests + +import ( + "encoding/hex" + "testing" + + "github.com/ingonyama-zk/icicle/v2/wrappers/golang/core" + cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime" + "github.com/ingonyama-zk/icicle/v2/wrappers/golang/hash/keccak" + + "github.com/stretchr/testify/assert" +) + +func createHostSliceFromHexString(hexString string) core.HostSlice[uint8] { + byteArray, err := hex.DecodeString(hexString) + if err != nil { + panic("Not a hex string") + } + return core.HostSliceFromElements([]uint8(byteArray)) +} + +func TestSimpleHash256(t *testing.T) { + input := createHostSliceFromHexString("1725b6679911bfe75ad7e248fbeec4a01034feace99aca43cd95d338a37db8d04b4aa5d83c8f8f5bdb8f7f98cec9a658f7f8061a6de07dcfd61db169cc7e666e1971adb4c7e97c43361c9a9eed8bb7b5c33cfe576a383a7440803996fd28148b") + outHost := make(core.HostSlice[uint8], 32) + + cfg := keccak.GetDefaultKeccakConfig() + e := keccak.Keccak256(input, int32(input.Len()), 1, outHost, &cfg) + assert.Equal(t, e.CudaErrorCode, cr.CudaSuccess, "Hashing failed") + t.Log(outHost) + assert.Equal(t, outHost, createHostSliceFromHexString("10fd4a3df6046e32f282cad3ac78e1566304339e7a6696826af023a55ab42048")) +} + +func TestBatchHash256(t *testing.T) { + input := createHostSliceFromHexString("1725b6679911bfe75ad7e248fbeec4a01034feace99aca43cd95d338a37db8d04b4aa5d83c8f8f5bdb8f7f98cec9a658f7f8061a6de07dcfd61db169cc7e666e1971adb4c7e97c43361c9a9eed8bb7b5c33cfe576a383a7440803996fd28148b") + outHost := make(core.HostSlice[uint8], 32*2) + + cfg := keccak.GetDefaultKeccakConfig() + e := keccak.Keccak256(input, int32(input.Len()/2), 2, outHost, &cfg) + assert.Equal(t, e.CudaErrorCode, cr.CudaSuccess, "Hashing failed") + t.Log(outHost) + assert.Equal(t, outHost[:32], createHostSliceFromHexString("7983fbc4cb4539cc90731205c44f74ca74e0a49ad1032a7a1429b1e443e66f45")) + assert.Equal(t, outHost[32:64], createHostSliceFromHexString("2952c2491c75338d28943231a492e9ab684a6820e4af1d74c8c1976759f7bf4b")) +} + +func TestSimpleHash512(t *testing.T) { + input := createHostSliceFromHexString("1725b6679911bfe75ad7e248fbeec4a01034feace99aca43cd95d338a37db8d04b4aa5d83c8f8f5bdb8f7f98cec9a658f7f8061a6de07dcfd61db169cc7e666e1971adb4c7e97c43361c9a9eed8bb7b5c33cfe576a383a7440803996fd28148b") + outHost := make(core.HostSlice[uint8], 64) + + cfg := keccak.GetDefaultKeccakConfig() + e := keccak.Keccak512(input, int32(input.Len()), 1, outHost, &cfg) + assert.Equal(t, e.CudaErrorCode, cr.CudaSuccess, "Hashing failed") + t.Log(outHost) + assert.Equal(t, outHost, createHostSliceFromHexString("1da4e0264dc755bc0b3a3318d2496e11c72322104693b68dbddfa66aa6e8b95526e95a7684a55ea831202f475f3d6a322ed86360d7e0e80f4a129f15d59dd403")) +} + +func TestBatchHash512(t *testing.T) { + input := createHostSliceFromHexString("1725b6679911bfe75ad7e248fbeec4a01034feace99aca43cd95d338a37db8d04b4aa5d83c8f8f5bdb8f7f98cec9a658f7f8061a6de07dcfd61db169cc7e666e1971adb4c7e97c43361c9a9eed8bb7b5c33cfe576a383a7440803996fd28148b") + outHost := make(core.HostSlice[uint8], 64*2) + + cfg := keccak.GetDefaultKeccakConfig() + e := keccak.Keccak512(input, int32(input.Len()/2), 2, outHost, &cfg) + assert.Equal(t, e.CudaErrorCode, cr.CudaSuccess, "Hashing failed") + t.Log(outHost) + assert.Equal(t, outHost[:64], createHostSliceFromHexString("709974f0dc1df1461fcbc2275e968fcb510c947d38837d577d661b6b40249c6b348e33092e4795faad7d2829403bd70fe860207f40a84a23e03c4610ca7927a9")) + assert.Equal(t, outHost[64:128], createHostSliceFromHexString("b8e46caa6cf7fbe6858deb28d4d9e58b768333b1260f5386656c0ae0d0850262bf6aa00293ef0979c37903fb5d2b784a02a4a227725a2b091df182abda03231d")) +} diff --git a/wrappers/rust/icicle-hash/src/keccak/mod.rs b/wrappers/rust/icicle-hash/src/keccak/mod.rs index 2fbd52ce..7d5ea2b6 100644 --- a/wrappers/rust/icicle-hash/src/keccak/mod.rs +++ b/wrappers/rust/icicle-hash/src/keccak/mod.rs @@ -51,7 +51,7 @@ extern "C" { input_block_size: i32, number_of_blocks: i32, output: *mut u8, - config: KeccakConfig, + config: &KeccakConfig, ) -> CudaError; pub(crate) fn keccak512_cuda( @@ -59,7 +59,7 @@ extern "C" { input_block_size: i32, number_of_blocks: i32, output: *mut u8, - config: KeccakConfig, + config: &KeccakConfig, ) -> CudaError; } @@ -68,7 +68,7 @@ pub fn keccak256( input_block_size: i32, number_of_blocks: i32, output: &mut (impl HostOrDeviceSlice + ?Sized), - config: KeccakConfig, + config: &KeccakConfig, ) -> IcicleResult<()> { unsafe { keccak256_cuda( @@ -87,7 +87,7 @@ pub fn keccak512( input_block_size: i32, number_of_blocks: i32, output: &mut (impl HostOrDeviceSlice + ?Sized), - config: KeccakConfig, + config: &KeccakConfig, ) -> IcicleResult<()> { unsafe { keccak512_cuda(