#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 1

#define	NEKO_EB_BOUNDS(NT) __launch_bounds__((NT))

#define	NEKO_EB_WAVE_BARRIER 0

#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)

Functions
__device__ __host__ constexpr bool	neko_eb_slice_in_wave (unsigned wave, int lx)

template<int LX>
__device__ __forceinline__ void	neko_eb_kstep_barrier ()

Macro Definition Documentation

◆ NEKO_CHUNKS

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

Definition at line 294 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 274 of file elem_block.h.

◆ NEKO_CHUNKS_NTHRDS

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

Definition at line 295 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 296 of file elem_block.h.

◆ NEKO_EB

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

Definition at line 253 of file elem_block.h.

◆ NEKO_EB_BOUNDS

#define NEKO_EB_BOUNDS ( NT ) __launch_bounds__((NT))

Definition at line 135 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 1

Definition at line 127 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 255 of file elem_block.h.

◆ NEKO_EB_NTHRDS

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

Definition at line 254 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 257 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.

◆ NEKO_EB_WAVE_BARRIER

#define NEKO_EB_WAVE_BARRIER 0

Barrier for the kstep k loop

A kstep block is dim3(LX, LX, EB), and every LDS buffer the k loop touches – shu, shur, shus, and the nine of the vector variant – is indexed from eb*LX*LX. One private slice per element, nothing shared between slices. The only cross slice data is shdx/shdy/shdz, written by eb == 0 before the loop, and the __syncthreads() guarding that stays where it is.

So when a slice cannot straddle a wavefront, the two barriers inside the loop are synchronising waves that share nothing, and a wave scoped barrier is enough. The alignment test is exact rather than conservative: the linear thread id of slice z runs [z*LX*LX, (z+1)*LX*LX), so slices land on wavefront boundaries exactly when LX*LX divides the wave. On wave64 that admits LX = 2, 4 and 8.

Note LX = 4, where one wave holds four whole slices: that is also where the packing has something to gain, 16 of 64 lanes busy unblocked. This is not a knob that is only legal where it is useless – unlike readfirstlane on the element base, whose validity condition picks out exactly the orders with no upside.

OFF BY DEFAULT, and MEASURED TO BUY NOTHING at lx = 8. On MI250X, 13824 elements: kstep eb=4 runs 2254, 2280 and 2240 us over three builds, one of which has this on. If the 2*LX workgroup barriers were what the blocked variants were paying for, removing them should have collapsed eb=4 back toward eb=1 – a factor of five. It did not move outside the noise.

Calibrate against that noise before reading anything into a single pair. Over three nominally identical runs the spread is 1.7% on kstep eb=1, 1.8% on eb=4, but 8.1% on the MFMA kernels and 8.4% on kstep eb=8 – and the MFMA kernels contain no barrier at all, so 8% swings there are the machine, not the code. Anything under ~10% on the low occupancy kernels is not a result.

So the residual eb cost is NOT barrier scope – look to the divergent element base (nelv enters the address, defeating the gfx9 SADDR form) and to the hardware co-residency floor at EB = 8, where 8 wave64s on 4 SIMDs cap occupancy at 2 waves/SIMD regardless of any launch bound.

Before trusting that refutation, confirm the mechanism actually engaged: the s_barrier count in the ISA must drop between a build with this off and one with it on. Equal counts mean the barrier was never removed and the experiment said nothing.

Kept because it is off, cheap and documented, and because lx = 4 is untested – though the prior is now weak, since lx = 4 blocking is also 4 wavefronts, the same shape that showed nothing at lx = 8.

It weakens a synchronisation scope, so it fails as wrong results rather than as an error, and wrong results inside a Krylov solve present as poor convergence rather than as anything that looks like a bug. The benchmark reports no correctness signal at all. Validate by pinning two strategies and requiring bit identical output. Never by checking that the solver still converges.

Ordering is spelled out rather than left to an inline asm memory clobber. A clobber constrains the compiler and says nothing to the memory model, and __builtin_amdgcn_wave_barrier() is itself only a scheduling barrier – it emits no instruction. The release/acquire pair around it is what states the requirement, the same shape the backend lowers a workgroup barrier into, at wavefront scope instead.

Definition at line 203 of file elem_block.h.

Function Documentation

◆ neko_eb_kstep_barrier()

template<int LX>

__device__ __forceinline__ void neko_eb_kstep_barrier ( )

Definition at line 230 of file elem_block.h.

Here is the call graph for this function:

◆ neko_eb_slice_in_wave()

__device__ __host__ constexpr bool neko_eb_slice_in_wave	(	unsigned	wave,
		int	lx
	)

constexpr

Definition at line 223 of file elem_block.h.