Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

bw6761 golang wrapper #2

Merged
merged 26 commits into from
Oct 18, 2023
Merged

Conversation

ImmanuelSegol
Copy link

Describe the changes

This PR...

Linked Issues

Resolves #

ImmanuelSegol and others added 7 commits September 20, 2023 10:00
Update icicle crate name in examples and benches
* refactor

* lint

* Typo and spacing

---------

Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>
* Add cmake tests for cpp primitives

* Add cpp/cuda formatting

* Add conditional steps based on files changed for faster required checks

* Update runs-on for check files changed
@@ -3,6 +3,9 @@
#include "../../primitives/field.cuh"
#include "../../primitives/projective.cuh"
#if defined(G2_DEFINED)
#include "../../primitives/extension_field.cuh"
Copy link

@DmytroTym DmytroTym Sep 27, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think you need lines 6-8. I assume they were added to solve a compilation error but the problem was in the other place (in projective.cu). G2 curve isn't over the extension field and we should be able to speed up compilation quite a bit by removing this unnecessary import.

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done

@@ -21,6 +24,7 @@ namespace BW6_761 {
typedef Projective<point_field_t, scalar_t, b, gen_x, gen_y> projective_t;
typedef Affine<point_field_t> affine_t;
#if defined(G2_DEFINED)
typedef ExtensionField<PARAMS_BW6_761::fq_config> g2_point_field_t;

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This should be removed as well

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done, i removed the wrong imports as well from 377

@@ -62,6 +62,12 @@
#include "../curves/bls12_381/curve_config.cuh"
#include "../curves/bn254/curve_config.cuh"

__device__ void Error_UnsupportedType()

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@vhnatyk pls take a look, is this how Unresolved extern function '_Z21Error_UnsupportedTypev' should be resolved? Maybe #include "../curves/bw6_671/curve_config.cuh" can help?

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

while seems this may trick a linker - for compatibility at least with ntt templates - the file should be extended with like other curves

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

like it's 4years+ but seems there is no better workaround for templated dynamic shared memory other than Lei Mao blog's

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think this worked, im building it now

Copy link

@DmytroTym DmytroTym Oct 8, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But you haven't included struct SharedMemory<BW6_671::scalar_t> {...} and struct SharedMemory<BW6_671::projective_t> {...}, is it ok? @vhnatyk could you pls take a look again?

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

while seems this may trick a linker - for compatibility at least with ntt templates - the file should be extended with like other curves

@ImmanuelSegol - this may build with that stub function, but ntt should not pass correctly if shared mem kernel is used

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

pls remove as it just disables linking error and hides the issue

__device__ void Error_UnsupportedType() 
{
    // This function does not need a body, as it's just to throw a linker error.
}

and add the actual implementation like

template <>
struct SharedMemory<BW6_671::scalar_t> {
  __device__ BW6_671::scalar_t* getPointer()
  {
    extern __shared__ BW6_671::scalar_t s_scalar_t_bw6_671[];
    return s_scalar_t_bw6_671;
  }
};

template <>
struct SharedMemory<BW6_671::projective_t> {
  __device__ BW6_671::projective_t* getPointer()
  {
    extern __shared__ BW6_671::projective_t s_projective_t_bw6_671[];
    return s_projective_t_bw6_671;
  }
};
#endif //_SHAREDMEM_H_

}

type G2PointAffine struct {
X, Y ExtentionField

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That's not how G2 works for this curve unfortunately, the coordinates of points are from the same field as G1 points and not extension field elements

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done

return G2Point{
X: p.X,
Y: p.Y,
Z: ExtentionField{

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In principle it's probably better to use functions like projective_from_affine from CUDA, it's better to have a single source of truth. Like in this case, because G2 is actually not over extension field, and the size of limbs is wrong too, this function might work incorrectly

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks, this is a left over function that I didnt update.

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

removed it

Copy link

@DmytroTym DmytroTym left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Good job, I think you're almost done but I'd prefer that you delete newly auto-generated src/curves/bw6_761.rs and src/test_bw6_761.rs, they don't really work. Rust interface is not an immediate priority with this curve and we should introduce it as we refactor APIs on our feature branch IMO.

@@ -62,6 +62,12 @@
#include "../curves/bls12_381/curve_config.cuh"
#include "../curves/bn254/curve_config.cuh"

__device__ void Error_UnsupportedType()
Copy link

@DmytroTym DmytroTym Oct 8, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But you haven't included struct SharedMemory<BW6_671::scalar_t> {...} and struct SharedMemory<BW6_671::projective_t> {...}, is it ok? @vhnatyk could you pls take a look again?

!((point1->x == BW6_761::g2_point_field_t::zero()) && (point1->y == BW6_761::g2_point_field_t::zero()) &&
(point1->z == BW6_761::g2_point_field_t::zero())) &&
!((point2->x == BW6_761::g2_point_field_t::zero()) && (point2->y == BW6_761::g2_point_field_t::zero()) &&
(point2->z == BW6_761::g2_point_field_t::zero()));

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Trying to #include "projective.cu" leads to a compile error as the compiler doesn't know about g2_point_field_t (you deleted it from curve_config.cuh). I think you should revert to using point_field_t here.

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Well not #include but rather add to build

NttBatch(&nttResult, false, count, 0)
assert.NotEqual(t, nttResult, scalars)

assert.Equal(t, nttResult, nttResult)
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

asserting the equality of the same objects?

@@ -1,4 +1,3 @@
pub mod bls12_377;
pub mod bls12_381;
pub mod bn254;
pub mod bw6_761;
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

why is it removed?

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I asked to remove it during the previous review as I removed the (very much incorrectly working) autogenerated Rust wrappers for BW6-671

Copy link

@vhnatyk vhnatyk left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

just a few issues in the comments should be addressed, otherwise looks good 👍🏻

@liuxiaobleach liuxiaobleach merged commit 89a5081 into celer-network:main Oct 18, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
6 participants