ax_helm_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	ax_helm_kernel_1d (T *__restrict__ w, const T *__restrict__ u, const T *__restrict__ dx, const T *__restrict__ dy, const T *__restrict__ dz, const T *__restrict__ dxt, const T *__restrict__ dyt, const T *__restrict__ dzt, const T *__restrict__ h1, const T *__restrict__ g11, const T *__restrict__ g22, const T *__restrict__ g33, const T *__restrict__ g12, const T *__restrict__ g13, const T *__restrict__ g23)
template<typename T , const int LX>
__global__ void	__launch_bounds__ (LX *LX, 3) ax_helm_kernel_kstep(T *__restrict__ w
	for (int k=0;k< LX;++k)
	__syncthreads ()
template<typename T >
__global__ void	ax_helm_kernel_vector_part2 (T *__restrict__ au, T *__restrict__ av, T *__restrict__ aw, const T *__restrict__ u, const T *__restrict__ v, const T *__restrict__ w, const T *__restrict__ h2, const T *__restrict__ B, const int n)

Variables
__global__ void const T *__restrict__	u
__global__ void const T *__restrict__ const T *__restrict__	dx
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__	dy
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	dz
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	h1
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	g11
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	g22
__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__	g33
__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__	g12
__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__ const T *__restrict__	g13
__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__ const T *__restrict__ const T *__restrict__	g23
__shared__ T	shdy [LX *LX] = dy[ij]
__shared__ T	shdz [LX *LX] = dz[ij]
__shared__ T	shu [LX *LX]
__shared__ T	shur [LX *LX]
__shared__ T	shus [LX *LX]
T	ru [LX]
T	rw [LX]
T	rut
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
	shdx [ij] = dx[ij]
const int	ij_p = i + j*(LX+1)
__global__ void T *__restrict__	av
__global__ void T *__restrict__ T *__restrict__	aw
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__	v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__	w
__shared__ T	shv [LX *LX]
__shared__ T	shvr [LX *LX]
__shared__ T	shvs [LX *LX]
__shared__ T	shw [LX *LX]
__shared__ T	shwr [LX *LX]
__shared__ T	shws [LX *LX]
T	rv [LX]
T	ruw [LX]
T	rvw [LX]
T	rww [LX]
T	rvt
T	rwt

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:

ax_helm_kernel_1d()

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

__global__ void ax_helm_kernel_1d	(	T *__restrict__	w,
		const T *__restrict__	u,
		const T *__restrict__	dx,
		const T *__restrict__	dy,
		const T *__restrict__	dz,
		const T *__restrict__	dxt,
		const T *__restrict__	dyt,
		const T *__restrict__	dzt,
		const T *__restrict__	h1,
		const T *__restrict__	g11,
		const T *__restrict__	g22,
		const T *__restrict__	g33,
		const T *__restrict__	g12,
		const T *__restrict__	g13,
		const T *__restrict__	g23
	)

Device kernels for Ax helm

Definition at line 42 of file ax_helm_kernel.h.

Here is the call graph for this function:

ax_helm_kernel_vector_part2()

template<typename T >

__global__ void ax_helm_kernel_vector_part2	(	T *__restrict__	au,
		T *__restrict__	av,
		T *__restrict__	aw,
		const T *__restrict__	u,
		const T *__restrict__	v,
		const T *__restrict__	w,
		const T *__restrict__	h2,
		const T *__restrict__	B,
		const int	n
	)

Definition at line 726 of file ax_helm_kernel.h.

for()

for

(

)

Definition at line 186 of file ax_helm_kernel.h.

Variable Documentation

av

__global__ void T *__restrict__ av

Definition at line 359 of file ax_helm_kernel.h.

aw

__global__ void T *__restrict__ T *__restrict__ aw

Definition at line 360 of file ax_helm_kernel.h.

dx

__global__ void T* __restrict__ T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dx

Definition at line 152 of file ax_helm_kernel.h.

dy

__global__ void T* __restrict__ T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dy

Definition at line 153 of file ax_helm_kernel.h.

dz

__global__ void T* __restrict__ T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dz

Definition at line 154 of file ax_helm_kernel.h.

e

const int e = blockIdx.x

Definition at line 175 of file ax_helm_kernel.h.

ele

const int ele = e*LX*LX*LX

Definition at line 179 of file ax_helm_kernel.h.

g11

__global__ void T *__restrict__ 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__ g11

Definition at line 156 of file ax_helm_kernel.h.

g12

__global__ void T *__restrict__ 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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g12

Definition at line 159 of file ax_helm_kernel.h.

g13

__global__ void T *__restrict__ 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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g13

Definition at line 160 of file ax_helm_kernel.h.

g22

__global__ void T *__restrict__ 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__ const T *__restrict__ g22

Definition at line 157 of file ax_helm_kernel.h.

g23

__global__ void T *__restrict__ 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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g23

Initial value:

{
 
  __shared__ T shdx[LX * LX]

Definition at line 161 of file ax_helm_kernel.h.

g33

__global__ void T *__restrict__ 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__ const T *__restrict__ const T *__restrict__ g33

Definition at line 158 of file ax_helm_kernel.h.

h1

__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ h1

Definition at line 155 of file ax_helm_kernel.h.

i

const int i = threadIdx.x

Definition at line 177 of file ax_helm_kernel.h.

ij

const int ij = i + j*LX

Definition at line 178 of file ax_helm_kernel.h.

ij_p

const int ij_p = i + j*(LX+1)

Definition at line 283 of file ax_helm_kernel.h.

j

const int j = threadIdx.y

Definition at line 176 of file ax_helm_kernel.h.

ru

T ru[LX]

Definition at line 171 of file ax_helm_kernel.h.

rut

T rut

Definition at line 173 of file ax_helm_kernel.h.

ruw

T ruw

Definition at line 395 of file ax_helm_kernel.h.

rv

T rv[LX]

Definition at line 392 of file ax_helm_kernel.h.

rvt

T rvt

Definition at line 400 of file ax_helm_kernel.h.

rvw

T rvw

Definition at line 396 of file ax_helm_kernel.h.

rw

T rw[LX]

Definition at line 172 of file ax_helm_kernel.h.

rwt

T rwt

Definition at line 401 of file ax_helm_kernel.h.

rww

T rww

Definition at line 397 of file ax_helm_kernel.h.

shdx

shdx[ij_p] = dx[ij]

Definition at line 181 of file ax_helm_kernel.h.

shdy

shdy[ij_p] = dy[ij]

Definition at line 164 of file ax_helm_kernel.h.

shdz

shdz[ij_p] = dz[ij]

Definition at line 165 of file ax_helm_kernel.h.

shu

__shared__ T shu

Definition at line 167 of file ax_helm_kernel.h.

shur

__shared__ T shur

Definition at line 168 of file ax_helm_kernel.h.

shus

__shared__ T shus

Definition at line 169 of file ax_helm_kernel.h.

shv

__shared__ T shv[LX *(LX+1)]

Definition at line 383 of file ax_helm_kernel.h.

shvr

__shared__ T shvr

Definition at line 384 of file ax_helm_kernel.h.

shvs

__shared__ T shvs

Definition at line 385 of file ax_helm_kernel.h.

shw

__shared__ T shw[LX *(LX+1)]

Definition at line 387 of file ax_helm_kernel.h.

shwr

__shared__ T shwr

Definition at line 388 of file ax_helm_kernel.h.

shws

__shared__ T shws

Definition at line 389 of file ax_helm_kernel.h.

u

__global__ void T* __restrict__ T* __restrict__ const T* __restrict__ u

Definition at line 151 of file ax_helm_kernel.h.

v

__global__ void T* __restrict__ T* __restrict__ const T* __restrict__ const T* __restrict__ v

Definition at line 362 of file ax_helm_kernel.h.

w

__global__ void T* __restrict__ T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ w

Definition at line 363 of file ax_helm_kernel.h.