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_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__, const int)
 

Detailed Description

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

Device kernel for the convective term on the fp64 tensor cores

Structurally the same as opgrad_dmma_elem(): the same four DMMA_P^3 cubes, the same stage-u-then-contract-three-axes, and a pointwise tail. Only u is contracted – the convecting velocity and the metrics enter pointwise – so three contractions serve thirteen streamed factor cubes.

That gives conv1 20480 bytes per contraction at lx = 8, the highest of any operator in the tree and 3.3x axhelm's 6144. Since the staging tax is paid per contraction and the memory time hiding it is paid per byte, this is the most favourable of that trade anywhere, which is the whole reason the strategy is offered here. See dmma_kernel.h for the padded staging, the per axis matrix views and the arch and LX bounds.

Definition at line 402 of file conv1_kernel.h.

Member Function Documentation

◆ run()

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

Definition at line 403 of file conv1_kernel.h.

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