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

#include <conv1_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__, 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 conv1_dmma_tma_dispatch< T, LX, NW >

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

Same three contractions and the same cube layout as conv1_dmma_elem(); at lx == DMMA_P the padding is empty and the cube offset is the flat point index, so the arithmetic is identical and only the staging differs.

All fourteen input copies are issued together at kernel entry – u on its own mbarrier and waited immediately, the other thirteen on a second and waited only at the pointwise step – so 84% of the read traffic is in flight underneath the contractions. Batching every read into one issue is the arrangement the scalar axhelm TMA kernel won with. The single output leaves as one bulk store through the staged input, which is dead by then.

Seventeen cubes is 71184 B, so the allocation is dynamic and its layout is conv1_tma_smem in dmma_tma_kernel.h. That is three resident blocks per SM rather than the four the batched axhelm and the opgrad variants get, which is the reason this one is a real question rather than an expected win: it spends a block to buy the largest batch in the tree, against the highest bytes per contraction in the tree. Which way that lands is a measurement, and it is a tuner candidate like every other one.

Definition at line 642 of file conv1_kernel.h.

Member Function Documentation

◆ run()

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

Definition at line 643 of file conv1_kernel.h.

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