Functions

◆ small_wg()

template<class T>

int batchlas::ops::gemm::small_wg	(	bool	nn,
		int	max_dim )

constexpr

The work-group small launches: the batched kernel's 128 lanes, or the float NN tiled leg's (NB/4)^2 with NB = 48 or 56 (small_batched.hh).

Definition at line 101 of file choice.hh.

◆ small_needs_sg32()

template<class T>

bool batchlas::ops::gemm::small_needs_sg32	(	bool	nn,
		int	max_dim )

constexpr

small's batched leg is [[sycl::reqd_sub_group_size(32)]]; only the float NN tiled leg is not.

Definition at line 116 of file choice.hh.

◆ small_fits()

template<class T>

bool batchlas::ops::gemm::small_fits	(	const select::Device &	d,
		bool	nn,
		std::int64_t	max_dim )

The shape-and-device half of small's can_run, here so tests can probe synthetic devices.

Definition at line 121 of file choice.hh.

◆ candidates()

template<class T>

auto batchlas::ops::gemm::candidates ( )

constexpr

Every compiled choice, once, in tie-break order (§6.3): simpler first, vendor last.

Definition at line 129 of file choice.hh.

Types

◆ GemmChoice

using batchlas::ops::gemm::GemmChoice = std::variant<Direct, Tiled, Small, Reg, Wide, Vendor>

natives: GPU or vendor-free host, Default precision

Definition at line 45 of file choice.hh.

Variables

◆ reg_configs

std::array<RegCfg, 10> batchlas::ops::gemm::reg_configs

inlineconstexpr

Initial value:

{{
    {32, 32, 8, 1, 2, 2, 1, {true, false, false, false}, false},
    {64, 64, 8, 1, 4, 4, 1, {true, false, false, false}, false},
    {64, 64, 16, 1, 4, 4, 1, {true, true, true, true}, false},
    {128, 32, 16, 1, 4, 4, 2, {true, true, true, true}, false},
    {128, 32, 32, 1, 4, 4, 2, {true, true, true, true}, true},
    {128, 64, 16, 1, 4, 4, 1, {false, true, true, true}, false},
    {32, 128, 16, 1, 4, 4, 1, {true, false, true, true}, false},
    {128, 64, 32, 4, 8, 4, 2, {true, false, false, false}, true},
    {128, 64, 32, 2, 8, 4, 2, {true, false, false, false}, true},
    {128, 128, 8, 1, 8, 8, 2, {true, false, false, false}, true},  
}}

Definition at line 65 of file choice.hh.

◆ wide_configs

std::array<WideCfg, 5> batchlas::ops::gemm::wide_configs

inlineconstexpr

Initial value:

{{
    {64, 64, 16, 4, 4, true, true, true},
    {128, 32, 16, 4, 4, false, false, true},  
    {32, 128, 16, 4, 4, false, true, false},  
    {32, 32, 16, 4, 4, true, false, false},
    {16, 16, 16, 2, 2, true, false, false},
}}

Definition at line 85 of file choice.hh.

◆ kSmallMaxDim

int batchlas::ops::gemm::kSmallMaxDim = 64

inlineconstexpr

Definition at line 94 of file choice.hh.

◆ kSmallWg

int batchlas::ops::gemm::kSmallWg = 128

inlineconstexpr

Definition at line 95 of file choice.hh.

◆ kSmallTiledMaxDim

int batchlas::ops::gemm::kSmallTiledMaxDim = 56

inlineconstexpr

Definition at line 96 of file choice.hh.

◆ kDirectWg

int batchlas::ops::gemm::kDirectWg = 64

inlineconstexpr

Definition at line 106 of file choice.hh.

◆ kTiledWg

int batchlas::ops::gemm::kTiledWg = 256

inlineconstexpr

Definition at line 107 of file choice.hh.

◆ kMaxGridBatch

std::int64_t batchlas::ops::gemm::kMaxGridBatch = 65535

inlineconstexpr

direct, tiled, reg and wide put the batch in SYCL dim 0 = CUDA grid z (65535); small is 1-D.

Definition at line 109 of file choice.hh.

◆ is_complex_v

template<class T>

bool batchlas::ops::gemm::is_complex_v = !std::is_same_v<T, float> && !std::is_same_v<T, double>

inlineconstexpr

Definition at line 112 of file choice.hh.

◆ last_resort

std::array<std::string_view, 2> batchlas::ops::gemm::last_resort {"direct", "vendor"}

inlineconstexpr

Generality order (§5.5): direct serves every GPU shape, vendor everything else (CPU, precision).

Definition at line 152 of file choice.hh.

◆ spec

select::OpSpec batchlas::ops::gemm::spec {Op::gemm, select::Lib::level3, {last_resort}}

inlineconstexpr

op, vendor library, rules

Definition at line 153 of file choice.hh.

◆ key_names

std::array<std::string_view, 7> batchlas::ops::gemm::key_names

inlineconstexpr

Initial value:

{"ta:exact", "tb:exact", "layout:exact", "m:log",

"n:log", "k:log", "batch:log"}

Table keys: ta, tb (N|T|C; C folds to T for a real scalar) and layout (packed = A, B, C contiguous with 16-byte bases) exact; m, n, k, batch log, each weight 1 (work ~ m n k batch).

Definition at line 157 of file choice.hh.

◆ grid_square

std::array<int, 14> batchlas::ops::gemm::grid_square {8, 16, 24, 32, 48, 64, 96, 128, 192, 256, 384, 512, 768, 1024}

inlineconstexpr

The tuner's demand-driven grid; the transcriber (tuned/README.md) spells it again.

Squares for every form and both layouts; panels and skinny shapes for the issued forms only, packed panels from m, n >= 128. evidence: docs/design/flat-kernel-selection-phase3-plan.md (§3)

Definition at line 163 of file choice.hh.

◆ grid_panel_mn

std::array<int, 7> batchlas::ops::gemm::grid_panel_mn {32, 64, 128, 256, 512, 1024, 2048}

inlineconstexpr

Definition at line 164 of file choice.hh.

◆ grid_panel_k

std::array<int, 6> batchlas::ops::gemm::grid_panel_k {8, 16, 32, 64, 96, 128}

inlineconstexpr

Definition at line 165 of file choice.hh.

◆ grid_skinny_mn

std::array<int, 6> batchlas::ops::gemm::grid_skinny_mn {64, 128, 256, 512, 1024, 2048}

inlineconstexpr

Definition at line 166 of file choice.hh.

◆ grid_skinny_k

std::array<int, 2> batchlas::ops::gemm::grid_skinny_k {256, 1024}

inlineconstexpr

Definition at line 167 of file choice.hh.

◆ grid_batch

std::array<int, 3> batchlas::ops::gemm::grid_batch {128, 2048, 32768}

inlineconstexpr

Definition at line 168 of file choice.hh.

◆ grid_packed_panel_min

int batchlas::ops::gemm::grid_packed_panel_min = 128

inlineconstexpr

Definition at line 169 of file choice.hh.

◆ grid_real_forms

std::array<std::string_view, 4> batchlas::ops::gemm::grid_real_forms {"NN", "NT", "TN", "TT"}

inlineconstexpr

Definition at line 170 of file choice.hh.

◆ grid_real_panel_forms

std::array<std::string_view, 3> batchlas::ops::gemm::grid_real_panel_forms {"NN", "NT", "TN"}

inlineconstexpr

Definition at line 171 of file choice.hh.

◆ grid_complex_forms

std::array<std::string_view, 6> batchlas::ops::gemm::grid_complex_forms {"NN", "NT", "TN", "NC", "CN", "CT"}

inlineconstexpr

Definition at line 172 of file choice.hh.

◆ grid_complex_panel_forms

std::array<std::string_view, 5> batchlas::ops::gemm::grid_complex_panel_forms {"NN", "NT", "NC", "TN", "CN"}

inlineconstexpr

Definition at line 173 of file choice.hh.

Classes
struct	Direct
	gemm_direct: one work-item per element, any form; max_wg >= 64 More...
struct	Tiled
	gemm_tiled: 16x16 shared-memory tile, any form; max_wg >= 256 More...
struct	Small
	gemm_small: several matrices per work-group; real, see small_fits More...
struct	Reg
	gemm_reg, float only: m x n x k macro tile with k unroll u; must name a reg_configs entry that instantiates the call's transpose form, and max_wg must cover its threads(). More...
struct	Wide
	gemm_wide, every scalar: 16-byte-granule m x n x k tiles; must name a wide_configs entry that instantiates the call's form, and max_wg must cover its threads(). More...
struct	Vendor
	backend::gemm_vendor; needs the level-3 library More...
struct	Forms
	Transpose forms an instantiation exists for, after the real-scalar C->T fold. More...
struct	RegCfg
	One compiled register config: tr/tc thread tile, stages, the forms instantiated, and whether NN takes the unpredicated leg when the layout allows (derived per call, never a gate). More...
struct	WideCfg
	One compiled wide config. More...

Functions

◆ small_wg()

◆ small_needs_sg32()

◆ small_fits()

◆ candidates()

Types

◆ GemmChoice

Variables

◆ reg_configs

◆ wide_configs

◆ kSmallMaxDim

◆ kSmallWg

◆ kSmallTiledMaxDim

◆ kDirectWg

◆ kTiledWg

◆ kMaxGridBatch

◆ is_complex_v

◆ last_resort

◆ spec

◆ key_names

◆ grid_square

◆ grid_panel_mn

◆ grid_panel_k

◆ grid_skinny_mn

◆ grid_skinny_k

◆ grid_batch

◆ grid_packed_panel_min

◆ grid_real_forms

◆ grid_real_panel_forms

◆ grid_complex_forms

◆ grid_complex_panel_forms

Classes