Kernelsf.cu File Reference

Kernels for CUDA PO calculations. More...

#include "InterfaceCUDA.h"

Functions
__host__ std::array< dim3, 2 >	initCUDA (float k, float epsilon, int gt, int gs, float t_direction, int nBlocks, int nThreads)
__device__ void	fieldAtPoint (float *d_xs, float *d_ys, float *d_zs, float *d_nxs, float *d_nys, float *d_nzs, cuFloatComplex *d_Jx, cuFloatComplex *d_Jy, cuFloatComplex *d_Jz, cuFloatComplex *d_Mx, cuFloatComplex *d_My, cuFloatComplex *d_Mz, float(&point)[3], float *d_A, cuFloatComplex(&d_ei)[3], cuFloatComplex(&d_hi)[3])
__global__ void	GpropagateBeam_0_PEC (float *d_xs, float *d_ys, float *d_zs, float *d_A, float *d_xt, float *d_yt, float *d_zt, float *d_nxs, float *d_nys, float *d_nzs, float *d_nxt, float *d_nyt, float *d_nzt, cuFloatComplex *d_Jx, cuFloatComplex *d_Jy, cuFloatComplex *d_Jz, cuFloatComplex *d_Mx, cuFloatComplex *d_My, cuFloatComplex *d_Mz, cuFloatComplex *d_Jxt, cuFloatComplex *d_Jyt, cuFloatComplex *d_Jzt, cuFloatComplex *d_Mxt, cuFloatComplex *d_Myt, cuFloatComplex *d_Mzt)
__global__ void	GpropagateBeam_0 (float *d_xs, float *d_ys, float *d_zs, float *d_A, float *d_xt, float *d_yt, float *d_zt, float *d_nxs, float *d_nys, float *d_nzs, float *d_nxt, float *d_nyt, float *d_nzt, cuFloatComplex *d_Jx, cuFloatComplex *d_Jy, cuFloatComplex *d_Jz, cuFloatComplex *d_Mx, cuFloatComplex *d_My, cuFloatComplex *d_Mz, cuFloatComplex *d_Jxt, cuFloatComplex *d_Jyt, cuFloatComplex *d_Jzt, cuFloatComplex *d_Mxt, cuFloatComplex *d_Myt, cuFloatComplex *d_Mzt)
__global__ void	GpropagateBeam_2_PEC (float *d_xs, float *d_ys, float *d_zs, float *d_A, float *d_xt, float *d_yt, float *d_zt, float *d_nxs, float *d_nys, float *d_nzs, float *d_nxt, float *d_nyt, float *d_nzt, cuFloatComplex *d_Jx, cuFloatComplex *d_Jy, cuFloatComplex *d_Jz, cuFloatComplex *d_Mx, cuFloatComplex *d_My, cuFloatComplex *d_Mz, cuFloatComplex *d_Jxt, cuFloatComplex *d_Jyt, cuFloatComplex *d_Jzt, cuFloatComplex *d_Mxt, cuFloatComplex *d_Myt, cuFloatComplex *d_Mzt, cuFloatComplex *d_Ext, cuFloatComplex *d_Eyt, cuFloatComplex *d_Ezt, cuFloatComplex *d_Hxt, cuFloatComplex *d_Hyt, cuFloatComplex *d_Hzt)
__global__ void	GpropagateBeam_1 (float *d_xs, float *d_ys, float *d_zs, float *d_A, float *d_xt, float *d_yt, float *d_zt, float *d_nxs, float *d_nys, float *d_nzs, cuFloatComplex *d_Jx, cuFloatComplex *d_Jy, cuFloatComplex *d_Jz, cuFloatComplex *d_Mx, cuFloatComplex *d_My, cuFloatComplex *d_Mz, cuFloatComplex *d_Ext, cuFloatComplex *d_Eyt, cuFloatComplex *d_Ezt, cuFloatComplex *d_Hxt, cuFloatComplex *d_Hyt, cuFloatComplex *d_Hzt)
__global__ void	GpropagateBeam_2 (float *d_xs, float *d_ys, float *d_zs, float *d_A, float *d_xt, float *d_yt, float *d_zt, float *d_nxs, float *d_nys, float *d_nzs, float *d_nxt, float *d_nyt, float *d_nzt, cuFloatComplex *d_Jx, cuFloatComplex *d_Jy, cuFloatComplex *d_Jz, cuFloatComplex *d_Mx, cuFloatComplex *d_My, cuFloatComplex *d_Mz, cuFloatComplex *d_Jxt, cuFloatComplex *d_Jyt, cuFloatComplex *d_Jzt, cuFloatComplex *d_Mxt, cuFloatComplex *d_Myt, cuFloatComplex *d_Mzt, cuFloatComplex *d_Ext, cuFloatComplex *d_Eyt, cuFloatComplex *d_Ezt, cuFloatComplex *d_Hxt, cuFloatComplex *d_Hyt, cuFloatComplex *d_Hzt)
__global__ void	GpropagateBeam_3 (float *d_xs, float *d_ys, float *d_zs, float *d_A, float *d_xt, float *d_yt, float *d_zt, float *d_nxs, float *d_nys, float *d_nzs, float *d_nxt, float *d_nyt, float *d_nzt, cuFloatComplex *d_Jx, cuFloatComplex *d_Jy, cuFloatComplex *d_Jz, cuFloatComplex *d_Mx, cuFloatComplex *d_My, cuFloatComplex *d_Mz, cuFloatComplex *d_Ext, cuFloatComplex *d_Eyt, cuFloatComplex *d_Ezt, cuFloatComplex *d_Hxt, cuFloatComplex *d_Hyt, cuFloatComplex *d_Hzt, float *d_Prxt, float *d_Pryt, float *d_Przt)
__device__ void	farfieldAtPoint (float *d_xs, float *d_ys, float *d_zs, float *d_nxs, float *d_nys, float *d_nzs, cuFloatComplex *d_Jx, cuFloatComplex *d_Jy, cuFloatComplex *d_Jz, cuFloatComplex *d_Mx, cuFloatComplex *d_My, cuFloatComplex *d_Mz, float(&r_hat)[3], float *d_A, cuFloatComplex(&e)[3], cuFloatComplex(&h)[3])
void __global__	GpropagateBeam_4 (float *d_xs, float *d_ys, float *d_zs, float *d_nxs, float *d_nys, float *d_nzs, float *d_A, float *d_xt, float *d_yt, cuFloatComplex *d_Jx, cuFloatComplex *d_Jy, cuFloatComplex *d_Jz, cuFloatComplex *d_Mx, cuFloatComplex *d_My, cuFloatComplex *d_Mz, cuFloatComplex *d_Ext, cuFloatComplex *d_Eyt, cuFloatComplex *d_Ezt, cuFloatComplex *d_Hxt, cuFloatComplex *d_Hyt, cuFloatComplex *d_Hzt)
void __device__	scalarfieldAtPoint (float *d_xs, float *d_ys, float *d_zs, cuFloatComplex *d_sfs, float(&point)[3], float *d_A, cuFloatComplex &e)
void __global__	GpropagateBeam_5 (float *d_xs, float *d_ys, float *d_zs, float *d_A, float *d_xt, float *d_yt, float *d_zt, cuFloatComplex *d_sfs, cuFloatComplex *d_sft)
__host__ void	_arrC1ToCUDAC (float *rarr, float *iarr, cuFloatComplex *carr, int size)
__host__ void	_arrC3ToCUDAC (float *r1arr, float *r2arr, float *r3arr, float *i1arr, float *i2arr, float *i3arr, cuFloatComplex *c1arr, cuFloatComplex *c2arr, cuFloatComplex *c3arr, int size)
__host__ void	_arrCUDACToC1 (cuFloatComplex *carr, float *rarr, float *iarr, int size)
__host__ void	_arrCUDACToC3 (cuFloatComplex *c1arr, cuFloatComplex *c2arr, cuFloatComplex *c3arr, float *r1arr, float *r2arr, float *r3arr, float *i1arr, float *i2arr, float *i3arr, int size)
void	callKernelf_JM (c2Bundlef *res, reflparamsf source, reflparamsf target, reflcontainerf *cs, reflcontainerf *ct, c2Bundlef *currents, float k, float epsilon, float t_direction, int nBlocks, int nThreads)
void	callKernelf_EH (c2Bundlef *res, reflparamsf source, reflparamsf target, reflcontainerf *cs, reflcontainerf *ct, c2Bundlef *currents, float k, float epsilon, float t_direction, int nBlocks, int nThreads)
void	callKernelf_JMEH (c4Bundlef *res, reflparamsf source, reflparamsf target, reflcontainerf *cs, reflcontainerf *ct, c2Bundlef *currents, float k, float epsilon, float t_direction, int nBlocks, int nThreads)
void	callKernelf_EHP (c2rBundlef *res, reflparamsf source, reflparamsf target, reflcontainerf *cs, reflcontainerf *ct, c2Bundlef *currents, float k, float epsilon, float t_direction, int nBlocks, int nThreads)
void	callKernelf_FF (c2Bundlef *res, reflparamsf source, reflparamsf target, reflcontainerf *cs, reflcontainerf *ct, c2Bundlef *currents, float k, float epsilon, float t_direction, int nBlocks, int nThreads)
void	callKernelf_scalar (arrC1f *res, reflparamsf source, reflparamsf target, reflcontainerf *cs, reflcontainerf *ct, arrC1f *inp, float k, float epsilon, float t_direction, int nBlocks, int nThreads)

Variables
__constant__ cuFloatComplex	con [CSIZE]
__constant__ float	eye [3][3]
__constant__ int	g_s
__constant__ int	g_t

Detailed Description

Kernels for CUDA PO calculations.

Contains kernels for PO calculations. Multiple kernels are defined, each one optimized for a certain calculation.

Function Documentation

_arrC1ToCUDAC()

__host__ void _arrC1ToCUDAC	(	float *	rarr,
		float *	iarr,
		cuFloatComplex *	carr,
		int	size
	)

Convert 2 arrays of floats to 1 array of cuComplex

Parameters

rarr	Real part of complex array.
iarr	Real part of complex array.
carr	Array of cuFloatComplex, to be filled.
size	Size of arrays.

_arrC3ToCUDAC()

__host__ void _arrC3ToCUDAC	(	float *	r1arr,
		float *	r2arr,
		float *	r3arr,
		float *	i1arr,
		float *	i2arr,
		float *	i3arr,
		cuFloatComplex *	c1arr,
		cuFloatComplex *	c2arr,
		cuFloatComplex *	c3arr,
		int	size
	)

Convert 6 arrays of floats to 3 arrays of cuComplex

Parameters

r1arr	Real part of complex array.
r2arr	Real part of complex array.
r3arr	Real part of complex array.
i1arr	Imaginary part of complex array.
i2arr	Imaginary part of complex array.
i3arr	Imaginary part of complex array.
c1arr	Array of cuFloatComplex, to be filled.
c2arr	Array of cuFloatComplex, to be filled.
c3arr	Array of cuFloatComplex, to be filled.
size	Size of arrays.

_arrCUDACToC1()

__host__ void _arrCUDACToC1	(	cuFloatComplex *	carr,
		float *	rarr,
		float *	iarr,
		int	size
	)

Convert 1 array of cuComplex to 2 arrays of floats.

Parameters

carr	Array of cuFloatComplex.
rarr	Real part of complex array, to be filled.
iarr	Imaginary part of complex array, to be filled.
size	Size of arrays.

_arrCUDACToC3()

__host__ void _arrCUDACToC3	(	cuFloatComplex *	c1arr,
		cuFloatComplex *	c2arr,
		cuFloatComplex *	c3arr,
		float *	r1arr,
		float *	r2arr,
		float *	r3arr,
		float *	i1arr,
		float *	i2arr,
		float *	i3arr,
		int	size
	)

Convert 3 arrays of cuComplex to 6 arrays of floats.

Parameters

c1arr	Array of cuFloatComplex.
c2arr	Array of cuFloatComplex.
c3arr	Array of cuFloatComplex.
r1arr	Real part of complex array, to be filled.
r2arr	Real part of complex array, to be filled.
r3arr	Real part of complex array, to be filled.
i1arr	Imaginary part of complex array, to be filled.
i2arr	Imaginary part of complex array, to be filled.
i3arr	Imaginary part of complex array, to be filled.
size	Size of arrays.

callKernelf_EH()

void callKernelf_EH	(	c2Bundlef *	res,
		reflparamsf	source,
		reflparamsf	target,
		reflcontainerf *	cs,
		reflcontainerf *	ct,
		c2Bundlef *	currents,
		float	k,
		float	epsilon,
		float	t_direction,
		int	nBlocks,
		int	nThreads
	)

Call EH kernel.

Calculate E, H fields on a target surface using CUDA.

Parameters

res	Pointer to c2Bundlef object.
source	reflparamsf object containing source surface parameters.
target	reflparamsf object containing target surface parameters.
cs	Pointer to reflcontainerf object containing source grids.
ct	Pointer to reflcontainerf object containing target grids.
currents	Pointer to c2Bundlef object containing source currents.
k	Wavenumber of radiation in 1 /mm.
epsilon	Relative permittivity of source surface.
t_direction	Time direction (experimental!).
nBlocks	Number of blocks in GPU grid.
nThreads	Number of threads in a block.

See also

c2Bundlef

reflparamsf

reflcontainerf

callKernelf_EHP()

void callKernelf_EHP	(	c2rBundlef *	res,
		reflparamsf	source,
		reflparamsf	target,
		reflcontainerf *	cs,
		reflcontainerf *	ct,
		c2Bundlef *	currents,
		float	k,
		float	epsilon,
		float	t_direction,
		int	nBlocks,
		int	nThreads
	)

Call EHP kernel.

Calculate reflected E, H fields and P, the reflected Poynting vectorfield, on a target surface using CUDA.

Parameters

res	Pointer to c2rBundlef object.
source	reflparamsf object containing source surface parameters.
target	reflparamsf object containing target surface parameters.
cs	Pointer to reflcontainerf object containing source grids.
ct	Pointer to reflcontainerf object containing target grids.
currents	Pointer to c2Bundlef object containing source currents.
k	Wavenumber of radiation in 1 /mm.
epsilon	Relative permittivity of source surface.
t_direction	Time direction (experimental!).
nBlocks	Number of blocks in GPU grid.
nThreads	Number of threads in a block.

See also

c2rBundlef

c2Bundlef

reflparamsf

reflcontainerf

callKernelf_FF()

void callKernelf_FF	(	c2Bundlef *	res,
		reflparamsf	source,
		reflparamsf	target,
		reflcontainerf *	cs,
		reflcontainerf *	ct,
		c2Bundlef *	currents,
		float	k,
		float	epsilon,
		float	t_direction,
		int	nBlocks,
		int	nThreads
	)

Call FF kernel.

Calculate E, H fields on a far-field target surface using CUDA.

Parameters

res	Pointer to c2Bundlef object.
source	reflparamsf object containing source surface parameters.
target	reflparamsf object containing target surface parameters.
cs	Pointer to reflcontainerf object containing source grids.
ct	Pointer to reflcontainerf object containing target grids.
currents	Pointer to c2Bundlef object containing source currents.
k	Wavenumber of radiation in 1 /mm.
epsilon	Relative permittivity of source surface.
t_direction	Time direction (experimental!).
nBlocks	Number of blocks in GPU grid.
nThreads	Number of threads in a block.

See also

c2Bundlef

reflparamsf

reflcontainerf

callKernelf_JM()

void callKernelf_JM	(	c2Bundlef *	res,
		reflparamsf	source,
		reflparamsf	target,
		reflcontainerf *	cs,
		reflcontainerf *	ct,
		c2Bundlef *	currents,
		float	k,
		float	epsilon,
		float	t_direction,
		int	nBlocks,
		int	nThreads
	)

Call JM kernel.

Calculate J, M currents on a target surface using CUDA.

Parameters

res	Pointer to c2Bundlef object.
source	reflparamsf object containing source surface parameters.
target	reflparamsf object containing target surface parameters.
cs	Pointer to reflcontainerf object containing source grids.
ct	Pointer to reflcontainerf object containing target grids.
currents	Pointer to c2Bundlef object containing source currents.
k	Wavenumber of radiation in 1 /mm.
epsilon	Relative permittivity of source surface.
t_direction	Time direction (experimental!).
nBlocks	Number of blocks in GPU grid.
nThreads	Number of threads in a block.

See also

c2Bundlef

reflparamsf

reflcontainerf

callKernelf_JMEH()

void callKernelf_JMEH	(	c4Bundlef *	res,
		reflparamsf	source,
		reflparamsf	target,
		reflcontainerf *	cs,
		reflcontainerf *	ct,
		c2Bundlef *	currents,
		float	k,
		float	epsilon,
		float	t_direction,
		int	nBlocks,
		int	nThreads
	)

Call JMEH kernel.

Calculate J, M currents and E, H fields on a target surface using CUDA.

Parameters

res	Pointer to c4Bundlef object.
source	reflparamsf object containing source surface parameters.
target	reflparamsf object containing target surface parameters.
cs	Pointer to reflcontainerf object containing source grids.
ct	Pointer to reflcontainerf object containing target grids.
currents	Pointer to c2Bundlef object containing source currents.
k	Wavenumber of radiation in 1 /mm.
epsilon	Relative permittivity of source surface.
t_direction	Time direction (experimental!).
nBlocks	Number of blocks in GPU grid.
nThreads	Number of threads in a block.

See also

c4Bundlef

c2Bundlef

reflparamsf

reflcontainerf

callKernelf_scalar()

void callKernelf_scalar	(	arrC1f *	res,
		reflparamsf	source,
		reflparamsf	target,
		reflcontainerf *	cs,
		reflcontainerf *	ct,
		arrC1f *	inp,
		float	k,
		float	epsilon,
		float	t_direction,
		int	nBlocks,
		int	nThreads
	)

Call scalar kernel.

Calculate scalar field on a target surface using CUDA.

Parameters

res	Pointer to arrC1f object.
source	reflparamsf object containing source surface parameters.
target	reflparamsf object containing target surface parameters.
cs	Pointer to reflcontainerf object containing source grids.
ct	Pointer to reflcontainerf object containing target grids.
inp	Pointer to arrC1f object containing source field.
k	Wavenumber of radiation in 1 /mm.
epsilon	Relative permittivity of source surface.
t_direction	Time direction (experimental!).
nBlocks	Number of blocks in GPU grid.
nThreads	Number of threads in a block.

See also

arrC1f

reflparamsf

reflcontainerf

farfieldAtPoint()

__device__ void farfieldAtPoint	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_nxs,
		float *	d_nys,
		float *	d_nzs,
		cuFloatComplex *	d_Jx,
		cuFloatComplex *	d_Jy,
		cuFloatComplex *	d_Jz,
		cuFloatComplex *	d_Mx,
		cuFloatComplex *	d_My,
		cuFloatComplex *	d_Mz,
		float(&)	r_hat[3],
		float *	d_A,
		cuFloatComplex(&)	e[3],
		cuFloatComplex(&)	h[3]
	)

Calculate total E and H field at point on far-field target.

Parameters

d_xs	Array containing source point's x-coordinate.
d_ys	Array containing source point's y-coordinate.
d_zs	Array containing source point's z-coordinate.
d_nxs	Array containing source point's normal x-component.
d_nys	Array containing source point's y-component.
d_nzs	Array containing source point's z-component.
d_Jx	Array containing source J x-component.
d_Jy	Array containing source J y-component.
d_Jz	Array containing source J z-component.
d_Mx	Array containing source M x-component.
d_My	Array containing source M y-component.
d_Mz	Array containing source M z-component.
r_hat	Array of 3 float, containing xyz coordinates of target point direction.
d_A	Array containing area elements.
e	Array of 3 cuFloatComplex, to be filled with E-field at point.

fieldAtPoint()

__device__ void fieldAtPoint	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_nxs,
		float *	d_nys,
		float *	d_nzs,
		cuFloatComplex *	d_Jx,
		cuFloatComplex *	d_Jy,
		cuFloatComplex *	d_Jz,
		cuFloatComplex *	d_Mx,
		cuFloatComplex *	d_My,
		cuFloatComplex *	d_Mz,
		float(&)	point[3],
		float *	d_A,
		cuFloatComplex(&)	d_ei[3],
		cuFloatComplex(&)	d_hi[3]
	)

Calculate total E and H field at point on target.

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_nxs	Array containing source normals x-component.
d_nys	Array containing source normals y-component.
d_nzs	Array containing source normals z-component.
d_Jx	Array containing source J x-component.
d_Jy	Array containing source J y-component.
d_Jz	Array containing source J z-component.
d_Mx	Array containing source M x-component.
d_My	Array containing source M y-component.
d_Mz	Array containing source M z-component.
point	Array of 3 float, containing xyz coordinates of target point.
d_A	Array containing area elements.
d_ei	Array of 3 cuFloatComplex, to be filled with E-field at point.
d_hi	Array of 3 cuFloatComplex, to be filled with H-field at point.

GpropagateBeam_0()

__global__ void GpropagateBeam_0	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_A,
		float *	d_xt,
		float *	d_yt,
		float *	d_zt,
		float *	d_nxs,
		float *	d_nys,
		float *	d_nzs,
		float *	d_nxt,
		float *	d_nyt,
		float *	d_nzt,
		cuFloatComplex *	d_Jx,
		cuFloatComplex *	d_Jy,
		cuFloatComplex *	d_Jz,
		cuFloatComplex *	d_Mx,
		cuFloatComplex *	d_My,
		cuFloatComplex *	d_Mz,
		cuFloatComplex *	d_Jxt,
		cuFloatComplex *	d_Jyt,
		cuFloatComplex *	d_Jzt,
		cuFloatComplex *	d_Mxt,
		cuFloatComplex *	d_Myt,
		cuFloatComplex *	d_Mzt
	)

Calculate JM on target.

Kernel for calculating J, M currents on target surface.

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_A	Array containing area elements.
d_xt	Array containing target points x-coordinate.
d_yt	Array containing target points y-coordinate.
d_zt	Array containing target points z-coordinate.
d_nxs	Array containing source normals x-component.
d_nys	Array containing source normals y-component.
d_nzs	Array containing source normals z-component.
d_nxt	Array containing target norms x-component.
d_nyt	Array containing target norms y-component.
d_nzt	Array containing target norms z-component.
d_Jx	Array containing source J x-component.
d_Jy	Array containing source J y-component.
d_Jz	Array containing source J z-component.
d_Mx	Array containing source M x-component.
d_My	Array containing source M y-component.
d_Mz	Array containing source M z-component.
d_Jxt	Array to be filled with target J x-component.
d_Jyt	Array to be filled with target J y-component.
d_Jzt	Array to be filled with target J z-component.
d_Mxt	Array to be filled with target M x-component.
d_Myt	Array to be filled with target M y-component.
d_Mzt	Array to be filled with target M z-component.

GpropagateBeam_0_PEC()

__global__ void GpropagateBeam_0_PEC	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_A,
		float *	d_xt,
		float *	d_yt,
		float *	d_zt,
		float *	d_nxs,
		float *	d_nys,
		float *	d_nzs,
		float *	d_nxt,
		float *	d_nyt,
		float *	d_nzt,
		cuFloatComplex *	d_Jx,
		cuFloatComplex *	d_Jy,
		cuFloatComplex *	d_Jz,
		cuFloatComplex *	d_Mx,
		cuFloatComplex *	d_My,
		cuFloatComplex *	d_Mz,
		cuFloatComplex *	d_Jxt,
		cuFloatComplex *	d_Jyt,
		cuFloatComplex *	d_Jzt,
		cuFloatComplex *	d_Mxt,
		cuFloatComplex *	d_Myt,
		cuFloatComplex *	d_Mzt
	)

Calculate JM on PEC target

Kernel for calculating J, M currents on a perfectly conducting target surface. This reduces the calculation of the currents to Je = 2 n^Hinc and Jm = -2 n^Einc

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_A	Array containing area elements.
d_xt	Array containing target points x-coordinate.
d_yt	Array containing target points y-coordinate.
d_zt	Array containing target points z-coordinate.
d_nxs	Array containing source normals x-component.
d_nys	Array containing source normals y-component.
d_nzs	Array containing source normals z-component.
d_nxt	Array containing target norms x-component.
d_nyt	Array containing target norms y-component.
d_nzt	Array containing target norms z-component.
d_Jx	Array containing source J x-component.
d_Jy	Array containing source J y-component.
d_Jz	Array containing source J z-component.
d_Mx	Array containing source M x-component.
d_My	Array containing source M y-component.
d_Mz	Array containing source M z-component.
d_Jxt	Array to be filled with target J x-component.
d_Jyt	Array to be filled with target J y-component.
d_Jzt	Array to be filled with target J z-component.
d_Mxt	Array to be filled with target M x-component.
d_Myt	Array to be filled with target M y-component.
d_Mzt	Array to be filled with target M z-component.

GpropagateBeam_1()

__global__ void GpropagateBeam_1	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_A,
		float *	d_xt,
		float *	d_yt,
		float *	d_zt,
		float *	d_nxs,
		float *	d_nys,
		float *	d_nzs,
		cuFloatComplex *	d_Jx,
		cuFloatComplex *	d_Jy,
		cuFloatComplex *	d_Jz,
		cuFloatComplex *	d_Mx,
		cuFloatComplex *	d_My,
		cuFloatComplex *	d_Mz,
		cuFloatComplex *	d_Ext,
		cuFloatComplex *	d_Eyt,
		cuFloatComplex *	d_Ezt,
		cuFloatComplex *	d_Hxt,
		cuFloatComplex *	d_Hyt,
		cuFloatComplex *	d_Hzt
	)

Calculate EH on target.

Kernel for calculating E, H fields on target surface.

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_A	Array containing area elements.
d_xt	Array containing target points x-coordinate.
d_yt	Array containing target points y-coordinate.
d_zt	Array containing target points z-coordinate.
d_nxs	Array containing source normals x-component.
d_nys	Array containing source normals y-component.
d_nzs	Array containing source normals z-component.
d_Jx	Array containing source J x-component.
d_Jy	Array containing source J y-component.
d_Jz	Array containing source J z-component.
d_Mx	Array containing source M x-component.
d_My	Array containing source M y-component.
d_Mz	Array containing source M z-component.
d_Ext	Array to be filled with target E x-component.
d_Eyt	Array to be filled with target E y-component.
d_Ezt	Array to be filled with target E z-component.
d_Hxt	Array to be filled with target H x-component.
d_Hyt	Array to be filled with target H y-component.
d_Hzt	Array to be filled with target H z-component.

GpropagateBeam_2()

__global__ void GpropagateBeam_2	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_A,
		float *	d_xt,
		float *	d_yt,
		float *	d_zt,
		float *	d_nxs,
		float *	d_nys,
		float *	d_nzs,
		float *	d_nxt,
		float *	d_nyt,
		float *	d_nzt,
		cuFloatComplex *	d_Jx,
		cuFloatComplex *	d_Jy,
		cuFloatComplex *	d_Jz,
		cuFloatComplex *	d_Mx,
		cuFloatComplex *	d_My,
		cuFloatComplex *	d_Mz,
		cuFloatComplex *	d_Jxt,
		cuFloatComplex *	d_Jyt,
		cuFloatComplex *	d_Jzt,
		cuFloatComplex *	d_Mxt,
		cuFloatComplex *	d_Myt,
		cuFloatComplex *	d_Mzt,
		cuFloatComplex *	d_Ext,
		cuFloatComplex *	d_Eyt,
		cuFloatComplex *	d_Ezt,
		cuFloatComplex *	d_Hxt,
		cuFloatComplex *	d_Hyt,
		cuFloatComplex *	d_Hzt
	)

Calculate JM on target.

Kernel for calculating J, M currents on target surface.

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_A	Array containing area elements.
d_xt	Array containing target points x-coordinate.
d_yt	Array containing target points y-coordinate.
d_zt	Array containing target points z-coordinate.
d_nxs	Array containing source normals x-component.
d_nys	Array containing source normals y-component.
d_nzs	Array containing source normals z-component.
d_nxt	Array containing target norms x-component.
d_nyt	Array containing target norms y-component.
d_nzt	Array containing target norms z-component.
d_Jx	Array containing source J x-component.
d_Jy	Array containing source J y-component.
d_Jz	Array containing source J z-component.
d_Mx	Array containing source M x-component.
d_My	Array containing source M y-component.
d_Mz	Array containing source M z-component.
d_Jxt	Array to be filled with target J x-component.
d_Jyt	Array to be filled with target J y-component.
d_Jzt	Array to be filled with target J z-component.
d_Mxt	Array to be filled with target M x-component.
d_Myt	Array to be filled with target M y-component.
d_Mzt	Array to be filled with target M z-component.
d_Ext	Array to be filled with target E x-component.
d_Eyt	Array to be filled with target E y-component.
d_Ezt	Array to be filled with target E z-component.
d_Hxt	Array to be filled with target H x-component.
d_Hyt	Array to be filled with target H y-component.
d_Hzt	Array to be filled with target H z-component.

GpropagateBeam_2_PEC()

__global__ void GpropagateBeam_2_PEC	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_A,
		float *	d_xt,
		float *	d_yt,
		float *	d_zt,
		float *	d_nxs,
		float *	d_nys,
		float *	d_nzs,
		float *	d_nxt,
		float *	d_nyt,
		float *	d_nzt,
		cuFloatComplex *	d_Jx,
		cuFloatComplex *	d_Jy,
		cuFloatComplex *	d_Jz,
		cuFloatComplex *	d_Mx,
		cuFloatComplex *	d_My,
		cuFloatComplex *	d_Mz,
		cuFloatComplex *	d_Jxt,
		cuFloatComplex *	d_Jyt,
		cuFloatComplex *	d_Jzt,
		cuFloatComplex *	d_Mxt,
		cuFloatComplex *	d_Myt,
		cuFloatComplex *	d_Mzt,
		cuFloatComplex *	d_Ext,
		cuFloatComplex *	d_Eyt,
		cuFloatComplex *	d_Ezt,
		cuFloatComplex *	d_Hxt,
		cuFloatComplex *	d_Hyt,
		cuFloatComplex *	d_Hzt
	)

Calculate JMEH on PEC target

Kernel for calculating J, M, E & H currents on a perfectly conducting target surface. This reduces the calculation of the currents to Je = 2 n^Hinc and Jm = 2 n^Einc

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_A	Array containing area elements.
d_xt	Array containing target points x-coordinate.
d_yt	Array containing target points y-coordinate.
d_zt	Array containing target points z-coordinate.
d_nxs	Array containing source normals x-component.
d_nys	Array containing source normals y-component.
d_nzs	Array containing source normals z-component.
d_nxt	Array containing target norms x-component.
d_nyt	Array containing target norms y-component.
d_nzt	Array containing target norms z-component.
d_Jx	Array containing source J x-component.
d_Jy	Array containing source J y-component.
d_Jz	Array containing source J z-component.
d_Mx	Array containing source M x-component.
d_My	Array containing source M y-component.
d_Mz	Array containing source M z-component.
d_Jxt	Array to be filled with target J x-component.
d_Jyt	Array to be filled with target J y-component.
d_Jzt	Array to be filled with target J z-component.
d_Mxt	Array to be filled with target M x-component.
d_Myt	Array to be filled with target M y-component.
d_Mzt	Array to be filled with target M z-component.
d_Ext	Array to be filled with target E x-component.
d_Eyt	Array to be filled with target E y-component.
d_Ezt	Array to be filled with target E z-component.
d_Hxt	Array to be filled with target H x-component.
d_Hyt	Array to be filled with target H y-component.
d_Hzt	Array to be filled with target H z-component.

GpropagateBeam_3()

__global__ void GpropagateBeam_3	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_A,
		float *	d_xt,
		float *	d_yt,
		float *	d_zt,
		float *	d_nxs,
		float *	d_nys,
		float *	d_nzs,
		float *	d_nxt,
		float *	d_nyt,
		float *	d_nzt,
		cuFloatComplex *	d_Jx,
		cuFloatComplex *	d_Jy,
		cuFloatComplex *	d_Jz,
		cuFloatComplex *	d_Mx,
		cuFloatComplex *	d_My,
		cuFloatComplex *	d_Mz,
		cuFloatComplex *	d_Ext,
		cuFloatComplex *	d_Eyt,
		cuFloatComplex *	d_Ezt,
		cuFloatComplex *	d_Hxt,
		cuFloatComplex *	d_Hyt,
		cuFloatComplex *	d_Hzt,
		float *	d_Prxt,
		float *	d_Pryt,
		float *	d_Przt
	)

Calculate reflected EH and P on target.

Kernel for calculating reflected E, H fields and P, the reflected Poynting vector field, on target surface.

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_A	Array containing area elements.
d_xt	Array containing target points x-coordinate.
d_yt	Array containing target points y-coordinate.
d_zt	Array containing target points z-coordinate.
d_nxs	Array containing source normals x-component.
d_nys	Array containing source normals y-component.
d_nzs	Array containing source normals z-component.
d_nxt	Array containing target norms x-component.
d_nyt	Array containing target norms y-component.
d_nzt	Array containing target norms z-component.
d_Jx	Array containing source J x-component.
d_Jy	Array containing source J y-component.
d_Jz	Array containing source J z-component.
d_Mx	Array containing source M x-component.
d_My	Array containing source M y-component.
d_Mz	Array containing source M z-component.
d_Ext	Array to be filled with target E x-component.
d_Eyt	Array to be filled with target E y-component.
d_Ezt	Array to be filled with target E z-component.
d_Hxt	Array to be filled with target H x-component.
d_Hyt	Array to be filled with target H y-component.
d_Hzt	Array to be filled with target H z-component.
d_Prxt	Array to be filled with target P x-component.
d_Pryt	Array to be filled with target P y-component.
d_Przt	Array to be filled with target P z-component.

GpropagateBeam_4()

void __global__ GpropagateBeam_4	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_nxs,
		float *	d_nys,
		float *	d_nzs,
		float *	d_A,
		float *	d_xt,
		float *	d_yt,
		cuFloatComplex *	d_Jx,
		cuFloatComplex *	d_Jy,
		cuFloatComplex *	d_Jz,
		cuFloatComplex *	d_Mx,
		cuFloatComplex *	d_My,
		cuFloatComplex *	d_Mz,
		cuFloatComplex *	d_Ext,
		cuFloatComplex *	d_Eyt,
		cuFloatComplex *	d_Ezt,
		cuFloatComplex *	d_Hxt,
		cuFloatComplex *	d_Hyt,
		cuFloatComplex *	d_Hzt
	)

Calculate EH on far-field target.

Kernel for calculating E, H fields on a far-field target.

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_A	Array containing area elements.
d_xt	Array containing target direction x-coordinate.
d_yt	Array containing target direction y-coordinate.
d_Jx	Array containing source J x-component.
d_Jy	Array containing source J y-component.
d_Jz	Array containing source J z-component.
d_Mx	Array containing source M x-component.
d_My	Array containing source M y-component.
d_Mz	Array containing source M z-component.
d_Ext	Array to be filled with target E x-component.
d_Eyt	Array to be filled with target E y-component.
d_Ezt	Array to be filled with target E z-component.
d_Hxt	Array to be filled with target H x-component.
d_Hyt	Array to be filled with target H y-component.
d_Hzt	Array to be filled with target H z-component.

GpropagateBeam_5()

void __global__ GpropagateBeam_5	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		float *	d_A,
		float *	d_xt,
		float *	d_yt,
		float *	d_zt,
		cuFloatComplex *	d_sfs,
		cuFloatComplex *	d_sft
	)

Calculate scalar field on target.

Kernel for calculating scalar field on a target, given a scattered field and source surface..

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_A	Array containing area elements.
d_xt	Array containing target direction x-coordinate.
d_yt	Array containing target direction y-coordinate.
d_zt	Array containing target direction z-coordinate.
d_sfs	Array containing source scalar field.
d_sft	Array to be filled with target scalar field.

initCUDA()

__host__ std::array<dim3, 2> initCUDA	(	float	k,
		float	epsilon,
		int	gt,
		int	gs,
		float	t_direction,
		int	nBlocks,
		int	nThreads
	)

Initialize CUDA.

Instantiate program and populate constant memory.

Parameters

k	Wavenumber of incoming field in 1 / mm.
epsilon	Relative electric permittivity of source.
gt	Number of cells on target.
gs	Number of cells on source.
t_direction	Time direction (experimental!).
nBlocks	Number of blocks per grid.
nThreads	Number of threads per block.

Returns

BT Array of two dim3 objects, containing number of blocks per grid and number of threads per block.

scalarfieldAtPoint()

void __device__ scalarfieldAtPoint	(	float *	d_xs,
		float *	d_ys,
		float *	d_zs,
		cuFloatComplex *	d_sfs,
		float(&)	point[3],
		float *	d_A,
		cuFloatComplex &	e
	)

Calculate scalarfield on target.

Kernel for calculating scalarfield on a target.

Parameters

d_xs	Array containing source points x-coordinate.
d_ys	Array containing source points y-coordinate.
d_zs	Array containing source points z-coordinate.
d_sfs	Array containing source scalarfield.
point	Array containing target point.
d_A	Array containing area elements.
e	Array to be filled with results.