#include "elem_block.h"
#include "mfma_kernel.h"

Include dependency graph for ax_helm_kernel.h:

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

Classes
struct	ax_helm_mfma_dispatch< T, LX, NWF, TILE >

struct	ax_helm_mfma_vector_dispatch< T, LX, NWF, TILE >

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, const int EB>
__global__ void	__launch_bounds__ ((LX LX EB)) ax_helm_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 NWF, const int TILE>
__global__ void	__launch_bounds__ (64 NWF) ax_helm_kernel_mfma(T __restrict__ w

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

__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	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 = ebLXLX

const int	ele = eLXLX*LX

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

const int	sh_p = ebLX(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 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

◆ __launch_bounds__() [2/2]

template<typename T , const int LX, const int NWF, const int TILE>

__global__ void __launch_bounds__ ( 64 * NWF )

Here is the call graph for this function:

◆ __syncthreads()

__syncthreads ( )

◆ 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 46 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
	)