#include <ax_helm_kernel.h>

Static Public Member Functions
static __device__ void	run (T __restrict__, 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__)

Detailed Description

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

Device kernel for the vector axhelm on the fp64 tensor cores

The three components share one set of geometric factors, which is the whole point of the vector operator: reading them once and applying them to u, v and w moves 13 arrays per element instead of the 27 that three scalar calls would. Twelve staged cubes will not fit in a block's shared memory, so the components are run one after another through the same four cubes and the factors are hoisted into registers instead – PPT points per thread, held across all three passes. Everything else is the scalar kernel; see dmma_kernel.h for the padded staging and the per axis matrix views.

Note the padded entries of a staged cube only have to be finite, not zero. A padded row of the derivative matrix is zero and kills the contribution of a padded row of the cube, but 0 * NaN would not, so the cube is zeroed once before its first use and then left to carry whatever the contractions put there. That is why the component loop below re-stages only the LX^3 real points and never clears the padding again.

Measured on GH200 at lx = 8 this loses to the kstep variant: holding the factors costs enough occupancy (~60 registers of metrics at nw = 4, against the scalar kernel's handful) to give back more than the access pattern wins. It is kept for the low order end, where PPT falls to 1 and the register cost with it, while the shared memory footprint stays flat. The autotuner decides.

Definition at line 1572 of file ax_helm_kernel.h.

Member Function Documentation

◆ run()

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

static __device__ void ax_helm_dmma_vector_dispatch< T, LX, NW >::run	(	T *	__restrict__,
		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__
	)

inlinestatic

Definition at line 1573 of file ax_helm_kernel.h.

The documentation for this struct was generated from the following file:

/home/runner/work/neko/neko/src/math/bcknd/device/cuda/ax_helm_kernel.h

Static Public Member Functions

Detailed Description

Member Function Documentation

◆ run()