Skip to content

Commit

Permalink
Merge branch 'supranational:main' into dev
Browse files Browse the repository at this point in the history
  • Loading branch information
@winston-h-zhang
winston-h-zhang authored Feb 5, 2024
2 parents f74b7ea + f2887c4 commit efe394a
Show file tree
Hide file tree
Showing 17 changed files with 283 additions and 283 deletions.
55 changes: 49 additions & 6 deletions .github/workflows/ci.yml
View file
Original file line number Diff line number Diff line change
Expand Up @@ -47,7 +47,7 @@ jobs:
sudo apt-get -y install cuda-minimal-build-12-3
[ -d /usr/local/cuda-12.3/bin ]
- name: Test-build poc/ntt-cuda
- name: Test-build poc/ntt-cuda with cuda-12.3
shell: bash
run: |
rustc --version --verbose
Expand All @@ -70,13 +70,9 @@ jobs:
fi
cargo clean -p ntt-cuda
cargo clean -p ntt-cuda --release
rm -rf target/.rustc_info.json
rm -rf target/package
rm -rf target/{debug,release}/incremental
rm -rf target/*/{debug,release}/incremental
)
- name: Test-build poc/msm-cuda
- name: Test-build poc/msm-cuda with cuda-12.3
shell: bash
run: |
rustc --version --verbose
Expand All @@ -99,6 +95,53 @@ jobs:
fi
cargo clean -p msm-cuda
cargo clean -p msm-cuda --release
)
- name: Install cuda-minimal-build-11-8
shell: bash
run: |
sudo apt-get -y install cuda-minimal-build-11-8
[ -d /usr/local/cuda-11.8/bin ]
- name: Test-build poc/ntt-cuda with cuda-11.8
shell: bash
run: |
rustc --version --verbose
export PATH=$PATH:/usr/local/cuda-11.8/bin
( cd poc/ntt-cuda
cargo update
cargo test --no-run --release --features=bls12_381
cargo test --no-run --release --features=gl64
cargo test --no-run --release --features=bb31
cargo clean -p ntt-cuda
cargo clean -p ntt-cuda --release
)
- name: Test-build poc/msm-cuda with cuda-11.8
shell: bash
run: |
rustc --version --verbose
export PATH=$PATH:/usr/local/cuda-11.8/bin
( cd poc/msm-cuda
sed "s/^crit/#crit/" Cargo.toml > Cargo.$$.toml && \
mv Cargo.$$.toml Cargo.toml
cargo update
cargo test --no-run --release --features=bls12_381,quiet
cargo test --no-run --release --features=bn254,quiet
cargo clean -p msm-cuda
cargo clean -p msm-cuda --release
)
- name: Clean up
shell: bash
run: |
( cd poc/ntt-cuda
rm -rf target/.rustc_info.json
rm -rf target/package
rm -rf target/{debug,release}/incremental
rm -rf target/*/{debug,release}/incremental
)
( cd poc/msm-cuda
rm -rf target/.rustc_info.json
rm -rf target/package
rm -rf target/{debug,release}/incremental
Expand Down
14 changes: 8 additions & 6 deletions ec/jacobian_t.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,8 @@
#ifndef __SPPARK_EC_JACOBIAN_T_HPP__
#define __SPPARK_EC_JACOBIAN_T_HPP__

template<class field_t> class jacobian_t {
template<class field_t, const field_t* a4 = nullptr>
class jacobian_t {
field_t X, Y, Z;

inline operator const void*() const { return this; }
Expand Down Expand Up @@ -84,8 +85,7 @@ template<class field_t> class jacobian_t {
* infinity by virtue of Z3 = (U2-U1)*zz = H*zz = 0*zz == 0.
*/
static void dadd(jacobian_t& out, const jacobian_t& p1,
const jacobian_t& p2,
const field_t* a4 = nullptr)
const jacobian_t& p2)
{
jacobian_t p3; /* starts as (U1, S1, zz) from addition side */
struct { field_t H, R, sx; } add, dbl;
Expand Down Expand Up @@ -143,8 +143,8 @@ template<class field_t> class jacobian_t {
vec_select(&p3, &p1, &p3, sizeof(p3), p2inf);
vec_select(out, &p2, &p3, sizeof(p3), p1inf);
}
inline void dadd(const jacobian_t& p2, const field_t* a4 = nullptr)
{ dadd(*this, *this, p2, a4); }
inline void dadd(const jacobian_t& p2)
{ dadd(*this, *this, p2); }

/*
* Addition with affine point that can handle doubling [as well as
Expand Down Expand Up @@ -183,6 +183,8 @@ template<class field_t> class jacobian_t {
dbl.sx = p2.X + p2.X; /* sx = X2+X2 */
dbl.R = p2.X^2; /* X2^2 */
dbl.R += dbl.R + dbl.R; /* R = 3*X2^2 */
if (a4 != nullptr)
dbl.R += *a4;
dbl.H = p2.Y + p2.Y; /* H = 2*Y2 */

p1inf = p1.is_inf();
Expand Down Expand Up @@ -401,7 +403,7 @@ template<class field_t> class jacobian_t {
H = p2.X * Z1Z1; /* U2 = X2*Z1Z1 */
H -= U1; /* H = U2-U1 */

if (H.is_zero() & p3.Z.is_zero()) {
if ((int)H.is_zero() & (int)p3.Z.is_zero()) {
field_t A, B, C; /* double |p1| */

A = p1.X^2; /* A = X1^2 */
Expand Down
58 changes: 53 additions & 5 deletions ec/xyzz_t.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,9 +12,12 @@
# define __device__
# undef __noinline__
# define __noinline__
#else
# pragma nv_diag_suppress 284 // NULL reference is not allowed
#endif

template<class field_t, class field_h = typename field_t::mem_t>
template<class field_t, class field_h = typename field_t::mem_t,
const field_h* a4 = nullptr>
class xyzz_t {
field_t X, Y, ZZZ, ZZ;

Expand Down Expand Up @@ -69,7 +72,7 @@ class xyzz_t {
{ return (bool)(X.is_zero(Y)); }
#else
inline __host__ bool is_inf() const
{ return (bool)(X.is_zero() & Y.is_zero()); }
{ return (bool)((int)X.is_zero() & (int)Y.is_zero()); }
#endif

inline __host__ affine_t& operator=(const xyzz_t& a)
Expand Down Expand Up @@ -188,9 +191,11 @@ class xyzz_t {
#endif

#ifdef __CUDA_ARCH__
inline __device__ bool is_inf() const { return (bool)(ZZZ.is_zero(ZZ)); }
inline __device__ bool is_inf() const
{ return (bool)(ZZZ.is_zero(ZZ)); }
#else
inline __host__ bool is_inf() const { return (bool)(ZZZ.is_zero() & ZZ.is_zero()); }
inline __host__ bool is_inf() const
{ return (bool)((int)ZZZ.is_zero() & (int)ZZ.is_zero()); }
#endif
inline __host__ __device__ void inf() { ZZZ.zero(); ZZ.zero(); }
inline __host__ __device__ void cneg(bool neg) { ZZZ.cneg(neg); }
Expand Down Expand Up @@ -266,6 +271,16 @@ class xyzz_t {
S = p31.X * V; /* S = X1*V */
M = p31.X^2;
M = M + M + M; /* M = 3*X1^2[+a*ZZ1^2] */
if (a4 != nullptr) {
#ifdef __CUDA_ARCH__
U = *a4;
U *= p31.ZZ^2;
#else
U = p31.ZZ^2;
U *= *a4;
#endif
M += U;
}
p31.X = M^2;
p31.X -= S;
p31.X -= S; /* X3 = M^2-2*S */
Expand Down Expand Up @@ -327,7 +342,12 @@ class xyzz_t {
inf = A.is_zero(); /* X1==X2, not add |p1| and |p2| */
dbl = R.is_zero() & inf;
if (dbl) { /* X1==X2 && Y1==Y2, double |p2| */
A = p2.Y<<1; /* U = 2*Y1 */
if (a4 != nullptr) {
A = p2.ZZ;
pc = 16;
} else {
A = p2.Y<<1; /* U = 2*Y1 */
}
inf = false; /* don't set |p3| to infinity */
}
B = A;
Expand Down Expand Up @@ -387,6 +407,8 @@ class xyzz_t {
#define S PP
case 14:
A = A + A + A; /* M = 3*X1^2[+a*ZZ1^2] */
if (a4 != nullptr)
A += U;
B = A;
break;
case 15:
Expand All @@ -398,6 +420,19 @@ class xyzz_t {
done = true;
break;
#undef S
/*** account for a4 != nullptr when doubling ***/
case 17: /* ZZ1^2 */
if (a4 != nullptr)
B = *a4;
break;
case 18: /* ZZ1^2*a4 */
if (a4 != nullptr) {
U = A;
A = p2.Y<<1; /* U = 2*Y1 */
B = A;
}
pc = 3;
break;
}
} while (!done);
p31.Y = A - p31.Y; /* Y3 = R*(Q-X3)-S1*PPP */
Expand Down Expand Up @@ -468,6 +503,13 @@ class xyzz_t {
S = p2.X * p31.ZZ; /* S = X1*V */
M = p2.X^2;
M = M + M + M; /* M = 3*X1^2[+a] */
if (a4 != nullptr) {
#ifdef __CUDA_ARCH__
M += (U = *a4);
#else
M += *a4;
#endif
}
p31.X = M^2;
p31.X -= S;
p31.X -= S; /* X3 = M^2-2*S */
Expand Down Expand Up @@ -579,6 +621,8 @@ class xyzz_t {
#define S R
case 10: /* X1^2 */
A += A<<1; /* M = 3*X1^2[+a] */
if (a4 != nullptr)
A += (B = *a4);
B = A;
break;
case 11: /* M^2 */
Expand Down Expand Up @@ -617,4 +661,8 @@ class xyzz_t {
}
#endif
};

#ifdef __CUDACC__
# pragma nv_diag_default 284
#endif
#endif
2 changes: 1 addition & 1 deletion ff/bls12-377-fp2.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -426,7 +426,7 @@ class fp2_t {

friend inline fp2_t czero(const fp2_t& a, int set_z)
{ fp2_t ret;
const vec384x zero = { 0 };
const vec384x zero = {{0}};
vec_select(ret.val, zero, a.val, sizeof(ret), set_z);
return ret;
}
Expand Down
2 changes: 1 addition & 1 deletion ff/bls12-381-fp2.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -337,7 +337,7 @@ class fp2_t {

friend inline fp2_t czero(const fp2_t& a, int set_z)
{ fp2_t ret;
const vec384x zero = { 0 };
const vec384x zero = {{0}};
vec_select(ret.val, zero, a.val, sizeof(ret), set_z);
return ret;
}
Expand Down
100 changes: 0 additions & 100 deletions ntt/gen_twiddles.cu
View file

This file was deleted.

Loading

0 comments on commit efe394a

Please sign in to comment.