#include <stdint.h>
#include "fp.cuh"
#include "fr.cuh"

Include dependency graph for g1.cuh:

This graph shows which files directly or indirectly include this file:

Data Structures
struct	g1a_t
	G1 point in affine coordinates. More...

struct	g1p_t
	G1 point in projective coordinates. More...

Macros
#define	G1_CUH

#define	G1P_ANYINF 1

Functions
__device__ __host__ void	g1a_fromUint64 (g1a_t &a, const uint64_t x, const uint64_t y)
	Converts arrays of uint64_t into a G1 point in affine coordinates. Each array must be uint64_t. This function does not validate if the coordinates are a valid point in the curve. More...

__device__ __host__ void	g1a_fromFp (g1a_t &a, const fp_t &x, const fp_t &y)
	Converts Fp values into a point in G1 in affine coordinates. This function does not validate if the coordinates are a valid point in the curve. More...

__device__ void	g1a_fromG1p (g1a_t &a, const g1p_t &p)
	Converts a point in projective coordinates into affine coordinates. More...

__device__ __host__ void	g1a_print (const char *s, const g1a_t &a)
	Print a standard representation of a, preceded by the user-set string s. More...

__device__ __host__ void	g1a_cpy (g1a_t &a, const g1a_t &b)
	Copy from b into a. More...

__device__ void	g1p_toUint64 (const g1p_t &p, uint64_t x, uint64_t y, uint64_t *z)
	Converts G1 point into arrays of uint64_t. Each array must be uint64_t[6] This function does not validate if the coordinates are a valid point in the curve. More...

__device__ __host__ void	g1p_fromUint64 (g1p_t &p, const uint64_t x, const uint64_t y, const uint64_t *z)

__device__ __host__ void	g1p_fromFp (g1p_t &p, fp_t &x, fp_t &y, fp_t &z)

__device__ void	g1p_fromG1a (g1p_t &p, const g1a_t &a)
	Convert a point in affine coordinates to projective coordinates. More...

__device__ __host__ void	g1p_print (const char *s, const g1p_t &p)
	Print a standard representation of p, preceded by the user-set string s. More...

__device__ __host__ void	g1p_cpy (g1p_t &p, const g1p_t &q)
	Copy from q into p. More...

__device__ bool	g1p_eq (const g1p_t &p, const g1p_t &q)
	Compares two projective points returns true when equal. This function compares if both parameters represent the same point on the curve. The equality is given by comparing X and Y coordinates divided by Z coordinates (p.X/p.Z == q.X/q.Z) && (p.Y/p.Z == q.Y/q.Z). Code-wise it is done by cross multiplication which also works for Z==0: (p.Xq.Z == q.Xp.Z) && (p.Yq.Z == q.Yp.Z) More...

__device__ bool	g1p_neq (const g1p_t &p, const g1p_t &q)
	Compares two projective points, returns true when not equal. This function compares if both parameters represent the distinct points on the curve. The equality is given by comparing X and Y coordinates divided by Z coordinates (p.X/p.Z == q.X/q.Z) && (p.Y/p.Z == q.Y/q.Z). Code-wise it is done by cross multiplication which also works for Z==0: (p.Xq.Z == q.Xp.Z) && (p.Yq.Z == q.Yp.Z) More...

__device__ bool	g1p_isInf (const g1p_t &p)
	Check if the value stored in p is the the/any point at infinity. This implementation uses (0, 1, 0) as the point at infinity. Alternatively, the macro G1P_ANYINF allows the point at infinity to be represented as (0, y, 0) where y!=0. More...

__device__ bool	g1p_isPoint (const g1p_t &p)
	Check if the value stored in p is a valid point on the G1 curve. More...

__device__ void	g1p_neg (g1p_t &p)
	Computes the negative of the point p. Due to negation map automorphism on Elliptic Curves in Weierstrass form, this operation is done by computing the additive inverse of the Y coordinate. More...

__device__ void	g1p_scale (g1p_t &p, const fp_t &s)
	Scale the coordinates of a projective point. This operation multiplies each coordinate of p by s: (x, y, z) := (xs, ys, z*s). More...

__device__ void	g1p_dbl (g1p_t &p)
	G1 point doubling, with write back: p=2*p. More...

__device__ void	g1p_add (g1p_t &p, const g1p_t &q)
	Computes the sum of two points q into p, using projective coordinates. and stores in p. More...

__device__ void	g1p_sub (g1p_t &p, const g1p_t &q)
	Point subtraction using projective coordinates. p ← p-q. More...

__device__ void	g1p_addsub (g1p_t &p, g1p_t &q)
	Stores the sum and difference of p and q into p and q. Projective p and q, p,q ← p+q,p-q. More...

__device__ void	g1p_mul (g1p_t &p, const fr_t &x)
	p ← k·p Point multiplication by scalar, in projective coordinates. That result is stored back into p. More...

__device__ __host__ void	g1a_inf (g1a_t &a)
	Set a to the point-at-infinity (0,0) More...

__device__ __host__ void	g1a_gen (g1a_t &a)
	Sets a to the generator point G1 of bls12_381. More...

__device__ __host__ void	g1p_inf (g1p_t &p)
	Set p to the point-at-infinity (0,1,0) More...

__device__ __host__ void	g1p_gen (g1p_t &p)
	Sets p to the generator point G1 of bls12_381. More...

__device__ void	g1p_fft (g1p_t output, const g1p_t input)
	FFT of size 512 over G1 with projective coordinates. Input and output arrays may overlap. This function must be called with 256 threads per block, i.e. dim3(256,1,1). No interleaving of data for different FFTs. More...

__device__ void	g1p_ift (g1p_t output, const g1p_t input)
	Inverse FFT of size 512 over G1 with projective coordinates. Input and output arrays may overlap. This function must be called with 256 threads per block, i.e. dim3(256,1,1). No interleaving of data for different FFTs. More...

__global__ void	g1p_fft_wrapper (g1p_t output, const g1p_t input)
	wrapper for g1p_fft: FFT for arrays of g1p_t with length 512 More...

__global__ void	g1p_ift_wrapper (g1p_t output, const g1p_t input)
	wrapper for g1p_ift: inverse FFT for arrays of g1p_t with length 512 More...

__global__ void	g1a_fromG1p_wrapper (g1a_t a, size_t len, const g1p_t p)
	Kernel wrappers, host-callable conversion of points in projective coordinates into affine coordinates. More...

__global__ void	g1p_eq_wrapper (uint8_t output, size_t len, const g1p_t p, const g1p_t *q)
	Kernel wrapper, host-callable comparison of arrays of g1p_t. More...

Macro Definition Documentation

◆ G1_CUH

#define G1_CUH

Definition at line 6 of file g1.cuh.

◆ G1P_ANYINF

#define G1P_ANYINF 1

Definition at line 14 of file g1.cuh.

Function Documentation

◆ g1a_cpy()

__device__ __host__ void g1a_cpy	(	g1a_t &	a,
		const g1a_t &	b
	)

Copy from b into a.

Parameters

[out]	a
[in]	b

Returns: void

Definition at line 65 of file g1a.cu.

◆ g1a_fromFp()

__device__ __host__ void g1a_fromFp	(	g1a_t &	a,
		const fp_t &	x,
		const fp_t &	y
	)

Converts Fp values into a point in G1 in affine coordinates. This function does not validate if the coordinates are a valid point in the curve.

Parameters

[out]	a	point in G1 in affine representation
[in]	x	x-coordinate
[in]	y	y-coordinate

Returns: void

Definition at line 36 of file g1a.cu.

◆ g1a_fromG1p()

__device__ void g1a_fromG1p	(	g1a_t &	a,
		const g1p_t &	p
	)

Converts a point in projective coordinates into affine coordinates.

Parameters

[out]	a	point in G1 in affine form
[in]	p	point in G1 in projective form

Returns: void

Definition at line 48 of file g1a.cu.

◆ g1a_fromG1p_wrapper()

__global__ void g1a_fromG1p_wrapper	(	g1a_t *	a,
		size_t	count,
		const g1p_t *	p
	)

Kernel wrappers, host-callable conversion of points in projective coordinates into affine coordinates.

Parameters

[out]	a	Array g1a_t[count]
[in]	count	number of elements in the array
[in]	p	Array g1p_t[count]

Returns: void

Definition at line 166 of file g1p.cu.

◆ g1a_fromUint64()

__device__ __host__ void g1a_fromUint64	(	g1a_t &	a,
		const uint64_t *	x,
		const uint64_t *	y
	)

Converts arrays of uint64_t into a G1 point in affine coordinates. Each array must be uint64_t. This function does not validate if the coordinates are a valid point in the curve.

Parameters

[out]	a	point in G1 in affine representation
[in]	x	Array with the x coordinate (uint64_t[6])
[in]	y	Array with the y coordinate (uint64_t[6])

Returns: void

Definition at line 21 of file g1a.cu.

◆ g1a_gen()

__device__ __host__ void g1a_gen ( g1a_t & a )

Sets a to the generator point G1 of bls12_381.

Parameters

a

Returns: void

Definition at line 102 of file g1a.cu.

◆ g1a_inf()

__device__ __host__ void g1a_inf ( g1a_t & a )

Set a to the point-at-infinity (0,0)

Parameters

a

Returns: void

Definition at line 91 of file g1a.cu.

◆ g1a_print()

__device__ __host__ void g1a_print	(	const char *	s,
		const g1a_t &	a
	)

Print a standard representation of a, preceded by the user-set string s.

Parameters

[out]	s	message string
[out]	a	point in g1 in affine form

Returns: void

Definition at line 77 of file g1a.cu.

◆ g1p_add()

__device__ void g1p_add	(	g1p_t &	p,
		const g1p_t &	q
	)

Computes the sum of two points q into p, using projective coordinates. and stores in p.

p ← p+q

Parameters

[in,out]	p	accumulator
[in]	q	second operand

Returns: void

Definition at line 29 of file g1p_add.cu.

Here is the caller graph for this function:

◆ g1p_addsub()

__device__ void g1p_addsub	(	g1p_t &	p,
		g1p_t &	q
	)

Stores the sum and difference of p and q into p and q. Projective p and q, p,q ← p+q,p-q.

Parameters

[in,out]	p	First parameter, stores p+q
[in,out]	q	Second parameter, stores p-q

Returns: void

Definition at line 18 of file g1p_addsub.cu.

Here is the caller graph for this function:

◆ g1p_cpy()

__device__ __host__ void g1p_cpy	(	g1p_t &	p,
		const g1p_t &	q
	)

Copy from q into p.

Parameters

[out]	p
[in]	q

Returns: void

Definition at line 67 of file g1p.cu.

Here is the caller graph for this function:

◆ g1p_dbl()

__device__ void g1p_dbl ( g1p_t & p )

G1 point doubling, with write back: p=2*p.

Parameters

[in,out] p

Returns: void

Definition at line 23 of file g1p_dbl.cu.

Here is the call graph for this function:

◆ g1p_eq()

__device__ bool g1p_eq	(	const g1p_t &	p,
		const g1p_t &	q
	)

Compares two projective points returns true when equal. This function compares if both parameters represent the same point on the curve. The equality is given by comparing X and Y coordinates divided by Z coordinates (p.X/p.Z == q.X/q.Z) && (p.Y/p.Z == q.Y/q.Z). Code-wise it is done by cross multiplication which also works for Z==0: (p.X*q.Z == q.X*p.Z) && (p.Y*q.Z == q.Y*p.Z)

Parameters

[in]	p	Projective G1 point
[in]	q	Projective G1 point

Returns: bool 1 if equal, 0 otherwise

Definition at line 23 of file g1p_compare.cu.

Here is the caller graph for this function:

◆ g1p_eq_wrapper()

__global__ void g1p_eq_wrapper	(	uint8_t *	eq,
		size_t	count,
		const g1p_t *	p,
		const g1p_t *	q
	)

Kernel wrapper, host-callable comparison of arrays of g1p_t.

Parameters

[out]	eq	Array of size count, will store the result of the comparison. eq[i]==1 iff p[i]==q[i]
[in]	count	number of elements to be compared
[in]	p	Array g1p_t[count]
[in]	q	Array g1p_t[count]

Returns: void

Definition at line 140 of file g1p.cu.

◆ g1p_fft()

__device__ void g1p_fft	(	g1p_t *	output,
		const g1p_t *	input
	)

FFT of size 512 over G1 with projective coordinates. Input and output arrays may overlap. This function must be called with 256 threads per block, i.e. dim3(256,1,1). No interleaving of data for different FFTs.

Parameters

[out]	output	pointer to 512 elements array of g1p_t
[in]	input	pointer to 512 elements array of g1p_t

Returns: void

Definition at line 24 of file g1p_fft.cu.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ g1p_fft_wrapper()

__global__ void g1p_fft_wrapper	(	g1p_t *	output,
		const g1p_t *	input
	)

wrapper for g1p_fft: FFT for arrays of g1p_t with length 512

Executes an FFT over many arrays of arrays of g1p_t with length 512. One array per block. input and output can overlap without side effects. There is no interleaving of data for different FFTs.

Parameters

[out]	output	pointer to 512*blocksize elements array of g1p_t
[in]	input	pointer to 512*blocksize elements array of g1p_t

Returns: void

Definition at line 336 of file g1p_fft.cu.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ g1p_fromFp()

__device__ __host__ void g1p_fromFp	(	g1p_t &	p,
		fp_t &	x,
		fp_t &	y,
		fp_t &	z
	)

inline

Definition at line 39 of file g1.cuh.

Here is the call graph for this function:

◆ g1p_fromG1a()

__device__ void g1p_fromG1a	(	g1p_t &	p,
		const g1a_t &	a
	)

Convert a point in affine coordinates to projective coordinates.

Parameters

[out]	p	point in projective coordinates
[in]	a	point in affine coordinates

Returns: void

Definition at line 51 of file g1p.cu.

◆ g1p_fromUint64()

__device__ __host__ void g1p_fromUint64	(	g1p_t &	p,
		const uint64_t *	x,
		const uint64_t *	y,
		const uint64_t *	z
	)

Here is the caller graph for this function:

◆ g1p_gen()

__device__ __host__ void g1p_gen ( g1p_t & p )

Sets p to the generator point G1 of bls12_381.

Parameters

p

Returns: void

Definition at line 106 of file g1p.cu.

Here is the caller graph for this function:

◆ g1p_ift()

__device__ void g1p_ift	(	g1p_t *	output,
		const g1p_t *	input
	)

Inverse FFT of size 512 over G1 with projective coordinates. Input and output arrays may overlap. This function must be called with 256 threads per block, i.e. dim3(256,1,1). No interleaving of data for different FFTs.

Parameters

[out]	output	pointer to 512 elements array of g1p_t
[in]	input	pointer to 512 elements array of g1p_t

Returns: void

Definition at line 178 of file g1p_fft.cu.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ g1p_ift_wrapper()

__global__ void g1p_ift_wrapper	(	g1p_t *	output,
		const g1p_t *	input
	)

wrapper for g1p_ift: inverse FFT for arrays of g1p_t with length 512

Executes an Inverse FFT over many arrays of arrays of g1p_t with length 512. One array per block. input and output can overlap without side effects. There is no interleaving of data for different FFTs.

Parameters

[out]	output	pointer to 512*blocksize elements array of g1p_t
[in]	input	pointer to 512*blocksize elements array of g1p_t

Returns: void

Definition at line 349 of file g1p_fft.cu.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ g1p_inf()

__device__ __host__ void g1p_inf ( g1p_t & p )

Set p to the point-at-infinity (0,1,0)

Parameters

[in] p

Returns: void

Definition at line 93 of file g1p.cu.

Here is the caller graph for this function:

◆ g1p_isInf()

__device__ bool g1p_isInf ( const g1p_t & p )

Check if the value stored in p is the the/any point at infinity. This implementation uses (0, 1, 0) as the point at infinity. Alternatively, the macro G1P_ANYINF allows the point at infinity to be represented as (0, y, 0) where y!=0.

The algebra used in this library sets the point at infinity to (0, 1, 0) (or (0, y, 0)), instead of the usual (1, 1, 0) used in other libs, due to following the Algorithm 7 in eprint 2015-1060.

Parameters

[in] p

Returns: bool 1 if the p is the point at infinity.

Definition at line 20 of file g1p_ispoint.cu.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ g1p_isPoint()

__device__ bool g1p_isPoint ( const g1p_t & p )

Check if the value stored in p is a valid point on the G1 curve.

Parameters

[in] p

Returns: bool 1 if is in the curve, zero otherwise.

Definition at line 34 of file g1p_ispoint.cu.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ g1p_mul()

__device__ void g1p_mul	(	g1p_t &	p,
		const fr_t &	k
	)

p ← k·p Point multiplication by scalar, in projective coordinates. That result is stored back into p.

Parameters

[in,out]	p	Multiplicand (stores result after call)
[in]	k	Fr operand

Returns: void

Definition at line 19 of file g1p_mul.cu.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ g1p_neg()

__device__ void g1p_neg ( g1p_t & p )

Computes the negative of the point p. Due to negation map automorphism on Elliptic Curves in Weierstrass form, this operation is done by computing the additive inverse of the Y coordinate.

Parameters

[in,out] p

Returns: void

Definition at line 16 of file g1p_neg.cu.

Here is the call graph for this function:

◆ g1p_neq()

__device__ bool g1p_neq	(	const g1p_t &	p,
		const g1p_t &	q
	)

Compares two projective points, returns true when not equal. This function compares if both parameters represent the distinct points on the curve. The equality is given by comparing X and Y coordinates divided by Z coordinates (p.X/p.Z == q.X/q.Z) && (p.Y/p.Z == q.Y/q.Z). Code-wise it is done by cross multiplication which also works for Z==0: (p.X*q.Z == q.X*p.Z) && (p.Y*q.Z == q.Y*p.Z)

Parameters

[in]	p	Projective G1 point
[in]	q	Projective G1 point

Returns: bool 0 if equal, 1 otherwise

Definition at line 68 of file g1p_compare.cu.

Here is the caller graph for this function:

◆ g1p_print()

__device__ __host__ void g1p_print	(	const char *	s,
		const g1p_t &	p
	)

Print a standard representation of p, preceded by the user-set string s.

Parameters

[out]	s	Description string
[out]	p	Point in g1

Returns: void

Definition at line 80 of file g1p.cu.

Here is the caller graph for this function:

◆ g1p_scale()

__device__ void g1p_scale	(	g1p_t &	p,
		const fp_t &	s
	)

Scale the coordinates of a projective point. This operation multiplies each coordinate of p by s: (x, y, z) := (x*s, y*s, z*s).

Parameters

[in,out]	p	Point in G1 (stores result after call)
[in]	s	Multiplicand in Fp. Must be nonzero.

Returns: void

Definition at line 16 of file g1p_scale.cu.

Here is the call graph for this function:

◆ g1p_sub()

__device__ void g1p_sub	(	g1p_t &	p,
		const g1p_t &	q
	)

Point subtraction using projective coordinates. p ← p-q.

Parameters

[in,out]	p
[in]	q

Returns: void

Definition at line 17 of file g1p_sub.cu.

◆ g1p_toUint64()

__device__ void g1p_toUint64	(	const g1p_t &	p,
		uint64_t *	x,
		uint64_t *	y,
		uint64_t *	z
	)

Converts G1 point into arrays of uint64_t. Each array must be uint64_t[6] This function does not validate if the coordinates are a valid point in the curve.

Parameters

[in]	p	point in G1
[out]	x	Array with the x coordinate
[out]	y	Array with the y coordinate
[out]	z	Array with the z coordinate

Returns: void

Definition at line 21 of file g1p.cu.

Here is the call graph for this function:

Data Structures

Macros

Functions

Macro Definition Documentation

◆ G1_CUH

◆ G1P_ANYINF

Function Documentation

◆ g1a_cpy()

◆ g1a_fromFp()

◆ g1a_fromG1p()

◆ g1a_fromG1p_wrapper()

◆ g1a_fromUint64()

◆ g1a_gen()

◆ g1a_inf()

◆ g1a_print()

◆ g1p_add()

◆ g1p_addsub()

◆ g1p_cpy()

◆ g1p_dbl()

◆ g1p_eq()

◆ g1p_eq_wrapper()

◆ g1p_fft()

◆ g1p_fft_wrapper()

◆ g1p_fromFp()

◆ g1p_fromG1a()

◆ g1p_fromUint64()

◆ g1p_gen()

◆ g1p_ift()

◆ g1p_ift_wrapper()

◆ g1p_inf()

◆ g1p_isInf()

◆ g1p_isPoint()

◆ g1p_mul()

◆ g1p_neg()

◆ g1p_neq()

◆ g1p_print()

◆ g1p_scale()

◆ g1p_sub()

◆ g1p_toUint64()