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

Detailed Description

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

Device kernel for axhelm on the fp64 tensor cores

One element per block, NW warps, and the whole element resident in shared memory as four DMMA_P^3 cubes: the staged input (reused as the output), and the three reference derivatives. Unlike the kstep variants, which stream a k plane at a time and keep the k contraction in registers, all six contractions here are full D * U GEMMs handed to dmma_contract(), with the geometric factors applied pointwise in between. See dmma_kernel.h for the padded staging, the per axis matrix views and the arch and LX bounds.

Definition at line 605 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_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 int   
)
inlinestatic

Definition at line 606 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: