cdtp_kernel.h File Reference

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

Go to the source code of this file.

Functions
template<typename T , const int LX, const int CHUNKS>
__global__ void	cdtp_kernel_1d (T *__restrict__ dtx, const T *__restrict__ x, const T *__restrict__ dr, const T *__restrict__ ds, const T *__restrict__ dt, const T *__restrict__ dxt, const T *__restrict__ dyt, const T *__restrict__ dzt, const T *__restrict__ B, const T *__restrict__ jac)
template<typename T , const int LX>
__global__ void	__launch_bounds__ (LX *LX, 3) cdtp_kernel_kstep(T *__restrict__ dtx
	for (int k=0;k< LX;++k)
	__syncthreads ()

Variables
__global__ void const T *__restrict__	x
__global__ void const T *__restrict__ const T *__restrict__	dr
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__	ds
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	dt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	dxt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	dyt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	dzt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	B
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	jac
__shared__ T	shdyt [LX *LX] = dyt[ij]
__shared__ T	shdzt [LX *LX] = dzt[ij]
__shared__ T	shtar [LX *LX]
__shared__ T	shtas [LX *LX]
T	rtar [LX]
T	rtas [LX]
T	rtat [LX]
const int	e = blockIdx.x
const int	j = threadIdx.y
const int	i = threadIdx.x
const int	ij = i + j * LX
const int	ele = e*LX*LX*LX
	shdxt [ij] = dxt[ij]

Function Documentation

__launch_bounds__()

template<typename T , const int LX>

__global__ void __launch_bounds__	(	LX *	LX,
		3
	)

Device kernel for axhelm with padding in shared memory to remove bank conflicts when LX is a power of 2

__syncthreads()

__syncthreads

(

)

Here is the caller graph for this function:

cdtp_kernel_1d()

template<typename T , const int LX, const int CHUNKS>

__global__ void cdtp_kernel_1d	(	T *__restrict__	dtx,
		const T *__restrict__	x,
		const T *__restrict__	dr,
		const T *__restrict__	ds,
		const T *__restrict__	dt,
		const T *__restrict__	dxt,
		const T *__restrict__	dyt,
		const T *__restrict__	dzt,
		const T *__restrict__	B,
		const T *__restrict__	jac
	)

Device kernel for \( D^T x \)

Definition at line 42 of file cdtp_kernel.h.

Here is the call graph for this function:

for()

for

(

)

Definition at line 142 of file cdtp_kernel.h.

Variable Documentation

B

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ B

Definition at line 116 of file cdtp_kernel.h.

dr

__global__ void const T* __restrict__ const T* __restrict__ dr

Definition at line 110 of file cdtp_kernel.h.

ds

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ ds

Definition at line 111 of file cdtp_kernel.h.

dt

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dt

Definition at line 112 of file cdtp_kernel.h.

dxt

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dxt

Definition at line 113 of file cdtp_kernel.h.

dyt

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dyt

Definition at line 114 of file cdtp_kernel.h.

dzt

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dzt

Definition at line 115 of file cdtp_kernel.h.

e

const int e = blockIdx.x

Definition at line 130 of file cdtp_kernel.h.

ele

const int ele = e*LX*LX*LX

Definition at line 134 of file cdtp_kernel.h.

i

const int i = threadIdx.x

Definition at line 132 of file cdtp_kernel.h.

ij

const int ij = i + j * LX

Definition at line 133 of file cdtp_kernel.h.

j

const int j = threadIdx.y

Definition at line 131 of file cdtp_kernel.h.

jac

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ jac

Initial value:

{ 
  
  __shared__ T shdxt[LX * LX]

Definition at line 117 of file cdtp_kernel.h.

rtar

T rtar[LX]

Definition at line 126 of file cdtp_kernel.h.

rtas

T rtas[LX]

Definition at line 127 of file cdtp_kernel.h.

rtat

T rtat[LX]

Definition at line 128 of file cdtp_kernel.h.

shdxt

shdxt[ij] = dxt[ij]

Definition at line 136 of file cdtp_kernel.h.

shdyt

shdyt[ij] = dyt[ij]

Definition at line 120 of file cdtp_kernel.h.

shdzt

shdzt[ij] = dzt[ij]

Definition at line 121 of file cdtp_kernel.h.

shtar

__shared__ T shtar[LX *LX]

Definition at line 123 of file cdtp_kernel.h.

shtas

__shared__ T shtas[LX *LX]

Definition at line 124 of file cdtp_kernel.h.

x

__global__ void const T* __restrict__ x

Definition at line 109 of file cdtp_kernel.h.