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

Include dependency graph for ax_helm_compr_kernel.h:

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

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

struct	ax_helm_compr_mfma_vector_dispatch< T, LX, NWF, TILE >

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

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

const int	ele = eLXLX*LX

const int	gele = (ind[e] - 1)LXLX*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_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 kernels for Ax helm on compressed geometric factors

Derived mechanically from ax_helm_kernel.h, and a change to one belongs in the other. Every strategy there is here – 1d at every chunk candidate, kstep and kstep_padded at every elements per block candidate, the matrix core kernels at every wavefront and tile candidate, scalar and vector – with one difference: element e reads its geometric factors g11..g23 from copy ind[e] - 1, ind holding the 1-based copy indices of coef_t::compression_inds. u, h1 and the outputs stay indexed by e, and in the blocked and matrix core kernels each element of the block looks up its own copy.

Definition at line 56 of file ax_helm_compr_kernel.h.

Here is the call graph for this function: