KernelsRTf.cu File Reference

Kernel for CUDA RT calculations. More...

#include "InterfaceCUDA.h"

Functions
__device__ __inline__ float	common1 (float t, float xr, float yr, float dxr, float dyr)
__device__ __inline__ float	common2 (float t, float xr, float yr, float dxr, float dyr)
__device__ __inline__ float	gp (float t, float xr, float yr, float zr, float dxr, float dyr, float dzr)
__device__ __inline__ float	gh (float t, float xr, float yr, float zr, float dxr, float dyr, float dzr)
__device__ __inline__ float	ge (float t, float xr, float yr, float zr, float dxr, float dyr, float dzr)
__device__ __inline__ float	gpl (float t, float xr, float yr, float zr, float dxr, float dyr, float dzr)
__device__ __inline__ void	np (float xr, float yr, float zr, float(&out)[3])
__device__ __inline__ void	nh (float xr, float yr, float zr, float(&out)[3])
__device__ __inline__ void	ne (float xr, float yr, float zr, float(&out)[3])
__device__ __inline__ void	npl (float xr, float yr, float zr, float(&out)[3])
__device__ __inline__ void	transfRays (float *x, float *y, float *z, float *dx, float *dy, float *dz, int i, bool inv=false)
__host__ std::array< dim3, 2 >	_initCUDA (reflparamsf ctp, float epsilon, float t0, int _nTot, int nBlocks, int nThreads)
__global__ void	propagateRaysToTarget (float *xs, float *ys, float *zs, float *dxs, float *dys, float *dzs, float *xt, float *yt, float *zt, float *dxt, float *dyt, float *dzt)
void	callRTKernel (reflparamsf ctp, cframef *fr_in, cframef *fr_out, float epsilon, float t0, int nBlocks, int nThreads)

Variables
__constant__ float	conrt [CSIZERT]
__constant__ float	mat [16]
__constant__ int	nTot
__constant__ int	cflip
__constant__ int	ctype

Detailed Description

Kernel for CUDA RT calculations.

Contains kernel for RT calculations.

Function Documentation

_initCUDA()

__host__ std::array<dim3, 2> _initCUDA	(	reflparamsf	ctp,
		float	epsilon,
		float	t0,
		int	_nTot,
		int	nBlocks,
		int	nThreads
	)

Initialize CUDA.

Instantiate program and populate constant memory.

Parameters

ctp	reflparamsf object containing target surface parameters.
epsilon	Precision of NR method.
t0	Starting guess for NR method.
_nTot	Total number of rays in beam.
nBlocks	Number of blocks per grid.
nThreads	Number of threads per block.

Returns

BT Array of two dim3 objects.

callRTKernel()

void callRTKernel	(	reflparamsf	ctp,
		cframef *	fr_in,
		cframef *	fr_out,
		float	epsilon,
		float	t0,
		int	nBlocks,
		int	nThreads
	)

Call ray-trace Kernel.

Calculate a new frame of rays on a target, given an input frame of rays.

Parameters

ctp	reflparamsf object containing target surface parameters.
fr_in	Pointer to input cframef object.
fr_out	Pointer to output cframef object.
epsilon	Precision for NR method.
t0	Starting guess for NR method.
nBlocks	Number of blocks in GPU grid.
nThreads	Number of threads in block.

See also

reflparamsf

cframef

common1()

__device__ __inline__ float common1	(	float	t,
		float	xr,
		float	yr,
		float	dxr,
		float	dyr
	)

Calculate common factor 1.

Parameters

t	Scaling factor.
xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
dxr	x co-ordinate of ray direction
dyr	y co-ordinate of ray direction

common2()

__device__ __inline__ float common2	(	float	t,
		float	xr,
		float	yr,
		float	dxr,
		float	dyr
	)

Calculate common factor 2.

Parameters

t	Scaling factor.
xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
dxr	x co-ordinate of ray direction
dyr	y co-ordinate of ray direction

ge()

__device__ __inline__ float ge	(	float	t,
		float	xr,
		float	yr,
		float	zr,
		float	dxr,
		float	dyr,
		float	dzr
	)

Calculate ray-ellipsoid intersection.

Parameters

t	Scaling factor.
xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
zr	z co-ordinate of ray.
dxr	x co-ordinate of ray direction
dyr	y co-ordinate of ray direction
dzr	z co-ordinate of ray direction

gh()

__device__ __inline__ float gh	(	float	t,
		float	xr,
		float	yr,
		float	zr,
		float	dxr,
		float	dyr,
		float	dzr
	)

Calculate ray-hyperboloid intersection.

Parameters

t	Scaling factor.
xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
zr	z co-ordinate of ray.
dxr	x co-ordinate of ray direction
dyr	y co-ordinate of ray direction
dzr	z co-ordinate of ray direction

gp()

__device__ __inline__ float gp	(	float	t,
		float	xr,
		float	yr,
		float	zr,
		float	dxr,
		float	dyr,
		float	dzr
	)

Calculate ray-paraboloid intersection.

Parameters

t	Scaling factor.
xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
zr	z co-ordinate of ray.
dxr	x co-ordinate of ray direction
dyr	y co-ordinate of ray direction
dzr	z co-ordinate of ray direction

gpl()

__device__ __inline__ float gpl	(	float	t,
		float	xr,
		float	yr,
		float	zr,
		float	dxr,
		float	dyr,
		float	dzr
	)

Calculate ray-plane intersection.

Parameters

t	Scaling factor.
xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
zr	z co-ordinate of ray.
dxr	x co-ordinate of ray direction
dyr	y co-ordinate of ray direction
dzr	z co-ordinate of ray direction

ne()

__device__ __inline__ void ne	(	float	xr,
		float	yr,
		float	zr,
		float(&)	out[3]
	)

Calculate ellipsoid normals.

Parameters

xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
zr	z co-ordinate of ray.
out	Array of 3 float.

nh()

__device__ __inline__ void nh	(	float	xr,
		float	yr,
		float	zr,
		float(&)	out[3]
	)

Calculate hyperboloid normals.

Parameters

xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
zr	z co-ordinate of ray.
out	Array of 3 float.

np()

__device__ __inline__ void np	(	float	xr,
		float	yr,
		float	zr,
		float(&)	out[3]
	)

Calculate paraboloid normals.

Parameters

xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
zr	z co-ordinate of ray.
out	Array of 3 float.

npl()

__device__ __inline__ void npl	(	float	xr,
		float	yr,
		float	zr,
		float(&)	out[3]
	)

Calculate plane normals.

Parameters

xr	x co-ordinate of ray.
yr	y co-ordinate of ray.
zr	z co-ordinate of ray.
out	Array of 3 float.

propagateRaysToTarget()

__global__ void propagateRaysToTarget	(	float *	xs,
		float *	ys,
		float *	zs,
		float *	dxs,
		float *	dys,
		float *	dzs,
		float *	xt,
		float *	yt,
		float *	zt,
		float *	dxt,
		float *	dyt,
		float *	dzt
	)

Optimize ray-target distance.

Uses a Newton Rhapson (NR) method to find the point of ray-surface intersection.

Parameters

xs	Array of ray x co-ordinates.
ys	Array of ray y co-ordinates.
zs	Array of ray z co-ordinates.
dxs	Array of ray x directions.
dys	Array of ray y directions.
dzs	Array of ray z directions.
xt	Array of ray x co-ordinates, to be filled.
yt	Array of ray y co-ordinates, to be filled.
zt	Array of ray z co-ordinates, to be filled.
dxt	Array of ray x directions, to be filled.
dyt	Array of ray y directions, to be filled.
dzt	Array of ray z directions, to be filled.

transfRays()

__device__ __inline__ void transfRays	(	float *	x,
		float *	y,
		float *	z,
		float *	dx,
		float *	dy,
		float *	dz,
		int	i,
		bool	inv = false
	)

Transform rays to surface restframe.

Parameters

x	Array of ray x co-ordinates.
y	Array of ray y co-ordinates.
z	Array of ray z co-ordinates.
dx	Array of ray x directions.
dy	Array of ray y directions.
dz	Array of ray z directions.
i	Index of co-ordinate.
inv	Whether to apply inverse transformation.