BatchLAS Documentation BatchLASv0.1.0
Classes
batchlas::device::detail::subgroup Namespace Reference

Functions

◆ gemm()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::gemm ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
MatrixMatrixOperand< T > operand,
GeneralMatrixTransform transform )
inlineconstexpr

Definition at line 40 of file group_blas_gemm.hh.

◆ gemm_register_tiled()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::gemm_register_tiled ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
MatrixMatrixOperand< T > operand,
GeneralMatrixTransform transform,
detail::subgroup::GemmWorkspace< T > * gemm_workspace )
inlineconstexpr

Definition at line 96 of file group_blas_gemm.hh.

◆ gemm_aligned_nn_large()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::gemm_aligned_nn_large ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
MatrixMatrixOperand< T > operand,
detail::subgroup::GemmWorkspace< T > * gemm_workspace )
inlineconstexpr

Definition at line 220 of file group_blas_gemm.hh.

◆ rank2k_complex_tiled()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::rank2k_complex_tiled ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
MatrixMatrixOperand< T > operand,
SymmetricRank2kTransform transform,
ComplexRank2kWorkspace< T > * workspace )
inlineconstexpr

Definition at line 89 of file group_blas_rankk.hh.

◆ rankk_complex_tiled()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::rankk_complex_tiled ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
RankKOperand< T > operand,
SymmetricRankKTransform transform,
ComplexRank2kWorkspace< T > * workspace )
inlineconstexpr

Definition at line 165 of file group_blas_rankk.hh.

◆ rank2k_register_tiled()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::rank2k_register_tiled ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
MatrixMatrixOperand< T > operand,
SymmetricRank2kTransform transform,
RegisterMatrixWorkspace< T > * workspace )
inlineconstexpr

Definition at line 225 of file group_blas_rankk.hh.

◆ rankk_register_tiled()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::rankk_register_tiled ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
RankKOperand< T > operand,
SymmetricRankKTransform transform,
RegisterMatrixWorkspace< T > * workspace )
inlineconstexpr

Definition at line 299 of file group_blas_rankk.hh.

◆ static_for_impl()

template<int Count, typename Fn, std::size_t... I>
void batchlas::device::detail::subgroup::static_for_impl ( Fn && fn,
std::index_sequence< I... >  )
inlineconstexpr

Definition at line 94 of file group_blas_subgroup_common.hh.

◆ static_for()

template<int Count, typename Fn>
void batchlas::device::detail::subgroup::static_for ( Fn && fn)
inlineconstexpr

Definition at line 99 of file group_blas_subgroup_common.hh.

◆ supports_aligned_packet_loads()

template<typename T, int Width>
bool batchlas::device::detail::subgroup::supports_aligned_packet_loads ( const T * ptr,
int ld,
int stride )
inlineconstexpr

Definition at line 198 of file group_blas_subgroup_common.hh.

◆ packet_load_aligned()

template<typename T, int Width>
Packet< T, Width > batchlas::device::detail::subgroup::packet_load_aligned ( const T * ptr,
int offset )
inlineconstexpr

Definition at line 222 of file group_blas_subgroup_common.hh.

◆ register_matrix_lhs_stage()

int batchlas::device::detail::subgroup::register_matrix_lhs_stage ( int tile_idx)
inlineconstexpr

Definition at line 226 of file group_blas_subgroup_common.hh.

◆ register_matrix_rhs_stage()

int batchlas::device::detail::subgroup::register_matrix_rhs_stage ( int tile_idx)
inlineconstexpr

Definition at line 230 of file group_blas_subgroup_common.hh.

◆ matrix_tile_group_row() [1/2]

template<typename Item>
int batchlas::device::detail::subgroup::matrix_tile_group_row ( const Item & )
inlineconstexpr

Definition at line 235 of file group_blas_subgroup_common.hh.

◆ matrix_tile_group_col() [1/2]

template<typename Item>
int batchlas::device::detail::subgroup::matrix_tile_group_col ( const Item & )
inlineconstexpr

Definition at line 240 of file group_blas_subgroup_common.hh.

◆ matrix_tile_group_row_stride() [1/2]

template<typename Item>
int batchlas::device::detail::subgroup::matrix_tile_group_row_stride ( const Item & )
inlineconstexpr

Definition at line 245 of file group_blas_subgroup_common.hh.

◆ matrix_tile_group_col_stride() [1/2]

template<typename Item>
int batchlas::device::detail::subgroup::matrix_tile_group_col_stride ( const Item & )
inlineconstexpr

Definition at line 250 of file group_blas_subgroup_common.hh.

◆ matrix_tile_group_row() [2/2]

int batchlas::device::detail::subgroup::matrix_tile_group_row ( const sycl::nd_item< 3 > & item)
inlineconstexpr

Definition at line 254 of file group_blas_subgroup_common.hh.

◆ matrix_tile_group_col() [2/2]

int batchlas::device::detail::subgroup::matrix_tile_group_col ( const sycl::nd_item< 3 > & item)
inlineconstexpr

Definition at line 258 of file group_blas_subgroup_common.hh.

◆ matrix_tile_group_row_stride() [2/2]

int batchlas::device::detail::subgroup::matrix_tile_group_row_stride ( const sycl::nd_item< 3 > & item)
inlineconstexpr

Definition at line 262 of file group_blas_subgroup_common.hh.

◆ matrix_tile_group_col_stride() [2/2]

int batchlas::device::detail::subgroup::matrix_tile_group_col_stride ( const sycl::nd_item< 3 > & item)
inlineconstexpr

Definition at line 266 of file group_blas_subgroup_common.hh.

◆ subgroup_size() [1/2]

template<typename Item>
int batchlas::device::detail::subgroup::subgroup_size ( const Item & item)
inlineconstexpr

Definition at line 271 of file group_blas_subgroup_common.hh.

◆ subgroup_size() [2/2]

int batchlas::device::detail::subgroup::subgroup_size ( const DeviceBlasLaunchInfo & launch)
inlineconstexpr

Definition at line 275 of file group_blas_subgroup_common.hh.

◆ subgroup_local_id()

template<typename Item>
int batchlas::device::detail::subgroup::subgroup_local_id ( const Item & item)
inlineconstexpr

Definition at line 280 of file group_blas_subgroup_common.hh.

◆ subgroup_group_id()

template<typename Item>
int batchlas::device::detail::subgroup::subgroup_group_id ( const Item & item)
inlineconstexpr

Definition at line 285 of file group_blas_subgroup_common.hh.

◆ subgroup_count() [1/2]

template<typename Item>
int batchlas::device::detail::subgroup::subgroup_count ( const Item & item)
inlineconstexpr

Definition at line 290 of file group_blas_subgroup_common.hh.

◆ subgroup_count() [2/2]

int batchlas::device::detail::subgroup::subgroup_count ( const DeviceBlasLaunchInfo & launch)
inlineconstexpr

Definition at line 295 of file group_blas_subgroup_common.hh.

◆ is_nd_item_3d_launch() [1/2]

template<typename Exec>
bool batchlas::device::detail::subgroup::is_nd_item_3d_launch ( const Exec & )
inlineconstexpr

Definition at line 301 of file group_blas_subgroup_common.hh.

◆ is_nd_item_3d_launch() [2/2]

bool batchlas::device::detail::subgroup::is_nd_item_3d_launch ( const DeviceBlasLaunchInfo & launch)
inlineconstexpr

Definition at line 305 of file group_blas_subgroup_common.hh.

◆ rows_per_subgroup()

int batchlas::device::detail::subgroup::rows_per_subgroup ( int extent,
int total_sg,
int max_rows )
inlineconstexpr

Definition at line 309 of file group_blas_subgroup_common.hh.

◆ matrix_tile_k()

int batchlas::device::detail::subgroup::matrix_tile_k ( int contract_extent,
int sg_size,
Side side )
inlineconstexpr

Definition at line 319 of file group_blas_subgroup_common.hh.

◆ matrix_rows_per_subgroup()

int batchlas::device::detail::subgroup::matrix_rows_per_subgroup ( int row_extent,
int total_sg,
int sg_size,
int contract_extent,
Side side )
inlineconstexpr

Definition at line 329 of file group_blas_subgroup_common.hh.

◆ subgroup_policy_matches()

bool batchlas::device::detail::subgroup::subgroup_policy_matches ( DeviceBlasPolicy policy,
int actual_size )
inlineconstexpr

Definition at line 340 of file group_blas_subgroup_common.hh.

◆ can_use_matrix_fast_path()

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_matrix_fast_path ( const Item & item,
int row_extent,
int col_extent,
int contract_extent,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 355 of file group_blas_subgroup_common.hh.

◆ can_use_matrix_register_fast_path()

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_matrix_register_fast_path ( const Item & item,
int row_extent,
int col_extent,
int contract_extent,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 376 of file group_blas_subgroup_common.hh.

◆ can_use_matrix_aligned_nn_large_fast_path() [1/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_matrix_aligned_nn_large_fast_path ( const Item & item,
int row_extent,
int col_extent,
int contract_extent,
Transpose trans_a,
Transpose trans_b,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 446 of file group_blas_subgroup_common.hh.

◆ can_use_matrix_aligned_nn_large_fast_path() [2/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_matrix_aligned_nn_large_fast_path ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
const MatrixMatrixOperand< T > & operand,
GeneralMatrixTransform transform,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 410 of file group_blas_subgroup_common.hh.

◆ can_use_complex_rank2k_tiled_fast_path() [1/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_complex_rank2k_tiled_fast_path ( const Item & item,
int extent,
int contract_extent,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 521 of file group_blas_subgroup_common.hh.

◆ can_use_complex_rank2k_tiled_fast_path() [2/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_complex_rank2k_tiled_fast_path ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
const MatrixMatrixOperand< T > & operand,
SymmetricRank2kTransform transform,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 486 of file group_blas_subgroup_common.hh.

◆ can_use_complex_rankk_tiled_fast_path() [1/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_complex_rankk_tiled_fast_path ( const Item & item,
int extent,
int contract_extent,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 593 of file group_blas_subgroup_common.hh.

◆ can_use_complex_rankk_tiled_fast_path() [2/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_complex_rankk_tiled_fast_path ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
const RankKOperand< T > & operand,
SymmetricRankKTransform transform,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 558 of file group_blas_subgroup_common.hh.

◆ can_use_rank2k_register_fast_path() [1/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_rank2k_register_fast_path ( const Item & item,
int extent,
int contract_extent,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 665 of file group_blas_subgroup_common.hh.

◆ can_use_rank2k_register_fast_path() [2/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_rank2k_register_fast_path ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
const MatrixMatrixOperand< T > & operand,
SymmetricRank2kTransform transform,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 630 of file group_blas_subgroup_common.hh.

◆ can_use_rankk_register_fast_path() [1/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_rankk_register_fast_path ( const Item & item,
int extent,
int contract_extent,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 737 of file group_blas_subgroup_common.hh.

◆ can_use_rankk_register_fast_path() [2/2]

template<typename T, typename Item>
bool batchlas::device::detail::subgroup::can_use_rankk_register_fast_path ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
const RankKOperand< T > & operand,
SymmetricRankKTransform transform,
DeviceBlasPolicy policy )
inlineconstexpr

Definition at line 702 of file group_blas_subgroup_common.hh.

◆ triangular_begin()

int batchlas::device::detail::subgroup::triangular_begin ( int index,
int extent,
TriangularTransform transform,
Side side )
inlineconstexpr

Definition at line 767 of file group_blas_subgroup_common.hh.

◆ triangular_end()

int batchlas::device::detail::subgroup::triangular_end ( int index,
int extent,
TriangularTransform transform,
Side side )
inlineconstexpr

Definition at line 781 of file group_blas_subgroup_common.hh.

◆ triangular_tile_begin()

int batchlas::device::detail::subgroup::triangular_tile_begin ( int row_base,
int row_extent,
int col_base,
int col_extent,
int contract_extent,
TriangularTransform transform )
inlineconstexpr

Definition at line 795 of file group_blas_subgroup_common.hh.

◆ triangular_tile_end()

int batchlas::device::detail::subgroup::triangular_tile_end ( int row_base,
int row_extent,
int col_base,
int col_extent,
int contract_extent,
TriangularTransform transform )
inlineconstexpr

Definition at line 824 of file group_blas_subgroup_common.hh.

◆ load_register_matrix_stage()

template<typename Item, typename T, typename LhsLoader, typename RhsLoader>
void batchlas::device::detail::subgroup::load_register_matrix_stage ( const Item & item,
RegisterMatrixWorkspace< T > * workspace,
int linear_tid,
int subgroup_id,
int lhs_stage,
int rhs_stage,
bool load_lhs,
bool load_rhs,
int k_base,
int tile_extent,
LhsLoader && lhs_loader,
RhsLoader && rhs_loader )
inlineconstexpr

Definition at line 852 of file group_blas_subgroup_common.hh.

◆ accumulate_register_matrix_stage()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::accumulate_register_matrix_stage ( const Item & item,
const RegisterMatrixWorkspace< T > * workspace,
int subgroup_id,
int lhs_stage,
int rhs_stage,
int local_row,
int local_col,
int tile_extent,
RegisterMatrixAccumTile< T > & accum )
inlineconstexpr

Definition at line 914 of file group_blas_subgroup_common.hh.

◆ write_register_matrix_tile()

template<typename T>
void batchlas::device::detail::subgroup::write_register_matrix_tile ( MatrixMatrixOperand< T > operand,
int row_base,
int col_base,
int row_extent,
int col_extent,
const RegisterMatrixAccumTile< T > & accum )
inlineconstexpr

Definition at line 961 of file group_blas_subgroup_common.hh.

◆ write_rank2k_register_tile()

template<typename T, typename Transform>
void batchlas::device::detail::subgroup::write_rank2k_register_tile ( MatrixMatrixOperand< T > operand,
Transform transform,
int row_base,
int col_base,
int extent,
const RegisterMatrixAccumTile< T > & accum )
inlineconstexpr

Definition at line 985 of file group_blas_subgroup_common.hh.

◆ write_complex_rank2k_tile()

template<typename T, typename Transform>
void batchlas::device::detail::subgroup::write_complex_rank2k_tile ( MatrixMatrixOperand< T > operand,
Transform transform,
int row_base,
int col_base,
int extent,
const ComplexRank2kAccumTile< T > & accum )
inlineconstexpr

Definition at line 1015 of file group_blas_subgroup_common.hh.

◆ accumulate_rank2k_register_tiled_pass()

template<typename Item, typename T, typename LhsLoader, typename RhsLoader>
void batchlas::device::detail::subgroup::accumulate_rank2k_register_tiled_pass ( const Item & item,
RegisterMatrixWorkspace< T > * workspace,
int linear_tid,
int subgroup_id,
int contract_extent,
LhsLoader && lhs_loader,
RhsLoader && rhs_loader,
RegisterMatrixAccumTile< T > & accum )
inlineconstexpr

Definition at line 1045 of file group_blas_subgroup_common.hh.

◆ accumulate_complex_rank2k_tiled_pass_impl()

template<int TileK, int TileAStride, typename Workspace, typename Item, typename T, typename LhsLoader, typename RhsLoader>
void batchlas::device::detail::subgroup::accumulate_complex_rank2k_tiled_pass_impl ( const Item & item,
Workspace * workspace,
int linear_tid,
int subgroup_id,
int local_row,
int local_col,
int contract_extent,
LhsLoader && lhs_loader,
RhsLoader && rhs_loader,
ComplexRank2kAccumTile< T > & accum )
inlineconstexpr

Definition at line 1136 of file group_blas_subgroup_common.hh.

◆ symm()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::symm ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
MatrixMatrixOperand< T > operand,
SymmetricTransform transform )
inlineconstexpr

Definition at line 43 of file group_blas_symm.hh.

◆ symm_register_tiled()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::symm_register_tiled ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
MatrixMatrixOperand< T > operand,
SymmetricTransform transform,
RegisterMatrixWorkspace< T > * workspace )
inlineconstexpr

Definition at line 115 of file group_blas_symm.hh.

◆ trmm()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::trmm ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
MatrixMatrixOperand< T > operand,
TriangularTransform transform )
inlineconstexpr

Definition at line 58 of file group_blas_trmm.hh.

◆ trmm_register_tiled()

template<typename Item, typename T>
void batchlas::device::detail::subgroup::trmm_register_tiled ( const Item & item,
const KernelMatrixView< T, MatrixFormat::Dense > & a,
MatrixMatrixOperand< T > operand,
TriangularTransform transform,
RegisterMatrixWorkspace< T > * workspace )
inlineconstexpr

Definition at line 155 of file group_blas_trmm.hh.

Types

◆ Packet

template<typename T, int Width>
using batchlas::device::detail::subgroup::Packet = std::conditional_t<supports_packet_v<T, Width>, sycl::vec<T, Width>, PacketArray<T, Width>>

Definition at line 219 of file group_blas_subgroup_common.hh.

Variables

◆ kMaxSupportedSubgroupSize

int batchlas::device::detail::subgroup::kMaxSupportedSubgroupSize = 32
inlineconstexpr

Definition at line 11 of file group_blas_subgroup_common.hh.

◆ kMaxSubgroupsPerWorkGroup

int batchlas::device::detail::subgroup::kMaxSubgroupsPerWorkGroup = 8
inlineconstexpr

Definition at line 12 of file group_blas_subgroup_common.hh.

◆ kMaxMatrixRowsPerSubgroup

int batchlas::device::detail::subgroup::kMaxMatrixRowsPerSubgroup = 8
inlineconstexpr

Definition at line 13 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixTileM

int batchlas::device::detail::subgroup::kRegisterMatrixTileM = 128
inlineconstexpr

Definition at line 14 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixTileN

int batchlas::device::detail::subgroup::kRegisterMatrixTileN = 64
inlineconstexpr

Definition at line 15 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixTileK

int batchlas::device::detail::subgroup::kRegisterMatrixTileK = 32
inlineconstexpr

Definition at line 16 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixThreadTileRows

int batchlas::device::detail::subgroup::kRegisterMatrixThreadTileRows = 4
inlineconstexpr

Definition at line 17 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixThreadTileCols

int batchlas::device::detail::subgroup::kRegisterMatrixThreadTileCols = 8
inlineconstexpr

Definition at line 18 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixLocalRows

int batchlas::device::detail::subgroup::kRegisterMatrixLocalRows = kRegisterMatrixTileM / kRegisterMatrixThreadTileRows
inlineconstexpr

Definition at line 19 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixLocalCols

int batchlas::device::detail::subgroup::kRegisterMatrixLocalCols = kRegisterMatrixTileN / kRegisterMatrixThreadTileCols
inlineconstexpr

Definition at line 20 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixThreadsPerGroup

int batchlas::device::detail::subgroup::kRegisterMatrixThreadsPerGroup = kRegisterMatrixLocalRows * kRegisterMatrixLocalCols
inlineconstexpr

Definition at line 21 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixSubgroupRows

int batchlas::device::detail::subgroup::kRegisterMatrixSubgroupRows = kMaxSupportedSubgroupSize / kRegisterMatrixLocalCols
inlineconstexpr

Definition at line 22 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixSubgroupTileM

int batchlas::device::detail::subgroup::kRegisterMatrixSubgroupTileM = kRegisterMatrixSubgroupRows * kRegisterMatrixThreadTileRows
inlineconstexpr

Definition at line 23 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixTileAStride

int batchlas::device::detail::subgroup::kRegisterMatrixTileAStride = kRegisterMatrixTileK + 1
inlineconstexpr

Definition at line 24 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixTileBStride

int batchlas::device::detail::subgroup::kRegisterMatrixTileBStride = kRegisterMatrixTileN + 1
inlineconstexpr

Definition at line 25 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixLhsStages

int batchlas::device::detail::subgroup::kRegisterMatrixLhsStages = 1
inlineconstexpr

Definition at line 26 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixRhsStages

int batchlas::device::detail::subgroup::kRegisterMatrixRhsStages = 1
inlineconstexpr

Definition at line 27 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kTileM

int batchlas::device::detail::subgroup::kComplexRank2kTileM = 64
inlineconstexpr

Definition at line 28 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kTileN

int batchlas::device::detail::subgroup::kComplexRank2kTileN = 64
inlineconstexpr

Definition at line 29 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kTileK

int batchlas::device::detail::subgroup::kComplexRank2kTileK = 32
inlineconstexpr

Definition at line 30 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kInKernelTileK

int batchlas::device::detail::subgroup::kComplexRank2kInKernelTileK = 24
inlineconstexpr

Definition at line 31 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kThreadTileRows

int batchlas::device::detail::subgroup::kComplexRank2kThreadTileRows = 4
inlineconstexpr

Definition at line 32 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kThreadTileCols

int batchlas::device::detail::subgroup::kComplexRank2kThreadTileCols = 4
inlineconstexpr

Definition at line 33 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kLocalRows

int batchlas::device::detail::subgroup::kComplexRank2kLocalRows = kComplexRank2kTileM / kComplexRank2kThreadTileRows
inlineconstexpr

Definition at line 34 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kLocalCols

int batchlas::device::detail::subgroup::kComplexRank2kLocalCols = kComplexRank2kTileN / kComplexRank2kThreadTileCols
inlineconstexpr

Definition at line 35 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kThreadsPerGroup

int batchlas::device::detail::subgroup::kComplexRank2kThreadsPerGroup = kComplexRank2kLocalRows * kComplexRank2kLocalCols
inlineconstexpr

Definition at line 36 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kSubgroupRows

int batchlas::device::detail::subgroup::kComplexRank2kSubgroupRows = kMaxSupportedSubgroupSize / kComplexRank2kLocalCols
inlineconstexpr

Definition at line 37 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kSubgroupTileM

int batchlas::device::detail::subgroup::kComplexRank2kSubgroupTileM = kComplexRank2kSubgroupRows * kComplexRank2kThreadTileRows
inlineconstexpr

Definition at line 38 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kTileAStride

int batchlas::device::detail::subgroup::kComplexRank2kTileAStride = kComplexRank2kTileK + 1
inlineconstexpr

Definition at line 39 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kInKernelTileAStride

int batchlas::device::detail::subgroup::kComplexRank2kInKernelTileAStride = kComplexRank2kInKernelTileK + 1
inlineconstexpr

Definition at line 40 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kTileBStride

int batchlas::device::detail::subgroup::kComplexRank2kTileBStride = kComplexRank2kTileN + 1
inlineconstexpr

Definition at line 41 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmTileM

int batchlas::device::detail::subgroup::kOptimizedGemmTileM = 128
inlineconstexpr

Definition at line 42 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmTileN

int batchlas::device::detail::subgroup::kOptimizedGemmTileN = 32
inlineconstexpr

Definition at line 43 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmTileK

int batchlas::device::detail::subgroup::kOptimizedGemmTileK = 32
inlineconstexpr

Definition at line 44 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmThreadTileRows

int batchlas::device::detail::subgroup::kOptimizedGemmThreadTileRows = 4
inlineconstexpr

Definition at line 45 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmThreadTileCols

int batchlas::device::detail::subgroup::kOptimizedGemmThreadTileCols = 4
inlineconstexpr

Definition at line 46 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmVecA

int batchlas::device::detail::subgroup::kOptimizedGemmVecA = 4
inlineconstexpr

Definition at line 47 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmVecB

int batchlas::device::detail::subgroup::kOptimizedGemmVecB = 4
inlineconstexpr

Definition at line 48 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmUnrollK

int batchlas::device::detail::subgroup::kOptimizedGemmUnrollK = 2
inlineconstexpr

Definition at line 49 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmStages

int batchlas::device::detail::subgroup::kOptimizedGemmStages = 2
inlineconstexpr

Definition at line 50 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmLocalRows

int batchlas::device::detail::subgroup::kOptimizedGemmLocalRows = kOptimizedGemmTileM / kOptimizedGemmThreadTileRows
inlineconstexpr

Definition at line 51 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmLocalCols

int batchlas::device::detail::subgroup::kOptimizedGemmLocalCols = kOptimizedGemmTileN / kOptimizedGemmThreadTileCols
inlineconstexpr

Definition at line 52 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmThreadsPerGroup

int batchlas::device::detail::subgroup::kOptimizedGemmThreadsPerGroup = kOptimizedGemmLocalRows * kOptimizedGemmLocalCols
inlineconstexpr

Definition at line 53 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmTileAStride

int batchlas::device::detail::subgroup::kOptimizedGemmTileAStride = kOptimizedGemmTileM + 1
inlineconstexpr

Definition at line 54 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmTileBStride

int batchlas::device::detail::subgroup::kOptimizedGemmTileBStride = kOptimizedGemmTileK + 1
inlineconstexpr

Definition at line 55 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmStageASize

int batchlas::device::detail::subgroup::kOptimizedGemmStageASize = kOptimizedGemmTileAStride * kOptimizedGemmTileK
inlineconstexpr

Definition at line 56 of file group_blas_subgroup_common.hh.

◆ kOptimizedGemmStageBSize

int batchlas::device::detail::subgroup::kOptimizedGemmStageBSize = kOptimizedGemmTileBStride * kOptimizedGemmTileN
inlineconstexpr

Definition at line 57 of file group_blas_subgroup_common.hh.

◆ kSubgroupWorkspaceBudgetBytes

std::size_t batchlas::device::detail::subgroup::kSubgroupWorkspaceBudgetBytes
inlineconstexpr
Initial value:
=
static_cast<std::size_t>(::batchlas::device_limits::subgroup_workspace_budget_bytes())

Definition at line 58 of file group_blas_subgroup_common.hh.

◆ subgroup_limit_for_workspace_v

template<std::size_t FixedElements, std::size_t ElementsPerSubgroup, typename T>
int batchlas::device::detail::subgroup::subgroup_limit_for_workspace_v
inlineconstexpr
Initial value:
= []() {
constexpr std::size_t budget_elements = kSubgroupWorkspaceBudgetBytes / sizeof(T);
if constexpr (budget_elements <= FixedElements) {
return 0;
} else {
constexpr std::size_t remaining_elements = budget_elements - FixedElements;
constexpr std::size_t subgroup_limit = ElementsPerSubgroup == 0
? static_cast<std::size_t>(kMaxSubgroupsPerWorkGroup)
: (remaining_elements / ElementsPerSubgroup);
return static_cast<int>(std::min<std::size_t>(static_cast<std::size_t>(kMaxSubgroupsPerWorkGroup), subgroup_limit));
}
}()

Definition at line 62 of file group_blas_subgroup_common.hh.

◆ kRegisterMatrixWorkspaceMaxSubgroupsPerWorkGroup

template<typename T>
int batchlas::device::detail::subgroup::kRegisterMatrixWorkspaceMaxSubgroupsPerWorkGroup
inlineconstexpr
Initial value:
= subgroup_limit_for_workspace_v<
kRegisterMatrixRhsStages * kRegisterMatrixTileK * kRegisterMatrixTileBStride,
kRegisterMatrixLhsStages * kRegisterMatrixSubgroupTileM * kRegisterMatrixTileAStride,
T>

Definition at line 76 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kWorkspaceMaxSubgroupsPerWorkGroup

template<typename T>
int batchlas::device::detail::subgroup::kComplexRank2kWorkspaceMaxSubgroupsPerWorkGroup
inlineconstexpr
Initial value:
= subgroup_limit_for_workspace_v<
kComplexRank2kTileK * kComplexRank2kTileBStride,
kComplexRank2kSubgroupTileM * kComplexRank2kTileAStride,
T>

Definition at line 82 of file group_blas_subgroup_common.hh.

◆ kComplexRank2kInKernelWorkspaceMaxSubgroupsPerWorkGroup

template<typename T>
int batchlas::device::detail::subgroup::kComplexRank2kInKernelWorkspaceMaxSubgroupsPerWorkGroup
inlineconstexpr
Initial value:
= subgroup_limit_for_workspace_v<
kComplexRank2kInKernelTileK * kComplexRank2kTileBStride,
kComplexRank2kSubgroupTileM * kComplexRank2kInKernelTileAStride,
T>

Definition at line 88 of file group_blas_subgroup_common.hh.

◆ register_matrix_workspace_supported_v

template<typename T>
bool batchlas::device::detail::subgroup::register_matrix_workspace_supported_v = sizeof(RegisterMatrixWorkspace<T>) <= kSubgroupWorkspaceBudgetBytes
inlineconstexpr

Definition at line 173 of file group_blas_subgroup_common.hh.

◆ complex_rank2k_workspace_supported_v

template<typename T>
bool batchlas::device::detail::subgroup::complex_rank2k_workspace_supported_v = sizeof(ComplexRank2kWorkspace<T>) <= kSubgroupWorkspaceBudgetBytes
inlineconstexpr

Definition at line 176 of file group_blas_subgroup_common.hh.

◆ complex_rank2k_in_kernel_workspace_supported_v

template<typename T>
bool batchlas::device::detail::subgroup::complex_rank2k_in_kernel_workspace_supported_v
inlineconstexpr
Initial value:
=
sizeof(ComplexRank2kInKernelWorkspace<T>) <= kSubgroupWorkspaceBudgetBytes

Definition at line 179 of file group_blas_subgroup_common.hh.

◆ optimized_gemm_workspace_supported_v

template<typename T>
bool batchlas::device::detail::subgroup::optimized_gemm_workspace_supported_v = sizeof(OptimizedGemmWorkspace<T>) <= kSubgroupWorkspaceBudgetBytes
inlineconstexpr

Definition at line 183 of file group_blas_subgroup_common.hh.

◆ gemm_workspace_supported_v

template<typename T>
bool batchlas::device::detail::subgroup::gemm_workspace_supported_v = sizeof(GemmWorkspace<T>) <= kSubgroupWorkspaceBudgetBytes
inlineconstexpr

Definition at line 186 of file group_blas_subgroup_common.hh.

◆ supports_packet4_v

template<typename T>
bool batchlas::device::detail::subgroup::supports_packet4_v = std::is_same_v<T, float>
inlineconstexpr

Definition at line 189 of file group_blas_subgroup_common.hh.

◆ supports_packet_v

template<typename T, int Width>
bool batchlas::device::detail::subgroup::supports_packet_v = false
inlineconstexpr

Definition at line 192 of file group_blas_subgroup_common.hh.

◆ supports_packet_v< T, 4 >

template<typename T>
bool batchlas::device::detail::subgroup::supports_packet_v< T, 4 > = supports_packet4_v<T>
inlineconstexpr

Definition at line 195 of file group_blas_subgroup_common.hh.

Classes

struct  RegisterMatrixAccumTile
struct  ComplexRank2kAccumTile
struct  RegisterMatrixWorkspace
struct  ComplexRank2kWorkspace
struct  ComplexRank2kInKernelWorkspace
struct  OptimizedGemmWorkspace
struct  GemmWorkspace
struct  PacketArray
struct  RegisterMatrixThreadTileValues
Generated by doxygen 1.18.0