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 > 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>
struct ax_helm_mfma_dispatch< T, LX, NWF >

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.

Definition at line 563 of file ax_helm_kernel.h.

Member Function Documentation

◆ run()

template<typename T , const int LX, const int NWF>
static __device__ void ax_helm_mfma_dispatch< T, LX, NWF >::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 564 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: