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_dmma_tma_dispatch< T, LX, NW > Struct Template Reference

#include <ax_helm_kernel.h>

Static Public Member Functions

static __device__ void run (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_tma_dispatch< T, LX, NW >

Device kernel for axhelm on the fp64 tensor cores, with the element staged by the TMA engine

Same six contractions and the same padded cube layout as ax_helm_dmma_elem(), and at lx == DMMA_P the padding is empty and the cube offset is the flat point index, so the arithmetic is identical – the only thing that differs is how the eleven cubes get in and out of shared memory.

Here u arrives as one bulk copy on its own mbarrier and the seven geometric factors as seven more on a second, and only the first is waited on before the contractions start. The factors land in shared memory across the three contractions and are waited on at the pointwise step that consumes them, which is 78% of the read traffic moved out from between two barriers and put underneath the tensor cores. The result leaves as a single bulk store.

See dmma_tma_kernel.h for the primitives, the sm_90 and toolkit guards, the lx == DMMA_P bound and the occupancy trade this buys the overlap with.

Definition at line 863 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_tma_dispatch< T, LX, NW >::run ( 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 864 of file ax_helm_kernel.h.

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