ax_helm_compr_kernel.h File Reference

#include "elem_block.h"
#include "dmma_kernel.h"
#include "dmma_tma_kernel.h"

Include dependency graph for ax_helm_compr_kernel.h:

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

Go to the source code of this file.

Classes
struct	ax_helm_compr_dmma_dispatch< T, LX, NW >
struct	ax_helm_compr_dmma_tma_dispatch< T, LX, NW >
struct	ax_helm_compr_dmma_vector_dispatch< T, LX, NW >
struct	ax_helm_compr_dmma_tma_vector_dispatch< T, LX, NW >
struct	ax_helm_compr_dmma_tma_batch_dispatch< T, LX, NW >

Functions
template<typename T , const int LX, const int CHUNKS>
__global__ void	ax_helm_compr_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, const int *__restrict__ ind)
template<typename T , const int LX, const int EB>
__global__ void	__launch_bounds__ ((LX *LX *EB)) ax_helm_compr_kernel_kstep(T *__restrict__ w
	if (eb==0)
	for (int k=0;k< LX;++k)
	__syncthreads ()
	if (active)
template<typename T , const int LX, const int NW>
__global__ void	__launch_bounds__ ((32 *NW)) ax_helm_compr_kernel_dmma(T *__restrict__ w
template<typename T , const int LX, const int NW>
static bool	ax_helm_compr_dmma_tma_batch_optin ()

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
__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__ const int *__restrict__	ind
__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__ const int *__restrict__ const int	nelv
__shared__ T	shdy [LX *LX]
__shared__ T	shdz [LX *LX]
__shared__ T	shu [EB *LX *LX]
__shared__ T	shur [EB *LX *LX]
__shared__ T	shus [EB *LX *LX]
T	ru [LX]
T	rw [LX]
T	rut
const int	eb = (EB == 1) ? 0 : threadIdx.z
const int	e_blk = blockIdx.x * EB + eb
const bool	active = (EB == 1) ? true : (e_blk < nelv)
const int	e = active ? e_blk : (nelv - 1)
const int	j = threadIdx.y
const int	i = threadIdx.x
const int	ij = i + j*LX
const int	sh = eb*LX*LX
const int	ele = e*LX*LX*LX
const int	gele = (ind[e] - 1)*LX*LX*LX
const int	ij_p = i + j*(LX+1)
const int	sh_p = eb*LX*(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 [EB *LX *LX]
__shared__ T	shvr [EB *LX *LX]
__shared__ T	shvs [EB *LX *LX]
__shared__ T	shw [EB *LX *LX]
__shared__ T	shwr [EB *LX *LX]
__shared__ T	shws [EB *LX *LX]
T	rv [LX]
T	ruw [LX]
T	rvw [LX]
T	rww [LX]
T	rvt
T	rwt

Function Documentation

__launch_bounds__() [1/2]

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

__global__ void __launch_bounds__

(

(32 *NW)

)

Here is the call graph for this function:

__launch_bounds__() [2/2]

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

__global__ void __launch_bounds__

(

(LX *LX *EB)

)

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_compr_dmma_tma_batch_optin()

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

static bool ax_helm_compr_dmma_tma_batch_optin ( )

inlinestatic

Definition at line 2289 of file ax_helm_compr_kernel.h.

Here is the call graph for this function:

ax_helm_compr_kernel_1d()

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

__global__ void ax_helm_compr_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,
		const int *__restrict__	ind
	)

Device kernel for axhelm

Definition at line 88 of file ax_helm_compr_kernel.h.

Here is the call graph for this function:

for()

for

(

)

Definition at line 259 of file ax_helm_compr_kernel.h.

Here is the call graph for this function:

if() [1/2]

if

(

active

)

Definition at line 315 of file ax_helm_compr_kernel.h.

Here is the call graph for this function:

if() [2/2]

if

(

eb

= = 0

)

Definition at line 252 of file ax_helm_compr_kernel.h.

Variable Documentation

active

const bool active = (EB == 1) ? true : (e_blk < nelv)

Definition at line 243 of file ax_helm_compr_kernel.h.

av

__global__ void T *__restrict__ av

Definition at line 974 of file ax_helm_compr_kernel.h.

aw

__global__ void T *__restrict__ T *__restrict__ aw

Definition at line 975 of file ax_helm_compr_kernel.h.

dx

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

Definition at line 200 of file ax_helm_compr_kernel.h.

dy

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

Definition at line 201 of file ax_helm_compr_kernel.h.

dz

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

Definition at line 202 of file ax_helm_compr_kernel.h.

e

const int e = active ? e_blk : (nelv - 1)

Definition at line 244 of file ax_helm_compr_kernel.h.

e_blk

const int e_blk = blockIdx.x * EB + eb

Definition at line 242 of file ax_helm_compr_kernel.h.

eb

const int eb = (EB == 1) ? 0 : threadIdx.z

Definition at line 241 of file ax_helm_compr_kernel.h.

ele

const int ele = e*LX*LX*LX

Definition at line 249 of file ax_helm_compr_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 204 of file ax_helm_compr_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 207 of file ax_helm_compr_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 208 of file ax_helm_compr_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 205 of file ax_helm_compr_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:

{
 
  ax_helm_dmma_tma_dispatch< T, LX, NW >::run(w, u, dx, dy, dz, h1,
                                              g11, g22, g33, g12, g13, g23)

Definition at line 209 of file ax_helm_compr_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 206 of file ax_helm_compr_kernel.h.

gele

const int gele = (ind[e] - 1)*LX*LX*LX

Definition at line 250 of file ax_helm_compr_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 203 of file ax_helm_compr_kernel.h.

i

const int i = threadIdx.x

Definition at line 246 of file ax_helm_compr_kernel.h.

ij

const int ij = i + j*LX

Definition at line 247 of file ax_helm_compr_kernel.h.

ij_p

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

Definition at line 377 of file ax_helm_compr_kernel.h.

ind

__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__ const int *__restrict__ ind

Initial value:

{
 
  ax_helm_compr_dmma_tma_dispatch< T, LX, NW >::run(w, u, dx, dy, dz, h1, g11,
                                                    g22, g33, g12, g13, g23,
                                                    ind)

Definition at line 210 of file ax_helm_compr_kernel.h.

j

const int j = threadIdx.y

Definition at line 245 of file ax_helm_compr_kernel.h.

nelv

__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__ const T *__restrict__ const int nelv

Initial value:

{
 
  
  __shared__ T shdx[LX * LX]

Definition at line 211 of file ax_helm_compr_kernel.h.

ru

T ru

Definition at line 232 of file ax_helm_compr_kernel.h.

rut

T rut

Definition at line 234 of file ax_helm_compr_kernel.h.

ruw

T ruw

Definition at line 1029 of file ax_helm_compr_kernel.h.

rv

T rv

Definition at line 1026 of file ax_helm_compr_kernel.h.

rvt

T rvt

Definition at line 1034 of file ax_helm_compr_kernel.h.

rvw

T rvw

Definition at line 1030 of file ax_helm_compr_kernel.h.

rw

T rw

Definition at line 233 of file ax_helm_compr_kernel.h.

rwt

T rwt

Definition at line 1035 of file ax_helm_compr_kernel.h.

rww

T rww

Definition at line 1031 of file ax_helm_compr_kernel.h.

sh

const int sh = eb*LX*LX

Definition at line 248 of file ax_helm_compr_kernel.h.

sh_p

const int sh_p = eb*LX*(LX+1)

Definition at line 379 of file ax_helm_compr_kernel.h.

shdy

__shared__ T shdy

Definition at line 215 of file ax_helm_compr_kernel.h.

shdz

__shared__ T shdz

Definition at line 216 of file ax_helm_compr_kernel.h.

shu

__shared__ T shu

Definition at line 219 of file ax_helm_compr_kernel.h.

shur

__shared__ T shur

Definition at line 220 of file ax_helm_compr_kernel.h.

shus

__shared__ T shus

Definition at line 221 of file ax_helm_compr_kernel.h.

shv

__shared__ T shv

Definition at line 1002 of file ax_helm_compr_kernel.h.

shvr

__shared__ T shvr

Definition at line 1003 of file ax_helm_compr_kernel.h.

shvs

__shared__ T shvs

Definition at line 1004 of file ax_helm_compr_kernel.h.

shw

__shared__ T shw

Definition at line 1006 of file ax_helm_compr_kernel.h.

shwr

__shared__ T shwr

Definition at line 1007 of file ax_helm_compr_kernel.h.

shws

__shared__ T shws

Definition at line 1008 of file ax_helm_compr_kernel.h.

u

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

Definition at line 199 of file ax_helm_compr_kernel.h.

v

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

Definition at line 977 of file ax_helm_compr_kernel.h.

w

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

Definition at line 978 of file ax_helm_compr_kernel.h.