#include <opgrad_kernel.h>

Static Public Member Functions
static __device__ void	run (T , T , T , const T , const T , const T , const T , const T , const T , const T , const T , const T , const T , const T , const T , const T , const T *, const int)

Detailed Description

template<typename T, const int LX, const int NWF>
struct opgrad_mfma_dispatch< T, LX, NWF >

Device kernel for the weak gradient on the AMD matrix cores

The same staging and the same three contractions as dudxyz_mfma_elem(), with a wider tail: three outputs formed pointwise from nine metric cubes instead of one from three. See mfma_kernel.h for the tiles, the lane layouts, the wavefronts-per-element scheme and the arch and LX bounds.

w3 is a quadrature weight shared by every element, so it is indexed by the point within the element rather than by the global offset. Unlike the CUDA counterpart there is no padded cube here, so that index is simply p.

Definition at line 355 of file opgrad_kernel.h.

Member Function Documentation

◆ run()

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

static __device__ void opgrad_mfma_dispatch< T, LX, NWF >::run	(	T *	,
		T *	,
		T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const T *	,
		const int
	)

inlinestatic

Definition at line 356 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/hip/opgrad_kernel.h

Static Public Member Functions

Detailed Description

Member Function Documentation

◆ run()