#include <stdlib.h>
#include <string.h>
#include <hip/hip_runtime.h>
#include <device/device_config.h>
#include <device/hip/check.h>
#include "elem_block.h"

Include dependency graph for mfma_kernel.h:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Macros
#define	NEKO_MFMA_NWF_CANDIDATES 5

#define	NEKO_MFMA_TILE_CANDIDATES 2

#define	NEKO_MFMA_CANDIDATES (NEKO_MFMA_NWF_CANDIDATES * NEKO_MFMA_TILE_CANDIDATES)

#define	NEKO_MFMA_NWF(C) (1 << ((C) % NEKO_MFMA_NWF_CANDIDATES))

#define	NEKO_MFMA_TILE(C) ((C) / NEKO_MFMA_NWF_CANDIDATES)

#define	NEKO_MFMA_NTHRDS(C) dim3(64, NEKO_MFMA_NWF(C), 1)

#define	NEKO_MFMA_PAD_N(LX, SZ) ((((LX) % 4 == 0) && !((LX) == 4 && (SZ) == 4)) ? 1 : 0)

#define	NEKO_MFMA_SJ_N(LX, SZ) ((LX) + NEKO_MFMA_PAD_N(LX, SZ))

#define	NEKO_MFMA_SK_N(LX, SZ) (NEKO_MFMA_SJ_N(LX, SZ) * (LX))

#define	NEKO_MFMA_CUBE_N(LX, SZ) (NEKO_MFMA_SK_N(LX, SZ) * (LX))

#define	NEKO_MFMA_SD_N(LX, SZ) NEKO_MFMA_SJ_N(LX, SZ)

#define	NEKO_MFMA_DMAT_N(LX, SZ) (NEKO_MFMA_SD_N(LX, SZ) * (LX))

#define	NEKO_MFMA_NGROUPS(LX) (((LX) * (LX) + 15) / 16)

#define	NEKO_MFMA_LDS_N(EB, LX, SZ) ((3 * NEKO_MFMA_DMAT_N(LX, SZ) + 4 * (EB) * NEKO_MFMA_CUBE_N(LX, SZ)) * (SZ))

#define	NEKO_MFMA_LDS_FITS(EB, LX, SZ) (NEKO_MFMA_LDS_N(EB, LX, SZ) <= NEKO_EB_MAX_LDS)

#define	NEKO_MFMA_EB_N(NWF, LX, SZ)

#define	NEKO_MFMA_EB(LX, C) NEKO_MFMA_EB_N(NEKO_MFMA_NWF(C), LX, sizeof(real))

#define	NEKO_MFMA_WPE(LX, C) (NEKO_MFMA_NWF(C) / NEKO_MFMA_EB(LX, C))

#define	NEKO_MFMA_SPT_N(WPE, LX, SZ) ((NEKO_MFMA_CUBE_N(LX, SZ) + (WPE) * 64 - 1) / ((WPE) * 64))

#define	NEKO_MFMA_SPT(LX, C) NEKO_MFMA_SPT_N(NEKO_MFMA_WPE(LX, C), LX, sizeof(real))

#define	NEKO_MFMA_VECTOR_GREG_VGPRS 64

#define	NEKO_MFMA_VECTOR_GREG_N(SPT, SZ) ((7 * (SPT) * ((SZ) / 4)) <= NEKO_MFMA_VECTOR_GREG_VGPRS)

#define	NEKO_MFMA_VECTOR_GREG(LX, C) NEKO_MFMA_VECTOR_GREG_N(NEKO_MFMA_SPT(LX, C), sizeof(real))

#define	NEKO_MFMA_NBLCKS(NELV, LX, C) dim3(((NELV) + NEKO_MFMA_EB(LX, C) - 1) / NEKO_MFMA_EB(LX, C), 1, 1)

#define	NEKO_TUNE_LOG_MFMA(LX, T3)

#define	NEKO_TUNE_LOG_MFMA_VEC(LX, T3)

Functions
static __global__ void	hip_mfma_arch_probe (int *flag)

static bool	hip_have_mfma ()

template<const int LX>
static bool	mfma_lx_supported ()

static bool	mfma_tile_offered (const int c)

static const char *	mfma_tile_name (const int c)

static int	neko_mfma_sweep ()

static int	neko_mfma_pin ()

static int	neko_mfma_nwf_pin ()

static int	neko_mfma_candidates ()

Macro Definition Documentation

◆ NEKO_MFMA_CANDIDATES

#define NEKO_MFMA_CANDIDATES (NEKO_MFMA_NWF_CANDIDATES * NEKO_MFMA_TILE_CANDIDATES)

Definition at line 261 of file mfma_kernel.h.

◆ NEKO_MFMA_CUBE_N

#define NEKO_MFMA_CUBE_N	(	LX,
		SZ
	)	(NEKO_MFMA_SK_N(LX, SZ) * (LX))

Definition at line 330 of file mfma_kernel.h.

◆ NEKO_MFMA_DMAT_N

#define NEKO_MFMA_DMAT_N	(	LX,
		SZ
	)	(NEKO_MFMA_SD_N(LX, SZ) * (LX))

Definition at line 351 of file mfma_kernel.h.

◆ NEKO_MFMA_EB

#define NEKO_MFMA_EB	(	LX,
		C
	)	NEKO_MFMA_EB_N(NEKO_MFMA_NWF(C), LX, sizeof(real))

Definition at line 442 of file mfma_kernel.h.

◆ NEKO_MFMA_EB_N

#define NEKO_MFMA_EB_N	(	NWF,
		LX,
		SZ
	)

Value:

  ((NWF) / NEKO_MFMA_NGROUPS(LX) >= 16 && NEKO_MFMA_LDS_FITS(16, LX, SZ) ? 16 :\
   (NWF) / NEKO_MFMA_NGROUPS(LX) >= 8 && NEKO_MFMA_LDS_FITS(8, LX, SZ) ? 8 :  \
   (NWF) / NEKO_MFMA_NGROUPS(LX) >= 4 && NEKO_MFMA_LDS_FITS(4, LX, SZ) ? 4 :  \
   (NWF) / NEKO_MFMA_NGROUPS(LX) >= 2 && NEKO_MFMA_LDS_FITS(2, LX, SZ) ? 2 : 1)

Definition at line 437 of file mfma_kernel.h.

◆ NEKO_MFMA_LDS_FITS

#define NEKO_MFMA_LDS_FITS	(	EB,
		LX,
		SZ
	)	(NEKO_MFMA_LDS_N(EB, LX, SZ) <= NEKO_EB_MAX_LDS)

Definition at line 428 of file mfma_kernel.h.

◆ NEKO_MFMA_LDS_N

#define NEKO_MFMA_LDS_N	(	EB,
		LX,
		SZ
	)	((3 * NEKO_MFMA_DMAT_N(LX, SZ) + 4 * (EB) * NEKO_MFMA_CUBE_N(LX, SZ)) * (SZ))

Definition at line 426 of file mfma_kernel.h.

◆ NEKO_MFMA_NBLCKS

#define NEKO_MFMA_NBLCKS	(	NELV,
		LX,
		C
	)	dim3(((NELV) + NEKO_MFMA_EB(LX, C) - 1) / NEKO_MFMA_EB(LX, C), 1, 1)

Definition at line 495 of file mfma_kernel.h.

◆ NEKO_MFMA_NGROUPS

#define NEKO_MFMA_NGROUPS ( LX ) (((LX) * (LX) + 15) / 16)

Definition at line 418 of file mfma_kernel.h.

◆ NEKO_MFMA_NTHRDS

#define NEKO_MFMA_NTHRDS ( C ) dim3(64, NEKO_MFMA_NWF(C), 1)

Definition at line 265 of file mfma_kernel.h.

◆ NEKO_MFMA_NWF

#define NEKO_MFMA_NWF ( C ) (1 << ((C) % NEKO_MFMA_NWF_CANDIDATES))

Definition at line 263 of file mfma_kernel.h.

◆ NEKO_MFMA_NWF_CANDIDATES

#define NEKO_MFMA_NWF_CANDIDATES 5

Definition at line 259 of file mfma_kernel.h.

◆ NEKO_MFMA_PAD_N

#define NEKO_MFMA_PAD_N	(	LX,
		SZ
	)	((((LX) % 4 == 0) && !((LX) == 4 && (SZ) == 4)) ? 1 : 0)

Definition at line 324 of file mfma_kernel.h.

◆ NEKO_MFMA_SD_N

#define NEKO_MFMA_SD_N	(	LX,
		SZ
	)	NEKO_MFMA_SJ_N(LX, SZ)

Definition at line 350 of file mfma_kernel.h.

◆ NEKO_MFMA_SJ_N

#define NEKO_MFMA_SJ_N	(	LX,
		SZ
	)	((LX) + NEKO_MFMA_PAD_N(LX, SZ))

Definition at line 328 of file mfma_kernel.h.

◆ NEKO_MFMA_SK_N

#define NEKO_MFMA_SK_N	(	LX,
		SZ
	)	(NEKO_MFMA_SJ_N(LX, SZ) * (LX))

Definition at line 329 of file mfma_kernel.h.

◆ NEKO_MFMA_SPT

#define NEKO_MFMA_SPT	(	LX,
		C
	)	NEKO_MFMA_SPT_N(NEKO_MFMA_WPE(LX, C), LX, sizeof(real))

Definition at line 461 of file mfma_kernel.h.

◆ NEKO_MFMA_SPT_N

#define NEKO_MFMA_SPT_N	(	WPE,
		LX,
		SZ
	)	((NEKO_MFMA_CUBE_N(LX, SZ) + (WPE) * 64 - 1) / ((WPE) * 64))

Definition at line 459 of file mfma_kernel.h.

◆ NEKO_MFMA_TILE

#define NEKO_MFMA_TILE ( C ) ((C) / NEKO_MFMA_NWF_CANDIDATES)

Definition at line 264 of file mfma_kernel.h.

◆ NEKO_MFMA_TILE_CANDIDATES

#define NEKO_MFMA_TILE_CANDIDATES 2

Definition at line 260 of file mfma_kernel.h.

◆ NEKO_MFMA_VECTOR_GREG

#define NEKO_MFMA_VECTOR_GREG	(	LX,
		C
	)	NEKO_MFMA_VECTOR_GREG_N(NEKO_MFMA_SPT(LX, C), sizeof(real))

Definition at line 492 of file mfma_kernel.h.

◆ NEKO_MFMA_VECTOR_GREG_N

#define NEKO_MFMA_VECTOR_GREG_N	(	SPT,
		SZ
	)	((7 * (SPT) * ((SZ) / 4)) <= NEKO_MFMA_VECTOR_GREG_VGPRS)

Definition at line 490 of file mfma_kernel.h.

◆ NEKO_MFMA_VECTOR_GREG_VGPRS

#define NEKO_MFMA_VECTOR_GREG_VGPRS 64

Definition at line 488 of file mfma_kernel.h.

◆ NEKO_MFMA_WPE

#define NEKO_MFMA_WPE	(	LX,
		C
	)	(NEKO_MFMA_NWF(C) / NEKO_MFMA_EB(LX, C))

Definition at line 445 of file mfma_kernel.h.

◆ NEKO_TUNE_LOG_MFMA

#define NEKO_TUNE_LOG_MFMA	(	LX,
		T3
	)

Value:

  do {                                                                        \
    for (int c = 0; c < NEKO_MFMA_CANDIDATES; c++) {                          \
      if ((T3)[c] >= NEKO_TUNE_INIT) { continue; }                            \
      sprintf(neko_log_buf, "MFMA  %s %2dwf %-2de %-5s: %9.2f us/call",      \
              mfma_tile_name(c), NEKO_MFMA_NWF(c), NEKO_MFMA_EB(LX, c),       \
              NEKO_MFMA_PAD_N(LX, sizeof(real)) ? "pad" : "plain",            \
              NEKO_TUNE_US((T3)[c], iters));                                  \
      log_message(neko_log_buf);                                              \
    }                                                                         \
  } while (0)

Definition at line 602 of file mfma_kernel.h.

◆ NEKO_TUNE_LOG_MFMA_VEC

#define NEKO_TUNE_LOG_MFMA_VEC	(	LX,
		T3
	)

Value:

  do {                                                                        \
    for (int c = 0; c < NEKO_MFMA_CANDIDATES; c++) {                          \
      if ((T3)[c] >= NEKO_TUNE_INIT) { continue; }                            \
      sprintf(neko_log_buf, "MFMA  %s %2dwf %-2de %-5s %-4s: %9.2f us/call",  \
              mfma_tile_name(c), NEKO_MFMA_NWF(c), NEKO_MFMA_EB(LX, c),       \
              NEKO_MFMA_PAD_N(LX, sizeof(real)) ? "pad" : "plain",            \
              NEKO_MFMA_VECTOR_GREG(LX, c) ? "reg" : "glob",                  \
              NEKO_TUNE_US((T3)[c], iters));                                  \
      log_message(neko_log_buf);                                              \
    }                                                                         \
  } while (0)

Definition at line 623 of file mfma_kernel.h.

Function Documentation

◆ hip_have_mfma()

static bool hip_have_mfma ( )

inlinestatic

Returns true if the current device exposes the matrix cores used by the MFMA strategies (currently gfx90a / MI250X and gfx942 / MI300A / MI300X) and this build actually compiled the device code for them. The f64 4x4x4 / 16x16x4 and f32 16x16x4 instructions are all available on these arches, so a single check gates either precision. Result is cached after the first query.

Definition at line 136 of file mfma_kernel.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ hip_mfma_arch_probe()

static __global__ void hip_mfma_arch_probe ( int * flag )

static

Shared matrix-core (MFMA) primitives for the spectral-element tensor contractions, used by the Ax-helm autotuner strategy.

Recovered onto develop 2026-08-22 from the feature/mfma branch, scoped down to Ax-helm: the other operators it covered (dudxyz, cdtp, conv1, convect_scalar) are the four worst placed in the arithmetic intensity ranking and are not worth carrying without a measurement. The contraction primitives below are unchanged from the hardware-validated versions – see the layout note on mfma_contract_4x4 for why they must not be re-derived from host simulation alone.

Maps a reference derivative-matrix * field contraction onto the AMD matrix cores on CDNA2 (MI250X / gfx90a) and CDNA3 (MI300A / MI300X / gfx942). Double precision defaults to the batched 4x4x4 tile (v_mfma_f64_4x4x4f64, full M-utilisation for M = LX < 16; see mfma_contract_4x4), single precision uses the 16x16x4 tile (v_mfma_f32_16x16x4f32; no f32 4x4x4 equivalent). Double precision can be forced back onto 16x16x4 with -DMFMA_F64_USE_16X16. One thread block is launched as a single 64-lane wavefront (blockDim = (64,1,1), or (64, NWF, 1) for the multi-wavefront Ax-helm) and processes one element; each contraction is a D * U GEMM with M = LX, N = LX^2, K = LX, the partial tiles masked off.

Supported for single and double precision and 4 <= LX <= 12; the upper bound is set by the LDS needed to keep the cubes resident (operators stage four cubes plus three LX^2 derivative matrices, 63360 B of f64 at LX = 12 with the padded Ax-helm layout, 58752 B without it).

16x16x4 register/lane layout (used by the f32 tile and the f64 16x16x4 fallback; the 4x4x4 layout is documented at mfma_contract_4x4).

Verification status, since this tile came from the same source that got the 4x4x4 block selector and contraction index the wrong way round: CONFIRMED on gfx90a 2026-08-22, both accumulator packings. The f32 sweep of mfma_probe covers the A and B layouts (shared between precisions) and the i = 4*g + r packing at ~1e-7, fp32 epsilon; a second run with -DMFMA_F64_USE_16X16 covers the i = g + 4*r packing at ~1e-16. A layout error reads as O(1) here, so neither pass rests on a loose tolerance.

Note the f64 results are bit-identical between this tile and the 4x4x4 one. That is expected rather than suspicious: both decompose K into chunks of four with one MFMA per chunk accumulated in sequence, so the summation order is the same. It is not the signature of dead code, which is what identical results across NWF would be – NWF changes which wavefront takes which column group and so must perturb the schedule. D = A*B + C, wave of 64 lanes, g = lane/16 in 0..3, c = lane%16 in 0..15. A and B share the same layout for both precisions; only the accumulator packing of the 16x16 result differs: A[i][k] : lane holds A[i = c][k = g] B[k][j] : lane holds B[k = g][j = c] D[i][j] : lane holds D[j = c] in accumulator slot r = 0..3, with row i = g + 4*r for f64 (rows spread with stride 4) and i = 4*g + r for f32 (four contiguous rows per lane group)

Definition at line 120 of file mfma_kernel.h.

Here is the caller graph for this function:

◆ mfma_lx_supported()

template<const int LX>

static bool mfma_lx_supported ( )

inlinestatic

Compile-time predicate for the LX values that the MFMA strategy supports. Single or double precision and 4 <= LX <= 12 (bounded above by the LDS needed to keep one element resident). MUST match the dispatch specialisations in each operator's kernel header.

Definition at line 185 of file mfma_kernel.h.

Here is the call graph for this function:

◆ mfma_tile_name()

static const char * mfma_tile_name ( const int c )

inlinestatic

Definition at line 374 of file mfma_kernel.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ mfma_tile_offered()

static bool mfma_tile_offered ( const int c )

inlinestatic

Definition at line 362 of file mfma_kernel.h.

Here is the caller graph for this function:

◆ neko_mfma_candidates()

static int neko_mfma_candidates ( )

inlinestatic

Definition at line 593 of file mfma_kernel.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ neko_mfma_nwf_pin()

static int neko_mfma_nwf_pin ( )

static

Definition at line 570 of file mfma_kernel.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ neko_mfma_pin()

static int neko_mfma_pin ( )

static

Definition at line 533 of file mfma_kernel.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ neko_mfma_sweep()

static int neko_mfma_sweep ( )

static

Definition at line 511 of file mfma_kernel.h.

Here is the call graph for this function:

Here is the caller graph for this function:

Macros

Functions

Macro Definition Documentation

◆ NEKO_MFMA_CANDIDATES

◆ NEKO_MFMA_CUBE_N

◆ NEKO_MFMA_DMAT_N

◆ NEKO_MFMA_EB

◆ NEKO_MFMA_EB_N

◆ NEKO_MFMA_LDS_FITS

◆ NEKO_MFMA_LDS_N

◆ NEKO_MFMA_NBLCKS

◆ NEKO_MFMA_NGROUPS

◆ NEKO_MFMA_NTHRDS

◆ NEKO_MFMA_NWF

◆ NEKO_MFMA_NWF_CANDIDATES

◆ NEKO_MFMA_PAD_N

◆ NEKO_MFMA_SD_N

◆ NEKO_MFMA_SJ_N

◆ NEKO_MFMA_SK_N

◆ NEKO_MFMA_SPT

◆ NEKO_MFMA_SPT_N

◆ NEKO_MFMA_TILE

◆ NEKO_MFMA_TILE_CANDIDATES

◆ NEKO_MFMA_VECTOR_GREG

◆ NEKO_MFMA_VECTOR_GREG_N

◆ NEKO_MFMA_VECTOR_GREG_VGPRS

◆ NEKO_MFMA_WPE

◆ NEKO_TUNE_LOG_MFMA

◆ NEKO_TUNE_LOG_MFMA_VEC

Function Documentation

◆ hip_have_mfma()

◆ hip_mfma_arch_probe()

◆ mfma_lx_supported()

◆ mfma_tile_name()

◆ mfma_tile_offered()

◆ neko_mfma_candidates()

◆ neko_mfma_nwf_pin()

◆ neko_mfma_pin()

◆ neko_mfma_sweep()