Overview

Launch-description and operand builders, and workspace sizes.

Functions

◆ make_group_launch_info()

DeviceBlasLaunchInfo batchlas::device::make_group_launch_info ( int local_size )

inlineconstexpr

Launch description for a call that receives a sycl::group.

Parameters

local_size work-items in the executor; clamped to at least 1

Definition at line 57 of file group_blas_common.hh.

◆ make_nd_item_1d_launch_info()

DeviceBlasLaunchInfo batchlas::device::make_nd_item_1d_launch_info	(	int	local_size,
		int	subgroup_size )

inlineconstexpr

Launch description for a call that receives a sycl::nd_item<1>.

Parameters

local_size	work-group size; clamped to at least 1
subgroup_size	sub-group size the kernel runs with; clamped to at least 1

Definition at line 65 of file group_blas_common.hh.

◆ make_nd_item_3d_launch_info()

DeviceBlasLaunchInfo batchlas::device::make_nd_item_3d_launch_info	(	int	local_size,
		int	subgroup_size )

inlineconstexpr

Launch description for a call that receives a sycl::nd_item<3> (tiled output).

Parameters

local_size	work-group size; the register-tiled paths require 256
subgroup_size	sub-group size the kernel runs with; clamped to at least 1

Definition at line 73 of file group_blas_common.hh.

◆ make_matvec_operand()

template<typename T>

MatrixVectorOperand< T > batchlas::device::make_matvec_operand	(	const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1),
		T	beta = T(0) )

inlineconstexpr

Builds a MatrixVectorOperand.

Parameters

x	input vector
y	output vector
alpha	scale of the product
beta	scale of the input y

Definition at line 188 of file group_blas_common.hh.

◆ make_matmat_operand()

template<typename T>

MatrixMatrixOperand< T > batchlas::device::make_matmat_operand	(	const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0) )

inlineconstexpr

Builds a MatrixMatrixOperand.

Parameters

b	second input matrix
c	output matrix
alpha	scale of the product
beta	scale of the input C

Definition at line 202 of file group_blas_common.hh.

◆ make_rankk_operand()

template<typename T>

RankKOperand< T > batchlas::device::make_rankk_operand	(	const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0) )

inlineconstexpr

Builds a RankKOperand.

Parameters

c	square output matrix
alpha	scale of the product
beta	scale of the input C

Definition at line 215 of file group_blas_common.hh.

◆ make_rank1_update_operand()

template<typename T>

Rank1UpdateOperand< T > batchlas::device::make_rank1_update_operand	(	const VectorView< T > &	y,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		T	alpha = T(1) )

inlineconstexpr

Builds a Rank1UpdateOperand.

Parameters

y	right vector
a	matrix updated in place
alpha	scale of the outer product

Definition at line 227 of file group_blas_common.hh.

◆ gemm_workspace_elements() [1/2]

template<typename T, Transpose TransAV = Transpose::NoTrans, Transpose TransBV = Transpose::NoTrans>

std::size_t batchlas::device::gemm_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	row_extent,
		int	col_extent,
		int	contract_extent,
		bool	aligned_a = false,
		bool	aligned_b = false )

inlineconstexpr

Local memory, in elements of T, that gemm() can use for this launch and shape.

Allocate it as a sycl::local_accessor<T, 1> and pass its pointer; when the result is 0 pass nullptr. Use the same template arguments as the call.

Parameters

launch	description of the launch the call runs in
row_extent	m, rows of C
col_extent	n, columns of C
contract_extent	k
aligned_a	the caller guarantees A meets the aligned NN kernel's 4-element alignment
aligned_b	the same for B

Returns: element count, or 0 when no staged path applies

Definition at line 479 of file group_blas_gemm.hh.

◆ gemm_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Transpose TransAV = Transpose::NoTrans, Transpose TransBV = Transpose::NoTrans>

std::size_t batchlas::device::gemm_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	row_extent,
		int	col_extent,
		int	contract_extent,
		bool	aligned_a = false,
		bool	aligned_b = false )

inlineconstexpr

Local memory, in elements of T, that gemm() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
row_extent	rows of C
col_extent	columns of C
contract_extent	contraction length k
aligned_a	the caller guarantees A meets the aligned kernel's 4-element alignment
aligned_b	the same for B

Returns: element count, or 0 when no staged path applies

Definition at line 502 of file group_blas_gemm.hh.

◆ gemv_workspace_elements() [1/2]

template<typename T, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::gemv_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	rows,
		int	cols )

inlineconstexpr

Local memory, in elements of T, that gemv() can use; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
rows	rows of A
cols	columns of A

Returns: element count, or 0 when no staged path applies

Definition at line 238 of file group_blas_gemv.hh.

◆ gemv_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::gemv_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	rows,
		int	cols )

inlineconstexpr

Local memory, in elements of T, that gemv() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
rows	rows of A
cols	columns of A

Returns: element count, or 0 when no staged path applies

Definition at line 251 of file group_blas_gemv.hh.

◆ syrk_workspace_elements() [1/2]

template<typename T, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::syrk_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent,
		int	contract_extent )

inlineconstexpr

Local memory, in elements of T, that syrk() can use for this launch and shape.

Parameters

launch	description of the launch the call runs in
extent	n, the order of C
contract_extent	k

Returns: element count, or 0 when no staged path applies (then pass nullptr)

Definition at line 512 of file group_blas_rankk.hh.

◆ syrk_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::syrk_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent,
		int	contract_extent )

inlineconstexpr

Local memory, in elements of T, that syrk() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix
contract_extent	contraction length k

Returns: element count, or 0 when no staged path applies

Definition at line 528 of file group_blas_rankk.hh.

◆ herk_workspace_elements() [1/2]

template<typename T, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::herk_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent,
		int	contract_extent )

inlineconstexpr

Local memory, in elements of T, that herk() can use; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix
contract_extent	contraction length k

Returns: element count, or 0 when no staged path applies

Definition at line 597 of file group_blas_rankk.hh.

◆ herk_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::herk_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent,
		int	contract_extent )

inlineconstexpr

Local memory, in elements of T, that herk() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix
contract_extent	contraction length k

Returns: element count, or 0 when no staged path applies

Definition at line 613 of file group_blas_rankk.hh.

◆ syr2k_workspace_elements() [1/2]

template<typename T, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::syr2k_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent,
		int	contract_extent )

inlineconstexpr

Local memory, in elements of T, that syr2k() can use; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix
contract_extent	contraction length k

Returns: element count, or 0 when no staged path applies

Definition at line 684 of file group_blas_rankk.hh.

◆ syr2k_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::syr2k_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent,
		int	contract_extent )

inlineconstexpr

Local memory, in elements of T, that syr2k() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix
contract_extent	contraction length k

Returns: element count, or 0 when no staged path applies

Definition at line 700 of file group_blas_rankk.hh.

◆ her2k_workspace_elements() [1/2]

template<typename T, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::her2k_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent,
		int	contract_extent )

inlineconstexpr

Local memory, in elements of T, that her2k() can use; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix
contract_extent	contraction length k

Returns: element count, or 0 when no staged path applies

Definition at line 770 of file group_blas_rankk.hh.

◆ her2k_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans>

std::size_t batchlas::device::her2k_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent,
		int	contract_extent )

inlineconstexpr

Local memory, in elements of T, that her2k() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix
contract_extent	contraction length k

Returns: element count, or 0 when no staged path applies

Definition at line 786 of file group_blas_rankk.hh.

◆ symm_workspace_elements() [1/2]

template<typename T, Side SideV = Side::Left, Uplo UploV = Uplo::Upper>

std::size_t batchlas::device::symm_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	row_extent,
		int	col_extent )

inlineconstexpr

Local memory, in elements of T, that symm() can use; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
row_extent	rows of C
col_extent	columns of C

Returns: element count, or 0 when no staged path applies

Definition at line 334 of file group_blas_symm.hh.

◆ symm_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Side SideV = Side::Left, Uplo UploV = Uplo::Upper>

std::size_t batchlas::device::symm_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	row_extent,
		int	col_extent )

inlineconstexpr

Local memory, in elements of T, that symm() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
row_extent	rows of C
col_extent	columns of C

Returns: element count, or 0 when no staged path applies

Definition at line 347 of file group_blas_symm.hh.

◆ symv_workspace_elements() [1/2]

template<typename T, Uplo UploV = Uplo::Upper>

std::size_t batchlas::device::symv_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent )

inlineconstexpr

Local memory, in elements of T, that symv() can use; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix

Returns: element count, or 0 when no staged path applies

Definition at line 290 of file group_blas_symv.hh.

◆ symv_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Uplo UploV = Uplo::Upper>

std::size_t batchlas::device::symv_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent )

inlineconstexpr

Local memory, in elements of T, that symv() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix

Returns: element count, or 0 when no staged path applies

Definition at line 301 of file group_blas_symv.hh.

◆ hemv_workspace_elements() [1/2]

template<typename T, Uplo UploV = Uplo::Upper>

std::size_t batchlas::device::hemv_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent )

inlineconstexpr

Local memory, in elements of T, that hemv() can use; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix

Returns: element count, or 0 when no staged path applies

Definition at line 361 of file group_blas_symv.hh.

◆ hemv_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Uplo UploV = Uplo::Upper>

std::size_t batchlas::device::hemv_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	extent )

inlineconstexpr

Local memory, in elements of T, that hemv() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
extent	order of the matrix

Returns: element count, or 0 when no staged path applies

Definition at line 372 of file group_blas_symv.hh.

◆ trmm_workspace_elements() [1/2]

template<Side SideV = Side::Left, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans, Diag DiagV = Diag::NonUnit, typename T>

std::size_t batchlas::device::trmm_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	row_extent,
		int	col_extent,
		bool	aliased = false )

inlineconstexpr

Local memory, in elements of T, that trmm() can use for this launch and shape.

Nonzero only for an nd_item<3> launch description whose shape admits the register-tiled path. When the result is 0 pass nullptr.

Trap: T is the LAST template parameter here, after the four mode parameters, unlike every other *_workspace_elements and unlike the policy overload.

Parameters

launch	description of the launch the call runs in
row_extent	rows of C
col_extent	columns of C
aliased	B and C overlap; the call then takes no staged path, so this returns 0

Returns: element count, or 0 when no staged path applies

Definition at line 410 of file group_blas_trmm.hh.

◆ trmm_workspace_elements() [2/2]

template<typename T, DeviceBlasPolicy Policy, Side SideV = Side::Left, Uplo UploV = Uplo::Upper, Transpose TransV = Transpose::NoTrans, Diag DiagV = Diag::NonUnit>

std::size_t batchlas::device::trmm_workspace_elements	(	const DeviceBlasLaunchInfo &	launch,
		int	row_extent,
		int	col_extent,
		bool	aliased = false )

inlineconstexpr

Local memory, in elements of T, that trmm() can use under Policy; 0 means pass nullptr.

Parameters

launch	description of the launch the call runs in
row_extent	rows of C
col_extent	columns of C
aliased	B and C overlap; the call then takes no staged path, so this returns 0

Returns: element count, or 0 when no staged path applies

Definition at line 430 of file group_blas_trmm.hh.