elem_block.h File Reference

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

Go to the source code of this file.

Classes
struct	elem_block< LX, C >
struct	chunk_block< LX, C >

Macros
#define	NEKO_EB_MAX_BLOCK_THREADS   512
#define	NEKO_EB_MAX_SMEM   49152
#define	NEKO_EB_CANDIDATES   3
#define	NEKO_EB_BOUNDS(NT)   __launch_bounds__((NT))
#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 145 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 125 of file elem_block.h.

NEKO_CHUNKS_NTHRDS

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

Definition at line 146 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))

Definition at line 147 of file elem_block.h.

NEKO_EB

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

Definition at line 104 of file elem_block.h.

NEKO_EB_BOUNDS

#define NEKO_EB_BOUNDS

(

NT

)

__launch_bounds__((NT))

Definition at line 95 of file elem_block.h.

NEKO_EB_CANDIDATES

#define NEKO_EB_CANDIDATES   3

Definition at line 63 of file elem_block.h.

NEKO_EB_MAX_BLOCK_THREADS

#define NEKO_EB_MAX_BLOCK_THREADS   512

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

A kstep block is one (LX,LX) thread plane per element, which is half a warp at LX = 4 and two warps at LX = 8. Stacking EB elements along threadIdx.z packs the block up and lets the element independent derivative matrices be loaded once per block rather than once per element.

Whether that pays depends on the kernel's register footprint and on the architecture, so the variants are instantiated for the candidates below and picked between at runtime by each operator's autotuner. Measured on GH200 for ax_helm at lx = 4: 13.55 -> 8.34 us/call at EB = 4.

Candidate C selects EB = 2^C, clamped so a block stays within NEKO_EB_MAX_BLOCK_THREADS. Clamped candidates collapse onto the same specialisation, which just means the tuner times it twice.

Definition at line 55 of file elem_block.h.

NEKO_EB_MAX_SMEM

#define NEKO_EB_MAX_SMEM   49152

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

NEKO_EB_NTHRDS

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

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