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_dispatch< T, LX, NWF, TILE > Struct Template Reference

#include <ax_helm_kernel.h>

Static Public Member Functions

static __device__ void run (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_dispatch< T, LX, NWF, TILE >

Matrix-core (MFMA) device kernel for axhelm.

Additional autotuner strategy that maps the six spectral-element tensor contractions (3 gradient + 3 divergence) onto the AMD matrix cores via the precision-dispatched mfma_contract_sel primitive – double precision uses the batched 4x4x4 matrix core (full M-utilisation), single precision the 16x16x4 tile (see mfma_kernel.h for the precision traits, lane layout and supported (precision, LX) set). AX_HELM_MFMA_NWF cooperating wavefronts process one element, sharing the LX^3 field staged in LDS; decoupling the block shape from LX lets one implementation cover 4 <= LX <= 12 rather than only LX = 8. Unsupported (T, LX) instantiate to a no-op; the autotuner only launches this kernel for the supported set (see mfma_lx_supported() and hip_have_mfma() in mfma_kernel.h), so the no-op is never reached at runtime.

Note
Do not expect it to win. MEASURED on gfx90a in single precision at 8192 elements per rank: at lx = 4 it is chosen, 20.48 us/call against the 1d kernel's 20.87, a 1.9% margin; at lx = 8 it loses, 136.40 against 130.50 for 1d at 512 threads, 4.5% behind. Those two numbers are 1.107 and 1.157 TB/s of a 1.6 TB/s MI250X GCD, so both kernels are sitting on the memory roof and what separates them is streaming efficiency, not arithmetic. That is the whole difficulty with the strategy: the operator runs at roughly 1.6 flop/byte, tens of times below the ridge, so the tile utilisation a matrix core buys has to be paid for in staging traffic and usually is. Nothing has been measured on gfx942 at all.

ALL OF THOSE NUMBERS ARE STALE. They predate the elements per block rework (which changed what more than one wavefront means at lx = 4), the padded LDS layout, the B-hoisted 4x4x4 nest, the 16 wavefront rung and the geometry prefetch. Of those the first three are the ones expected to move the lx = 8 comparison: the bank model prices a whole lx = 8 f64 element – every LDS access of the six contractions and the three linear passes – at 2976 cycles before them and 1220 after, and the streaming efficiency those numbers are really measuring is what an 8-way LDS conflict costs. The same model puts lx = 12 f64 at 8964 -> 4752 and the orders it does not pad, lx = 10 and 11, at -30% and -27% from the hoisted operand alone. Re-measure before quoting any of it, and treat the 4.5% deficit as the thing to re-test rather than as a settled result.

Definition at line 674 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_dispatch< T, LX, NWF, TILE >::run ( 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 675 of file ax_helm_kernel.h.

Here is the caller graph for this function:
The documentation for this struct was generated from the following file: