#include <opgrad_kernel.h>

Static Public Member Functions
static __device__ void	run (T __restrict__, T __restrict__, 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 opgrad_dmma_dispatch< T, LX, NW >

Device kernel for the weak gradient on the fp64 tensor cores

Structurally this is dudxyz_dmma_elem() with a wider tail: the same four DMMA_P^3 cubes, the same stage-u-then-contract-three-axes, and then three outputs formed pointwise from nine streamed metric cubes instead of one from three. See dmma_kernel.h for the padded staging, the per axis matrix views and the arch and LX bounds.

This is the best placed scalar operator in the tree for the strategy. What won axhelm was the access pattern rather than the tensor cores – 82% of achievable bandwidth at 7.8% of fp64 tensor peak – and the cost of the staging is paid per contraction while the memory time hiding it is paid per byte. At lx = 8 opgrad carries 17749 bytes per contraction against axhelm's 6144, so there is 2.9x as much memory time here to hide the same staging behind. It is also in the default pnpn path, three calls per timestep in the dealiased advection plus the pressure gradient, where dudxyz is not.

w3 is the quadrature weight and is shared by every element, so it is indexed by the point within the element rather than the global offset – dmma_idx::l rather than ::g. That is the only thing here the derivative kernel does not already do.

Definition at line 396 of file opgrad_kernel.h.

Member Function Documentation

◆ run()

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

static __device__ void opgrad_dmma_dispatch< T, LX, NW >::run	(	T *	__restrict__,
		T *	__restrict__,
		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 397 of file opgrad_kernel.h.

Here is the caller graph for this function:

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

/home/runner/work/neko/neko/src/math/bcknd/device/cuda/opgrad_kernel.h

Static Public Member Functions

Detailed Description

Member Function Documentation

◆ run()