Skip to content

Commit

Permalink
ec/xyzz_t.hpp: rethink mem_t, make it device-only.
Browse files Browse the repository at this point in the history
  • Loading branch information
dot-asm committed Feb 28, 2024
1 parent 768171f commit d3d17ad
Showing 1 changed file with 6 additions and 11 deletions.
17 changes: 6 additions & 11 deletions ec/xyzz_t.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@ class xyzz_t {
public:
static const unsigned int degree = field_t::degree;

#ifdef __NVCC__
#ifdef __CUDACC__
class mem_t { friend class xyzz_t;
field_h X, Y, ZZZ, ZZ;

Expand All @@ -44,6 +44,7 @@ class xyzz_t {
Y = p.Y;
ZZZ = p.ZZZ;
ZZ = p.ZZ;
return *this;
}
inline __device__ void inf() { ZZZ.zero(); ZZ.zero(); }
};
Expand Down Expand Up @@ -100,7 +101,7 @@ class xyzz_t {
return p;
}

#ifdef __NVCC__
#ifdef __CUDACC__
class mem_t {
field_h X, Y;

Expand All @@ -126,7 +127,7 @@ class xyzz_t {
{ return inf; }

public:
inline __device__ operator affine_t() const
inline __host__ __device__ operator affine_t() const
{
bool inf = is_inf();
affine_t p;
Expand All @@ -135,20 +136,14 @@ class xyzz_t {
return p;
}

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

inline __host__ __device__ bool is_inf() const
inline __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() const
{
Expand Down

0 comments on commit d3d17ad

Please sign in to comment.