kernels.h File Reference

GPU kernel definitions. More...

#include <stdio.h>

Include dependency graph for kernels.h:

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

Go to the source code of this file.

Functions
__device__ double2	subtract (double2 a, double2 b)
	subtraction operation for 2 double2 values More...
__device__ double2	add (double2 a, double2 b)
	addition operation for 2 double2 values More...
__device__ double2	pow (double2 a, int b)
	power operation for a double2 More...
__device__ double2	mult (double2 a, double b)
	multiplication operation for double2 and double values More...
__device__ double2	mult (double2 a, double2 b)
	multiplication operation for 2 double2 values More...
__global__ void	make_cufftDoubleComplex (double *in, double2 *out)
	transforms an array of doubles into double2's More...
__device__ unsigned int	getGid3d3d ()
	Indexing of threads on grid. More...
__device__ unsigned int	getBid3d3d ()
	Indexing of blocks on device. More...
__device__ unsigned int	getTid3d3d ()
	Indexing of threads in a block on device. More...
__global__ void	is_eq (bool *a, bool *b, bool *ans)
	checks to arrays to see if they are equal More...
__device__ double	complexMagnitude (double2 in)
	Calculates magnitude of complex number. $|a + ib|$. More...
__global__ void	complexMultiply (double2 *in1, double2 *in2, double2 *out)
	Complex multiplication of two input arrays. More...
__host__ __device__ double2	complexMultiply (double2 in1, double2 in2)
	Complex multiplication of two input values. More...
__device__ double2	make_complex (double in, int evolution_type)
	Transforms field value into operator. More...
__global__ void	complexAbsSum (double2 *in1, double2 *in2, double *out)
	Sums the absolute value of two complex arrays. More...
__global__ void	complexMagnitude (double2 *in, double *out)
	Complex magnitude of a double2 array. More...
__device__ double	complexMagnitudeSquared (double2 in)
	Return the squared magnitude of a complex number. $|(a+{i}b)*(a-{i}b)|$. More...
__global__ void	complexMagnitudeSquared (double2 *in, double *out)
	Complex magnitude of a double2 array. More...
__global__ void	complexMagnitudeSquared (double2 *in, double2 *out)
	Complex magnitude of a double2 array. More...
__device__ double2	conjugate (double2 in)
	Returns conjugate of the a complex number. More...
__device__ double2	realCompMult (double scalar, double2 comp)
	Multiply real scalar by a complex number. More...
__global__ void	cMult (double2 *in1, double2 *in2, double2 *out)
	Kernel for complex multiplication. More...
__global__ void	cMultPhi (double2 *in1, double *in2, double2 *out)
	Kernel for multiplcation with real array and complex array. More...
__global__ void	cMultDensity (double2 *in1, double2 *in2, double2 *out, double dt, double mass, int gstate, double gDenConst)
	Kernel for complex multiplication with nonlinear density term. More...
__global__ void	cMultDensity_ast (EqnNode_gpu *eqn, double2 *in, double2 *out, double dx, double dy, double dz, double time, int e_num, double dt, double mass, int gstate, double gDenConst)
	Kernel for complex multiplication with nonlinear density term. More...
__global__ void	pinVortex (double2 *in1, double2 *in2, double2 *out)
	Hold vortex at specified position. Not implemented. cMultPhi should implement required functionality. More...
__global__ void	vecMult (double2 *in, double *factor, double2 *out)
	Complex field scaling and renormalisation. Used mainly post-FFT. More...
__global__ void	l2_norm (double *in1, double *in2, double *in3, double *out)
	performs the l2 normalization of the provided terms More...
__global__ void	l2_norm (double2 *in1, double2 *in2, double2 *in3, double *out)
	performs the l2 normalization of the provided terms More...
__global__ void	l2_norm (double *in1, double *in2, double *out)
	performs the l2 normalization of the provided terms More...
__global__ void	l2_norm (double2 *in1, double2 *in2, double *out)
	performs the l2 normalization of the provided terms More...
__global__ void	scalarDiv (double2 *in, double factor, double2 *out)
	Complex field scaling and renormalisation. Used mainly post-FFT. More...
__global__ void	scalarDiv (double *in, double factor, double *out)
	Real field scaling and renormalisation. Used mainly post-FFT. More...
__global__ void	scalarMult (double2 *in, double factor, double2 *out)
	Complex field scaling and renormalisation. Used mainly post-FFT. More...
__global__ void	scalarPow (double2 *in, double param, double2 *out)
	Complex field raised to a power. More...
__global__ void	vecConjugate (double2 *in, double2 *out)
	Conjugate of double2*. More...
__global__ void	scalarDiv1D (double2 *, double2 *)
	Complex field scaling and renormalisation. Not implemented. Use scalarDiv. More...
__global__ void	scalarDiv2D (double2 *, double2 *)
	Complex field scaling and renormalisation. Not implemented. Use scalarDiv. More...
__global__ void	scalarDiv_wfcNorm (double2 *in, double dr, double *pSum, double2 *out)
	Used as part of multipass to renormalise the wavefucntion. More...
__global__ void	reduce (double2 *in, double *out)
	Not implemented. More...
__global__ void	thread_test (double *input, double *output)
	Performs wavefunction renormalisation using parallel summation and applying scalarDiv_wfcNorm. More...
__global__ void	multipass (double2 *input, double2 *output, int pass)
	Performs wavefunction renormalisation using parallel summation and applying scalarDiv_wfcNorm. More...
__global__ void	multipass (double *input, double *output)
	Performs parallel summation of double arrays. More...
__global__ void	angularOp (double omega, double dt, double2 *wfc, double *xpyypx, double2 *out)
	Calculates angular momentum. Not fully implemented. Handled in post-processing instead. More...
__global__ void	ast_mult (double *array, double *array_out, EqnNode_gpu *eqn, double dx, double dy, double dz, double time, int element_num)
	Multiplication of array with AST. More...
__global__ void	ast_cmult (double2 *array, double2 *array_out, EqnNode_gpu *eqn, double dx, double dy, double dz, double time, int element_num)
	Complex multiplication of array with AST. More...
__global__ void	ast_op_mult (double2 *array, double2 *array_out, EqnNode_gpu *eqn, double dx, double dy, double dz, double time, int element_num, int evolution_type, double dt)
	Multiplication of array with AST Operator. More...
__device__ double2	real_ast (double val, double dt)
	Function to find AST operator in real-time. More...
__device__ double2	im_ast (double val, double dt)
	Function to find AST operator in imaginary-time. More...
__global__ void	zeros (bool *in, bool *out)
	Sets boolean array to 0. More...
__global__ void	set_eq (double *in1, double *in2)
	Sets in2 to be equal to in1. More...
__global__ void	energyCalc (double2 *wfc, double2 *op, double dt, double2 *energy, int gnd_state, int op_space, double sqrt_omegaz_mass, double gDenConst)
	Calculates energy of the current state during evolution. Not implemented. More...
__device__ double2	braKetMult (double2 in1, double2 in2)
	Performs bra-ket state multiplication. Not fully implemented. More...
__global__ void	pSum (double *in1, double *output, int pass)
	Performs parallel sum. Not verified. I use multipass instead. More...

Detailed Description

GPU kernel definitions.

Author

Lee J. O'Riordan (mlxd)

Date

12/11/2015

Version

0.1

DESCRIPTION

Kernel definitions for all CUDA-enabled routines for solving GPE.

Definition in file kernels.h.

Function Documentation

add()

__device__ double2 add	(	double2	a,
		double2	b
	)

addition operation for 2 double2 values

angularOp()

__global__ void angularOp	(	double	omega,
		double	dt,
		double2 *	wfc,
		double *	xpyypx,
		double2 *	out
	)

Calculates angular momentum. Not fully implemented. Handled in post-processing instead.

Parameters

omega	Harmonic trap rotation frequency
dt	Time-step for evolution
wfc	Wavefunction
xpyypx	L_z operator
out	Output of calculation

ast_cmult()

__global__ void ast_cmult	(	double2 *	array,
		double2 *	array_out,
		EqnNode_gpu *	eqn,
		double	dx,
		double	dy,
		double	dz,
		double	time,
		int	element_num
	)

Complex multiplication of array with AST.

ast_mult()

__global__ void ast_mult	(	double *	array,
		double *	array_out,
		EqnNode_gpu *	eqn,
		double	dx,
		double	dy,
		double	dz,
		double	time,
		int	element_num
	)

Multiplication of array with AST.

ast_op_mult()

__global__ void ast_op_mult	(	double2 *	array,
		double2 *	array_out,
		EqnNode_gpu *	eqn,
		double	dx,
		double	dy,
		double	dz,
		double	time,
		int	element_num,
		int	evolution_type,
		double	dt
	)

Multiplication of array with AST Operator.

braKetMult()

__device__ double2 braKetMult ( double2  in1, double2  in2  )

inline

Performs bra-ket state multiplication. Not fully implemented.

Parameters

in1	Bra
in2	Ket

Returns

<Bra|Ket>

cMult()

__global__ void cMult	(	double2 *	in1,
		double2 *	in2,
		double2 *	out
	)

Kernel for complex multiplication.

Multiplication for linear, non-linear and phase-imprinting of the condensate.

Parameters

in1	Wavefunction input
in2	Evolution operator input
out	Pass by reference output for multiplcation result

cMultDensity()

__global__ void cMultDensity	(	double2 *	in1,
		double2 *	in2,
		double2 *	out,
		double	dt,
		double	mass,
		int	gstate,
		double	gDenConst
	)

Kernel for complex multiplication with nonlinear density term.

Parameters

in1	Wavefunction input
in2	Evolution operator input
out	Pass by reference output for multiplication result
dt	Timestep for evolution
mass	Atomic species mass
gState	If performing real (1) or imaginary (0) time evolution
gDenConst	a constant for evolution

cMultDensity_ast()

__global__ void cMultDensity_ast	(	EqnNode_gpu *	eqn,
		double2 *	in,
		double2 *	out,
		double	dx,
		double	dy,
		double	dz,
		double	time,
		int	e_num,
		double	dt,
		double	mass,
		int	gstate,
		double	gDenConst
	)

Kernel for complex multiplication with nonlinear density term.

Parameters

GPU	AST
in	Wavefunction input
out	Wavefunction output
dx
dy
dz
time
element	number in AST
dt	Timestep for evolution
mass	Atomic species mass
gState	If performing real (1) or imaginary (0) time evolution
gDenConst	a constant for evolution

cMultPhi()

__global__ void cMultPhi	(	double2 *	in1,
		double *	in2,
		double2 *	out
	)

Kernel for multiplcation with real array and complex array.

Parameters

in1	Wavefunction input
in2	Evolution operator input
out	Pass by reference output for multiplcation result

complexAbsSum()

__global__ void complexAbsSum	(	double2 *	in1,
		double2 *	in2,
		double *	out
	)

Sums the absolute value of two complex arrays.

Parameters

Array	1
Array	2
Output

complexMagnitude() [1/2]

__device__ double complexMagnitude

(

double2

in

)

Calculates magnitude of complex number. $|a + ib|$.

Helper functions for complex numbers

Parameters

in	Complex number

Returns

Magnitude of complex number

complexMagnitude() [2/2]

__global__ void complexMagnitude	(	double2 *	in,
		double *	out
	)

Complex magnitude of a double2 array.

complexMagnitudeSquared() [1/3]

__device__ double complexMagnitudeSquared

(

double2

in

)

Return the squared magnitude of a complex number. $|(a+{i}b)*(a-{i}b)|$.

Parameters

in	Complex number

Returns

Absolute-squared complex number

complexMagnitudeSquared() [2/3]

__global__ void complexMagnitudeSquared	(	double2 *	in,
		double *	out
	)

Complex magnitude of a double2 array.

complexMagnitudeSquared() [3/3]

__global__ void complexMagnitudeSquared	(	double2 *	in,
		double2 *	out
	)

Complex magnitude of a double2 array.

complexMultiply() [1/2]

__global__ void complexMultiply	(	double2 *	in1,
		double2 *	in2,
		double2 *	out
	)

Complex multiplication of two input arrays.

Parameters

Input	1
Input	2
Output

complexMultiply() [2/2]

__host__ __device__ double2 complexMultiply	(	double2	in1,
		double2	in2
	)

Complex multiplication of two input values.

Parameters

Input	1
Input	2

Returns

Output

conjugate()

__device__ double2 conjugate

(

double2

in

)

Returns conjugate of the a complex number.

Parameters

in	Number to be conjugated

Returns

Conjugated complex number

energyCalc()

__global__ void energyCalc	(	double2 *	wfc,
		double2 *	op,
		double	dt,
		double2 *	energy,
		int	gnd_state,
		int	op_space,
		double	sqrt_omegaz_mass,
		double	gDenConst
	)

Calculates energy of the current state during evolution. Not implemented.

Parameters

wfc	Wavefunction
op	Operator to calculate energy for.
dt	Time-step for evolution
energy	Energy result output
gnd_state	Wavefunction
op_space	Check if position space with non-linear term or not.
sqrt_omegaz_mass	sqrt(omegaZ/mass), part of the nonlin interaction term.

getBid3d3d()

__device__ unsigned int getBid3d3d

(

)

Indexing of blocks on device.

getGid3d3d()

__device__ unsigned int getGid3d3d

(

)

Indexing of threads on grid.

getTid3d3d()

__device__ unsigned int getTid3d3d

(

)

Indexing of threads in a block on device.

im_ast()

__device__ double2 im_ast	(	double	val,
		double	dt
	)

Function to find AST operator in imaginary-time.

is_eq()

__global__ void is_eq	(	bool *	a,
		bool *	b,
		bool *	ans
	)

checks to arrays to see if they are equal

l2_norm() [1/4]

__global__ void l2_norm	(	double *	in1,
		double *	in2,
		double *	in3,
		double *	out
	)

performs the l2 normalization of the provided terms

l2_norm() [2/4]

__global__ void l2_norm	(	double2 *	in1,
		double2 *	in2,
		double2 *	in3,
		double *	out
	)

performs the l2 normalization of the provided terms

l2_norm() [3/4]

__global__ void l2_norm	(	double *	in1,
		double *	in2,
		double *	out
	)

performs the l2 normalization of the provided terms

l2_norm() [4/4]

__global__ void l2_norm	(	double2 *	in1,
		double2 *	in2,
		double *	out
	)

performs the l2 normalization of the provided terms

make_complex()

__device__ double2 make_complex	(	double	in,
		int	evolution_type
	)

Transforms field value into operator.

Parameters

Input	value
Evolution	type (0 for imaginary, 1 for real)

Returns

complex output

make_cufftDoubleComplex()

__global__ void make_cufftDoubleComplex	(	double *	in,
		double2 *	out
	)

transforms an array of doubles into double2's

mult() [1/2]

__device__ double2 mult	(	double2	a,
		double	b
	)

multiplication operation for double2 and double values

mult() [2/2]

__device__ double2 mult	(	double2	a,
		double2	b
	)

multiplication operation for 2 double2 values

multipass() [1/2]

__global__ void multipass	(	double2 *	input,
		double2 *	output,
		int	pass
	)

Performs wavefunction renormalisation using parallel summation and applying scalarDiv_wfcNorm.

Parameters

input	Wavefunction to be renormalised
output	Pass by reference return of renormalised wavefunction
pass	Number of passes performed by routine

multipass() [2/2]

__global__ void multipass	(	double *	input,
		double *	output
	)

Performs parallel summation of double arrays.

Parameters

input	double array
input	double array after summation

pinVortex()

__global__ void pinVortex	(	double2 *	in1,
		double2 *	in2,
		double2 *	out
	)

Hold vortex at specified position. Not implemented. cMultPhi should implement required functionality.

Parameters

in1	Wavefunction input
in2	Evolution operator input
out	Pass by reference output for multiplcation result

pow()

__device__ double2 pow	(	double2	a,
		int	b
	)

power operation for a double2

Parameters

base	number
power

Referenced by sepAvg().

Here is the caller graph for this function:

pSum()

__global__ void pSum	(	double *	in1,
		double *	output,
		int	pass
	)

Performs parallel sum. Not verified. I use multipass instead.

Parameters

in1	That which must be summed
output	That which has been summed
pass	Number of passes

real_ast()

__device__ double2 real_ast	(	double	val,
		double	dt
	)

Function to find AST operator in real-time.

realCompMult()

__device__ double2 realCompMult	(	double	scalar,
		double2	comp
	)

Multiply real scalar by a complex number.

Parameters

scalar	Scalar multiplier
comp	Complex multiplicand

Returns

Result of scalar * comp

reduce()

__global__ void reduce	(	double2 *	in,
		double *	out
	)

Not implemented.

Parameters

in	Input field
out	Output values

scalarDiv() [1/2]

__global__ void scalarDiv	(	double2 *	in,
		double	factor,
		double2 *	out
	)

Complex field scaling and renormalisation. Used mainly post-FFT.

Parameters

in	Complex field to be scaled (divided, not multiplied)
factor	Scaling factor to be used
out	Pass by reference output for result

scalarDiv() [2/2]

__global__ void scalarDiv	(	double *	in,
		double	factor,
		double *	out
	)

Real field scaling and renormalisation. Used mainly post-FFT.

Parameters

in	Real field to be scaled (divided, not multiplied)
factor	Scaling factor to be used
out	Pass by reference output for result

scalarDiv1D()

__global__ void scalarDiv1D	(	double2 *	,
		double2 *
	)

Complex field scaling and renormalisation. Not implemented. Use scalarDiv.

scalarDiv2D()

__global__ void scalarDiv2D	(	double2 *	,
		double2 *
	)

Complex field scaling and renormalisation. Not implemented. Use scalarDiv.

scalarDiv_wfcNorm()

__global__ void scalarDiv_wfcNorm	(	double2 *	in,
		double	dr,
		double *	pSum,
		double2 *	out
	)

Used as part of multipass to renormalise the wavefucntion.

Parameters

in	Complex field to be renormalised
dr	Smallest area element of grid (dx*dy)
pSum	GPU array used to store intermediate results during parallel summation

scalarMult()

__global__ void scalarMult	(	double2 *	in,
		double	factor,
		double2 *	out
	)

Complex field scaling and renormalisation. Used mainly post-FFT.

Parameters

in	Complex field to be scaled (multiplied, not divided)
factor	Scaling factor to be used
out	Pass by reference output for result

scalarPow()

__global__ void scalarPow	(	double2 *	in,
		double	param,
		double2 *	out
	)

Complex field raised to a power.

Parameters

in	Complex field to be scaled (multiplied, not divided)
power	parameter
out	Pass by reference output for result

set_eq()

__global__ void set_eq	(	double *	in1,
		double *	in2
	)

Sets in2 to be equal to in1.

subtract()

__device__ double2 subtract	(	double2	a,
		double2	b
	)

subtraction operation for 2 double2 values

thread_test()

__global__ void thread_test	(	double *	input,
		double *	output
	)

Performs wavefunction renormalisation using parallel summation and applying scalarDiv_wfcNorm.

Parameters

input	Wavefunction to be renormalised
output	Pass by reference return of renormalised wavefunction
pass	Number of passes performed by routine

vecConjugate()

__global__ void vecConjugate	(	double2 *	in,
		double2 *	out
	)

Conjugate of double2*.

Parameters

in	Complex field to be conjugated
out	Pass by reference output for result

vecMult()

__global__ void vecMult	(	double2 *	in,
		double *	factor,
		double2 *	out
	)

Complex field scaling and renormalisation. Used mainly post-FFT.

Parameters

in	Complex field to be scaled (divided, not multiplied)
factor	Scaling vector to be used
out	Pass by reference output for result

zeros()

__global__ void zeros	(	bool *	in,
		bool *	out
	)

Sets boolean array to 0.