Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
Static Public Member Functions | List of all members
ax_helm_mfma_vector_dispatch< T, LX, NWF, TILE > Struct Template Reference

#include <ax_helm_kernel.h>

Static Public Member Functions

static __device__ void run (T *, T *, T *, const T *, const T *, const T *, const T *, const T *, const T *, const T *, const T *, const T *, const T *, const T *, const T *, const T *, const int)
 

Detailed Description

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

Matrix-core (MFMA) device kernel for the vector axhelm.

The matrix core counterpart of ax_helm_kernel_vector_kstep, and the HIP counterpart of the CUDA ax_helm_dmma_vector_elem. It runs the three components through the same six contractions as ax_helm_mfma_elem, one component at a time through one set of staged cubes, so the block geometry, the wavefront to element split, the LDS footprint and the tail handling are all the scalar kernel's – the strategy therefore covers the same (precision, LX) set, takes the same wavefronts per block candidates, and one autotuner sweep serves both.

What is genuinely different is where the geometric factors live. The vector operator exists to read them once for the three components rather than once each, and the register file is the obvious place to keep them: seven values for each of the PPT = ceil(LX^3 / (WPE * 64)) points a thread owns. That is cheap where several wavefronts cooperate on an element and ruinous where one wavefront covers a high order element on its own – 189 values, 378 VGPRs, at LX = 12 with one wavefront, against the 256 a lane addresses – so the choice is made per instantiation against a register budget, see NEKO_MFMA_VECTOR_GREG in mfma_kernel.h. Where it does not fit, the factors are re-read from global memory for each component; the element was read moments earlier, so the reads are expected to hit L2.

Note
On CUDA the register resident form is what measured badly: the DMMA vector kernel held ~60 registers of geometry at four warps and lost to its kstep kernel at lx = 8 on GH200, which is what motivated the TMA staged variant that AMD has no counterpart to. Expect the same ranking here, for the same reason as the scalar operator (see the note on ax_helm_kernel_mfma): the operator is far from compute bound, so a matrix core buys nothing that the access pattern does not give back.

Definition at line 1432 of file ax_helm_kernel.h.

Member Function Documentation

◆ run()

template<typename T , const int LX, const int NWF, const int TILE>
static __device__ void ax_helm_mfma_vector_dispatch< T, LX, NWF, TILE >::run ( T *  ,
T *  ,
T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const T *  ,
const int   
)
inlinestatic

Definition at line 1433 of file ax_helm_kernel.h.

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