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.

Sign up for GitHub

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

separate affine_t and add some utility to ec/ #29

Open
wants to merge 2 commits into
base: main
Choose a base branch
from
Open

separate affine_t and add some utility to ec/ #29

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
6b8656d
separate affine_t and add some utility to ec/
sandsentinel Nov 16, 2023
698fb7a
fix equality operators
sandsentinel Nov 27, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
190 changes: 190 additions & 0 deletions ec/affine_t.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,190 @@
// Copyright Supranational LLC
// Licensed under the Apache License, Version 2.0, see LICENSE for details.
// SPDX-License-Identifier: Apache-2.0

#ifndef __SPPARK_EC_AFFINE_T_HPP__
#define __SPPARK_EC_AFFINE_T_HPP__

template<class field_t, class field_h = typename field_t::mem_t> class Affine_t;
template<class field_t, class field_h = typename field_t::mem_t> class Affine_inf_t;
template<class field_t> class jacobian_t;
template<class field_t, class field_h = typename field_t::mem_t> class xyzz_t;

#ifndef __CUDACC__
# undef __host__
# define __host__
# undef __device__
# define __device__
# undef __noinline__
# define __noinline__
#endif

template<class field_t, class field_h>
class Affine_t {
friend class Affine_inf_t<field_t>;
friend class jacobian_t<field_t>;
friend class xyzz_t<field_t>;

field_t X, Y;

public:
Affine_t(const field_t& x, const field_t& y) : X(x), Y(y) {}
inline __host__ __device__ Affine_t() {}

#ifdef __CUDA_ARCH__
inline __device__ bool is_inf() const
{ return (bool)(X.is_zero(Y)); }
#else
inline __host__ bool is_inf() const
{ return (bool)(X.is_zero() & Y.is_zero()); }
#endif

inline __host__ Affine_t& operator=(const jacobian_t<field_t>& a)
{
Y = 1/a.Z;
X = Y^2; // 1/Z^2
Y *= X; // 1/Z^3
X *= a.X; // X/Z^2
Y *= a.Y; // Y/Z^3
return *this;
}
inline __host__ Affine_t(const jacobian_t<field_t>& a) { *this = a; }

inline __host__ Affine_t& operator=(const xyzz_t<field_t>& a)
{
Y = 1/a.ZZZ;
X = Y * a.ZZ; // 1/Z
X = X^2; // 1/Z^2
X *= a.X; // X/Z^2
Y *= a.Y; // Y/Z^3
return *this;
}
inline __host__ Affine_t(const xyzz_t<field_t>& a) { *this = a; }

inline __host__ __device__ operator jacobian_t<field_t>() const
{
jacobian_t<field_t> p;
p.X = X;
p.Y = Y;
p.Z = field_t::one(is_inf());
return p;
}

inline __host__ __device__ operator xyzz_t<field_t>() const
{
xyzz_t<field_t> p;
p.X = X;
p.Y = Y;
p.ZZZ = p.ZZ = field_t::one(is_inf());
return p;
}

#ifdef __NVCC__
class mem_t {
field_h X, Y;

public:
inline __device__ operator Affine_t() const
{
Affine_t p;
p.X = X;
p.Y = Y;
return p;
}
};
#else
using mem_t = Affine_t;
#endif

#ifndef NDEBUG
friend inline __host__ __device__ bool operator==(const Affine_t& a, const Affine_t& b)
{ return (a.X == b.X) && (a.Y == b.Y); }

friend inline __host__ __device__ bool operator!=(const Affine_t& a, const Affine_t& b)
{ return (a.X != b.X) || (a.Y != b.Y); }

# if defined(_GLIBCXX_IOSTREAM) || defined(_IOSTREAM_) // non-standard
friend __host__ std::ostream& operator<<(std::ostream& os, const Affine_t& p)
{
return os << "X: " << p.X << std::endl
<< "Y: " << p.Y;
}
# endif
#endif
};

template<class field_t, class field_h>
class Affine_inf_t {
field_t X, Y;
bool inf;

inline __host__ __device__ bool is_inf() const
{ return inf; }

public:
inline __host__ __device__ operator Affine_t<field_t>() const
{
bool inf = is_inf();
Affine_t<field_t> p;
p.X = czero(X, inf);
p.Y = czero(Y, inf);
return p;
}

inline __host__ __device__ operator jacobian_t<field_t>() const
{
jacobian_t<field_t> p;
p.X = X;
p.Y = Y;
p.Z = field_t::one(is_inf());
return p;
}

inline __host__ __device__ operator xyzz_t<field_t>() const
{
xyzz_t<field_t> p;
p.X = X;
p.Y = Y;
p.ZZZ = p.ZZ = field_t::one(is_inf());
return p;
}

#ifdef __NVCC__
class mem_t {
field_h X, Y;
#ifdef __CUDACC__
int inf[sizeof(field_t)%16 ? 2 : 4];

inline __host__ __device__ bool is_inf() const
{ return inf[0]&1 != 0; }
#else
bool inf;

inline __host__ __device__ bool is_inf() const
{ return inf; }
#endif
public:
inline __device__ operator Affine_t<field_t>() const
{
bool inf = is_inf();
Affine_t<field_t> p;
p.X = czero((field_t)X, inf);
p.Y = czero((field_t)Y, inf);
return p;
}

inline __device__ operator Affine_inf_t() const
{
bool inf = is_inf();
Affine_inf_t p;
p.X = czero((field_t)X, inf);
p.Y = czero((field_t)Y, inf);
p.inf = inf;
return p;
}
};
#else
using mem_t = Affine_inf_t;
#endif
};
#endif
51 changes: 28 additions & 23 deletions ec/jacobian_t.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,39 +5,27 @@
#ifndef __SPPARK_EC_JACOBIAN_T_HPP__
#define __SPPARK_EC_JACOBIAN_T_HPP__

template<class field_t> class jacobian_t {
#ifndef __SPPARK_EC_AFFINE_T_HPP__
# include "affine_t.hpp"
#endif

template<class field_t>
class jacobian_t {
friend class Affine_t<field_t>;
friend class Affine_inf_t<field_t>;

field_t X, Y, Z;

inline operator const void*() const { return this; }
inline operator void*() { return this; }

public:
using affine_t = Affine_t<field_t>;

jacobian_t() {}
jacobian_t(const field_t& x, const field_t& y, const field_t& z) :
X(x), Y(y), Z(z) {}

class affine_t { friend jacobian_t;
field_t X, Y;

public:
affine_t() {}
affine_t(const field_t& x, const field_t& y) : X(x), Y(y) {}

inline bool is_inf() const
{ return (bool)(X.is_zero() & Y.is_zero()); }

inline affine_t& operator=(const jacobian_t& a)
{
Y = 1/a.Z;
X = Y^2; // 1/Z^2
Y *= X; // 1/Z^3
X *= a.X; // X/Z^2
Y *= a.Y; // Y/Z^3
return *this;
}
inline affine_t(const jacobian_t& a) { *this = a; }
};

inline operator affine_t() const { return affine_t(*this); }

inline jacobian_t& operator=(const affine_t& a)
Expand Down Expand Up @@ -538,5 +526,22 @@ template<class field_t> class jacobian_t {
}
*this = p3;
}

#ifndef NDEBUG
friend inline bool operator==(const jacobian_t& a, const jacobian_t& b)
{ return (a.X == b.X) && (a.Y == b.Y) && (a.Z == b.Z); }
Copy link
Collaborator

Choose a reason for hiding this comment

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

This is not how you compare points in Jacobian coordinates for equality. See POINT_IS_EQUAL_IMPL for a template. Same naturally applies to the inequality test, which should read just !(a == b). Or maybe compiler does it for you without having to define operator!=?

Copy link
Collaborator Author

Choose a reason for hiding this comment

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

This is not how you compare points in Jacobian coordinates for equality.

Right, same thing with xyzz_t. I'll fix both.

Copy link
Collaborator Author

Choose a reason for hiding this comment

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

Should be fixed now for both jacobian_t and xyzz_t, thanks! I also removed __device__ from == and != because I realised mont_t didn't have them and thought it would be better to leave it to a later time.


friend inline bool operator!=(const jacobian_t& a, const jacobian_t& b)
{ return (a.X != b.X) || (a.Y != b.Y) || (a.Z != b.Z); }

# if defined(_GLIBCXX_IOSTREAM) || defined(_IOSTREAM_) // non-standard
friend std::ostream& operator<<(std::ostream& os, const jacobian_t& p)
{
return os << "X: " << p.X << std::endl
<< "Y: " << p.Y << std::endl
<< "Z: " << p.Z;
}
# endif
#endif
};
#endif
Loading