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

#include <dudxyz_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__)
 

Detailed Description

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

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

Same three contractions and the same cube layout as dudxyz_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 – only how the six cubes get in and out of shared memory differs.

u arrives as one bulk copy on its own mbarrier and the four factors dr, ds, dt and jacinv as four more on a second, waited on only at the pointwise step, so they are in flight underneath the three contractions. That is 80% of the read traffic moved out from between two barriers, the same share the axhelm variant moves. Five copies issued together is a smaller batch than that kernel's eight but the same shape – every one of the element's reads in one issue – which is what separates the scalar axhelm variant that won from the component-at-a-time vector one that did not.

The result is formed in the staged input, dead once the contractions are done, and leaves as a single bulk store. Nothing is contracted back, so the block is 33.8 kB against the axhelm variant's 45.5 – six resident blocks per SM rather than five.

See dmma_tma_kernel.h for the primitives, the sm_90 and toolkit guards and the lx == DMMA_P bound.

Definition at line 586 of file dudxyz_kernel.h.

Member Function Documentation

◆ run()

template<typename T , const int LX, const int NW>
static __device__ void dudxyz_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__ 
)
inlinestatic

Definition at line 587 of file dudxyz_kernel.h.

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