## Describe the changes
This PR fixes the issue of v2 ICICLE not being discovered by Go's
packaging service by adding the required "v2" to the module path:
https://go.dev/doc/modules/release-workflow#breaking
## Describe the changes
This PR fixes an underflow in `get_optimal_c` when bitsize is less than
16 and adds `max(1,NUM_THREADS)` to prevent div by 0 when calculating
NUM_THREADS
## Describe the changes
This PR adds a NTT ReleaseDomain API in Golang and Rust
## Linked Issues
Resolves #
---------
Co-authored-by: Yuval Shekel <yshekel@gmail.com>
## Describe the changes
This PR adds an extern C link to the transpose kernel, now in
vec_ops.cu.
Also Rust binding, and I updated the test check_ntt_batch to use the new
transpose function.
The test passes.
## Linked Issues
Resolves #
---------
Co-authored-by: LeonHibnik <leon@ingonyama.com>
## Describe the changes
This PR adds the capability to slice a DeviceSlice, allowing portions of
data that are already on the device to be reused.
Additionally, this PR removes the need for a HostSlice underlying type
to implement a Size function and uses unsafe.Sizeof instead. This
together with #407 will allow direct usage of gnark-crypto types with
HostSlice without the need for converting to ICICLE types
---------
Co-authored-by: nonam3e <timur@ingonyama.com>
## Describe the changes
Due to Rust's ownership rules, we can't run NTT inplace using the
[`ntt`](https://github.com/ingonyama-zk/icicle/blob/v1.9.1/wrappers/rust/icicle-core/src/ntt/mod.rs#L139)
function. Which is why we saw a need to add a separate function a couple
of times.
Incidentally an issue with radix-2 NTT was found when ran inplace,
`__syncthreads()` was used in reverse order kernel as if it was a global
barrier for all blocks and not block-local one. Thus data race happened
that is fixed by this PR.
…ing other log2 func
## Describe the changes
This PR adds "std" as prefix on log2 function of
icicle/appUtils/msm/msm.cu to explicitly use std::log2 for MSM
calculatation.
## Linked Issues
https://github.com/ingonyama-zk/icicle/issues/458
Resolves #
Co-authored-by: pierre <pierreuu@gmail.com>
This PR adds support for Keccak-256 and Keccak-512. It only adds them in
c++. There is no way of adding rust or golang wrappers rn as it requires
having an `icicle-common` create / mod
## Describe the changes
Speed up CUDA c++ compile time using multi-threaded compilation
(--split-compile flag).
The tests on 8 core machine show ~2x acceleration.
## Linked Issues
Compiling c++ takes long time
## Brief description
This PR adds pre-computation to the MSM, for some theory see
[this](https://youtu.be/KAWlySN7Hm8?si=XeR-htjbnK_ySbUo&t=1734) timecode
of Niall Emmart's talk.
In terms of public APIs, one method is added. It does the
pre-computation on-device leaving resulting data on-device as well. No
extra structures are added, only `precompute_factor` from `MSMConfig` is
now activated.
## Performance
While performance gains are for now often limited by our inflexibility
in choice of `c` (for example, very large MSMs get basically no speedup
from pre-compute because currently `c` cannot be larger than 16),
there's still a number of MSM sizes which get noticeable improvement:
| Pre-computation factor | bn254 size `2^20` MSM, ms. | bn254 size
`2^12` MSM, size `2^10` batch, ms. | bls12-381 size `2^20` MSM, ms. |
bls12-381 size `2^12` MSM, size `2^10` batch, ms. |
| ------------- | ------------- | ------------- | ------------- |
------------- |
| 1 | 14.1 | 82.8 | 25.5 | 136.7 |
| 2 | 11.8 | 76.6 | 20.3 | 123.8 |
| 4 | 10.9 | 73.8 | 18.1 | 117.8 |
| 8 | 10.6 | 73.7 | 17.2 | 116.0 |
Here for example pre-computation factor = 4 means that alongside each
original base point, we pre-compute and pass into the MSM 3 of its
"shifted" versions. Pre-computation factor = 1 means no pre-computation.
GPU used for benchmarks is a 3090Ti.
## TODOs and open questions
- Golang APIs are missing;
- I mentioned that to utilise pre-compute to its full potential we need
arbitrary choice of `c`. One issue with this is that pre-compute will
become dependent on `c`. For now this is not the case as `c` can only be
a power of 2 and powers of 2 can always share the same pre-computation.
So apparently we need to make `c` a parameter of the precompute function
to future-proof it from a breaking change. This is pretty unnatural and
counterintuitive as `c` is typically chosen in runtime after pre-compute
is done but I don't really see another way, pls let me know if you do.
UPD: `c` is added into pre-compute function, for now it's unused and
it's documented how it will change in the future.
Resolves https://github.com/ingonyama-zk/icicle/issues/147
Co-authored with @ChickenLover
---------
Co-authored-by: ChickenLover <romangg81@gmail.com>
Co-authored-by: nonam3e <timur@ingonyama.com>
Co-authored-by: nonam3e <71525212+nonam3e@users.noreply.github.com>
Co-authored-by: LeonHibnik <leon@ingonyama.com>
This PR adds the columns batch feature - enabling batch NTT computation
to be performed directly on the columns of a matrix without having to
transpose it beforehand, as requested in issue #264.
Also some small fixes to the reordering kernels were added and some
unnecessary parameters were removes from functions interfaces.
---------
Co-authored-by: DmytroTym <dmytrotym1@gmail.com>
## Describe the changes
This PR adds multi gpu support in the golang bindings.
Tha main changes are to DeviceSlice which now includes a `deviceId`
attribute specifying which device the underlying data resides on and
checks for correct deviceId and current device when using DeviceSlices
in any operation.
In Go, most concurrency can be done via Goroutines (described as
lightweight threads - in reality, more of a threadpool manager),
however, there is no guarantee that a goroutine stays on a specific host
thread. Therefore, a function `RunOnDevice` was added to the
cuda_runtime package which locks a goroutine into a specific host
thread, sets a current GPU device, runs a provided function, and unlocks
the goroutine from the host thread after the provided function finishes.
While the goroutine is locked to the hsot thread, the Go runtime will
not assign other goroutines to that host thread
## Describe the changes
Add a non-blocking `warmup` function to `CudaStream`
> when you run the benchmark (e.g. the msm example you have) the first
instance is always slow, with a constant overhead of 200~300ms cuda
stream warmup. and I want to get rid of that in my application by
warming it up in parallel while my host do something else.
## Describe the changes
This PR:
- Moves common crate attributes to the workspace Cargo.toml.
- Adds a manual release flow for bumping, tagging, and draft release
This PR is a compilation of small improvements
- Lock bindgen version for `icicle-cuda-runtime`
- Add an error message when trying to build on Mac (or any non
windows/linux machine)
- Add documentation and template files for adding new curve
- Add documentation on _params.cuh contents
- Add the script to bump all the rust crates versions to the same
version
Resolves#313
## Describe the changes
This PR adds the use of setup-go in the CI to ensure that `go` is
installed properly and caches dependencies and build outputs by default
