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 64

#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 137 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 117 of file elem_block.h.

◆ NEKO_CHUNKS_NTHRDS

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

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

◆ NEKO_EB

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

Definition at line 96 of file elem_block.h.

◆ NEKO_EB_BOUNDS

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

Definition at line 88 of file elem_block.h.

◆ NEKO_EB_CANDIDATES

#define NEKO_EB_CANDIDATES 3

Definition at line 65 of file elem_block.h.

◆ NEKO_EB_MAX_LDS

#define NEKO_EB_MAX_LDS 65536

Definition at line 62 of file elem_block.h.

◆ NEKO_EB_MIN_WAVES_EU

#define NEKO_EB_MIN_WAVES_EU 3

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

◆ NEKO_EB_NTHRDS

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

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

◆ NEKO_EB_WAVE

#define NEKO_EB_WAVE 64

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

A kstep block is one (LX,LX) thread plane per element. Against a 64 wide wavefront that masks off a lot of lanes for every LX whose square is not a multiple of the wave: 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). The sweep therefore defaults off on this backend. The machinery is kept because the ranking inverts on NVIDIA, where blocking measures a 1.6x win, and because the cause of the AMD register blow-up is not yet understood.

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

Definition at line 57 of file elem_block.h.

Classes

Macros