#include "wave.h"

Include dependency graph for elem_block.h:

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

Classes
struct	elem_block< LX, C >

struct	chunk_block< LX, C >

Macros
#define	NEKO_EB_WAVE NEKO_WAVE_SIZE_UNIFORM

#define	NEKO_EB_MAX_LDS 65536

#define	NEKO_EB_CANDIDATES 3

#define	NEKO_EB_MIN_WAVES_EU 3

#define	NEKO_EB_BOUNDS(NT) __launch_bounds__((NT), NEKO_EB_MIN_WAVES_EU)

#define	NEKO_EB(LX, C) (elem_block<LX, C>::value)

#define	NEKO_EB_NTHRDS(LX, C) dim3((LX), (LX), NEKO_EB(LX, C))

#define	NEKO_EB_NBLCKS(NELV, LX, C) dim3(((NELV) + NEKO_EB(LX, C) - 1)/NEKO_EB(LX, C), 1, 1)

#define	NEKO_EB_SEL(LX, SEL) ((SEL) == 0 ? NEKO_EB(LX, 0) : (SEL) == 1 ? NEKO_EB(LX, 1) : NEKO_EB(LX, 2))

#define	NEKO_CHUNKS_CANDIDATES 4

#define	NEKO_CHUNKS(LX, C) (chunk_block<LX, C>::value)

#define	NEKO_CHUNKS_NTHRDS(LX, C) dim3(NEKO_CHUNKS(LX, C), 1, 1)

#define	NEKO_CHUNKS_SEL(LX, SEL)

Macro Definition Documentation

◆ NEKO_CHUNKS

#define NEKO_CHUNKS	(	LX,
		C
	)	(chunk_block<LX, C>::value)

Definition at line 154 of file elem_block.h.

◆ NEKO_CHUNKS_CANDIDATES

#define NEKO_CHUNKS_CANDIDATES 4

Chunk size for the 1d kernels

CHUNKS is both the thread block size and the stride over the LX^3 points of an element. At the historical 1024 that is badly mismatched at low order: only 64 of 1024 threads do work at LX = 4, 125 at LX = 5, 512 at LX = 8. Shared memory is sized by LX rather than by CHUNKS, so a smaller block also raises the number of elements resident per SM at no extra cost.

CONSTRAINT: the 1d kernels stage the derivative matrices with a single if (iii < LX*LX) guard, so a block smaller than one matrix would leave part of it unwritten. Candidates below LX*LX are therefore rejected outright and fall back to 1024, which always satisfies it for LX <= 16.

Definition at line 134 of file elem_block.h.

◆ NEKO_CHUNKS_NTHRDS

#define NEKO_CHUNKS_NTHRDS	(	LX,
		C
	)	dim3(NEKO_CHUNKS(LX, C), 1, 1)

Definition at line 155 of file elem_block.h.

◆ NEKO_CHUNKS_SEL

#define NEKO_CHUNKS_SEL	(	LX,
		SEL
	)

Value:

((SEL) == 0 ? NEKO_CHUNKS(LX, 0) : (SEL) == 1 ? NEKO_CHUNKS(LX, 1) : \

(SEL) == 2 ? NEKO_CHUNKS(LX, 2) : NEKO_CHUNKS(LX, 3))

ale_add_kinematics_kernel

__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)

Definition ale_kinematics_kernel.h:52

NEKO_CHUNKS

#define NEKO_CHUNKS(LX, C)

Definition elem_block.h:145

Definition at line 156 of file elem_block.h.

◆ NEKO_EB

#define NEKO_EB	(	LX,
		C
	)	(elem_block<LX, C>::value)

Definition at line 113 of file elem_block.h.

◆ NEKO_EB_BOUNDS

#define NEKO_EB_BOUNDS ( NT ) __launch_bounds__((NT), NEKO_EB_MIN_WAVES_EU)

Definition at line 105 of file elem_block.h.

◆ NEKO_EB_CANDIDATES

#define NEKO_EB_CANDIDATES 3

Definition at line 82 of file elem_block.h.

◆ NEKO_EB_MAX_LDS

#define NEKO_EB_MAX_LDS 65536

Definition at line 79 of file elem_block.h.

◆ NEKO_EB_MIN_WAVES_EU

#define NEKO_EB_MIN_WAVES_EU 3

Definition at line 102 of file elem_block.h.

◆ NEKO_EB_NBLCKS

#define NEKO_EB_NBLCKS	(	NELV,
		LX,
		C
	)	dim3(((NELV) + NEKO_EB(LX, C) - 1)/NEKO_EB(LX, C), 1, 1)

Definition at line 115 of file elem_block.h.

◆ NEKO_EB_NTHRDS

#define NEKO_EB_NTHRDS	(	LX,
		C
	)	dim3((LX), (LX), NEKO_EB(LX, C))

Definition at line 114 of file elem_block.h.

◆ NEKO_EB_SEL

#define NEKO_EB_SEL	(	LX,
		SEL
	)	((SEL) == 0 ? NEKO_EB(LX, 0) : (SEL) == 1 ? NEKO_EB(LX, 1) : NEKO_EB(LX, 2))

Definition at line 117 of file elem_block.h.

◆ NEKO_EB_WAVE

#define NEKO_EB_WAVE NEKO_WAVE_SIZE_UNIFORM

Elements per thread block for the SEM operator kstep kernels (HIP)

A kstep block is one (LX,LX) thread plane per element. Against a wavefront that masks off a lot of lanes for every LX whose square is not a multiple of the wave: at the 64 lanes of CDNA, 25% useful at LX = 4, 39% at LX = 5, 63% at LX = 9. Stacking EB elements along threadIdx.z packs the block back up.

MEASURED TO LOSE on gfx90a and gfx942 for ax_helm: the blocked specialisations roughly double VGPR usage and spill to scratch (kstep_padded<double,8,4>: 168 VGPRs, 2596 B/lane scratch, against 110/0 unblocked), and the cause of the register blow-up is still not understood. A single precision run at lx = 8 puts numbers on it: 130.7 us/call at one element per block against 1212 us at four and 1213 at eight, a factor of 9.3.

The sweep is nevertheless on by default here as of 2026-08-22, as it is on CUDA. Encoding the verdict in a default meant the tuner could never re-measure it, and the ranking does invert on NVIDIA, where blocking measures a 1.6x win. The tuner rejects the blocked variants on its own; the price is tuning time, and at 9.3x slower those two candidates are most of it – roughly 0.7 s per polynomial order at the default sampling.

Candidate 0 is one element per block. Candidate C > 0 fits as many whole elements as it can into 2^(C+1) wavefronts.

NEKO_EB_WAVE is deliberately NEKO_WAVE_SIZE_UNIFORM and not NEKO_WAVE_SIZE: it sizes a kernel template argument and the matching launch geometry, which the host and device passes have to agree on, and only the command line form of the width is visible to both. On RDNA, pass -DNEKO_WAVE_SIZE=32 to size the candidates for a 32 lane wave; getting it wrong costs occupancy, not correctness, and this sweep defaults off on AMD regardless.

Definition at line 74 of file elem_block.h.

Classes

Macros