Overview

In-kernel work-group and sub-group BLAS templates (#include <batchlas/blas/device.hh>, needs -fsycl).

Lower-level functions (algorithm-specific variants, workspace queries, validation and helpers): Device-side group BLAS: lower-level API.

Functions

◆ fill() [1/2]

template<typename Group, typename T>

void batchlas::device::fill	(	const Group &	group,
		const VectorView< T > &	x,
		const T &	value )

inlineconstexpr

Sets every element of x to value, cooperatively across group.

Template Parameters

Group sycl::group, sycl::sub_group or an nd_item; every work-item must call

Parameters

group	executor whose work-items share the elements
x	single vector (batch_size() == 1), written in place
value	fill value

Definition at line 16 of file group_blas_fill.hh.

◆ fill() [2/2]

template<typename Group, typename T>

void batchlas::device::fill	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const T &	value )

inlineconstexpr

Sets every element of the dense matrix a to value; the padding rows between a.rows() and a.ld() are not touched.

Parameters

group	executor whose work-items share the elements
a	single matrix, written in place
value	fill value

Definition at line 33 of file group_blas_fill.hh.

◆ gemm() [1/4]

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

void batchlas::device::gemm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

General matrix-matrix product \( C := \alpha\,\mathrm{op}(A)\,\mathrm{op}(B) + \beta\,C \).

With \(\mathrm{op}(A)\) m x k and \(\mathrm{op}(B)\) k x n, C is m x n. Path selection, first eligible wins: the aligned NN kernel (float, a workspace, 256 work-items, m a multiple of 128, n of 64 and k of 32, 4-element-aligned A and B), the register-tiled kernel (float, a workspace, 256 work-items, m >= 4, n >= 8, k >= 32), the sub-group kernel (float, sub-group size admitted by the policy, m >= 8 and n, k at least the sub-group size), then the generic group reduction. The first three need an nd_item executor.

Template Parameters

TransAV	op() applied to A
TransBV	op() applied to B
Group	sycl::group, sycl::sub_group, sycl::nd_item<1> or sycl::nd_item<3>; every work-item must call

Parameters

group	executor. With nd_item<3>, work-group ids 1 and 2 tile the output, and without a fast path only tile-group (0, 0) computes.
a	single matrix, m x k (NoTrans) or k x m
operand	B, C (m x n, must not alias A or B), alpha, beta
workspace	local memory of gemm_workspace_elements() elements, or nullptr

Precondition: C holds finite values even when beta == 0; conforming shapes (checked by assert only).

Note: On the generic path each C element is written by the group leader; barrier before other work-items read C.

See also: Device-side group BLAS

Definition at line 446 of file group_blas_gemm.hh.

◆ gemm() [2/4]

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

void batchlas::device::gemm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As gemm(), with an explicit DeviceBlasPolicy.

Definition at line 459 of file group_blas_gemm.hh.

◆ gemm() [3/4]

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

void batchlas::device::gemm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As gemm(), with the operands passed separately instead of in an operand struct.

Definition at line 514 of file group_blas_gemm.hh.

◆ gemm() [4/4]

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

void batchlas::device::gemm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As gemm(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 530 of file group_blas_gemm.hh.

◆ gemv() [1/4]

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

void batchlas::device::gemv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixVectorOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

General matrix-vector product \( y := \alpha\,\mathrm{op}(A)\,x + \beta\,y \).

With a workspace, NoTrans, a non-Generic policy, an executor of at least 16 work-items and both extents at least the tile (16 or 32), A is staged through local memory tile by tile; otherwise each output element is a group reduction (or one work-item's loop when the inner extent is smaller than the group).

Template Parameters

TransV	op() applied to A
Group	sycl::group, sycl::sub_group or an nd_item (unwrapped to its work-group); every work-item must call

Parameters

group	executor
a	single m x n matrix
operand	x (length of op(A)'s columns), y (length of op(A)'s rows), alpha, beta
workspace	local memory of gemv_workspace_elements() elements, or nullptr

Precondition: y holds finite values even when beta == 0; y does not alias x or A.

Note: Issue a group barrier before other work-items read y.

See also: Device-side group BLAS

Definition at line 212 of file group_blas_gemv.hh.

◆ gemv() [2/4]

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

void batchlas::device::gemv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixVectorOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As gemv(), with an explicit DeviceBlasPolicy.

Definition at line 224 of file group_blas_gemv.hh.

◆ gemv() [3/4]

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

void batchlas::device::gemv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As gemv(), with the operands passed separately instead of in an operand struct.

Definition at line 259 of file group_blas_gemv.hh.

◆ gemv() [4/4]

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

void batchlas::device::gemv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As gemv(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 274 of file group_blas_gemv.hh.

◆ ger() [1/4]

template<bool ConjugateXV = false, bool ConjugateYV = false, typename Group, typename T>

void batchlas::device::ger	(	const Group &	group,
		const VectorView< T > &	x,
		Rank1UpdateOperand< T >	operand )

inlineconstexpr

Rank-1 update \( A := A + \alpha\,\tilde{x}\,\tilde{y}^{T} \), with \(\tilde{x}\), \(\tilde{y}\) optionally conjugated.

geru() and gerc() are the BLAS-named spellings (no conjugation, and conjugated y). Each element of A is owned by one work-item, so no reduction and no barrier is involved.

Template Parameters

ConjugateXV	use \(\bar{x}\)
ConjugateYV	use \(\bar{y}\)
Group	sycl::group, sycl::sub_group or an nd_item; every work-item must call

Parameters

group	executor whose work-items share the elements of A
x	left vector, length a.rows()
operand	y (length a.cols()), A (updated in place) and alpha

Precondition: single-problem views; sizes as above. Checked by assert only.

Definition at line 61 of file group_blas_ger.hh.

◆ ger() [2/4]

template<DeviceBlasPolicy Policy, bool ConjugateXV = false, bool ConjugateYV = false, typename Group, typename T>

void batchlas::device::ger	(	const Group &	group,
		const VectorView< T > &	x,
		Rank1UpdateOperand< T >	operand )

inlineconstexpr

As ger(), with an explicit DeviceBlasPolicy (ignored: one implementation).

Definition at line 73 of file group_blas_ger.hh.

◆ ger() [3/4]

template<bool ConjugateXV = false, bool ConjugateYV = false, typename Group, typename T>

void batchlas::device::ger	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		T	alpha = T(1) )

inlineconstexpr

As ger(), with the operands passed separately instead of in an operand struct.

Definition at line 82 of file group_blas_ger.hh.

◆ ger() [4/4]

template<DeviceBlasPolicy Policy, bool ConjugateXV = false, bool ConjugateYV = false, typename Group, typename T>

void batchlas::device::ger	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		T	alpha = T(1) )

inlineconstexpr

As ger(), with an explicit DeviceBlasPolicy and the operands passed separately (ignored: one implementation).

Definition at line 96 of file group_blas_ger.hh.

◆ geru() [1/4]

template<typename Group, typename T>

void batchlas::device::geru	(	const Group &	group,
		const VectorView< T > &	x,
		Rank1UpdateOperand< T >	operand )

inlineconstexpr

Unconjugated rank-1 update \( A := A + \alpha\,x\,y^{T} \); ger() without conjugation.

Parameters

group	executor whose work-items share the elements of A
x	left vector, length a.rows()
operand	y, A and alpha

Definition at line 110 of file group_blas_ger.hh.

◆ geru() [2/4]

template<DeviceBlasPolicy Policy, typename Group, typename T>

void batchlas::device::geru	(	const Group &	group,
		const VectorView< T > &	x,
		Rank1UpdateOperand< T >	operand )

inlineconstexpr

As geru(), with an explicit DeviceBlasPolicy (ignored: one implementation).

Definition at line 118 of file group_blas_ger.hh.

◆ geru() [3/4]

template<typename Group, typename T>

void batchlas::device::geru	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		T	alpha = T(1) )

inlineconstexpr

As geru(), with the operands passed separately instead of in an operand struct.

Definition at line 127 of file group_blas_ger.hh.

◆ geru() [4/4]

template<DeviceBlasPolicy Policy, typename Group, typename T>

void batchlas::device::geru	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		T	alpha = T(1) )

inlineconstexpr

As geru(), with an explicit DeviceBlasPolicy and the operands passed separately (ignored: one implementation).

Definition at line 137 of file group_blas_ger.hh.

◆ gerc() [1/4]

template<typename Group, typename T>

void batchlas::device::gerc	(	const Group &	group,
		const VectorView< T > &	x,
		Rank1UpdateOperand< T >	operand )

inlineconstexpr

Conjugated rank-1 update \( A := A + \alpha\,x\,y^{H} \); ger() with conjugated y.

Parameters

group	executor whose work-items share the elements of A
x	left vector, length a.rows()
operand	y, A and alpha

Definition at line 151 of file group_blas_ger.hh.

◆ gerc() [2/4]

template<DeviceBlasPolicy Policy, typename Group, typename T>

void batchlas::device::gerc	(	const Group &	group,
		const VectorView< T > &	x,
		Rank1UpdateOperand< T >	operand )

inlineconstexpr

As gerc(), with an explicit DeviceBlasPolicy (ignored: one implementation).

Definition at line 159 of file group_blas_ger.hh.

◆ gerc() [3/4]

template<typename Group, typename T>

void batchlas::device::gerc	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		T	alpha = T(1) )

inlineconstexpr

As gerc(), with the operands passed separately instead of in an operand struct.

Definition at line 168 of file group_blas_ger.hh.

◆ gerc() [4/4]

template<DeviceBlasPolicy Policy, typename Group, typename T>

void batchlas::device::gerc	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		T	alpha = T(1) )

inlineconstexpr

As gerc(), with an explicit DeviceBlasPolicy and the operands passed separately (ignored: one implementation).

Definition at line 178 of file group_blas_ger.hh.

◆ syrk() [1/4]

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

void batchlas::device::syrk	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		RankKOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

Symmetric rank-k update \( C := \alpha\,A\,A^{T} + \beta\,C \) (NoTrans) or \( C := \alpha\,A^{T} A + \beta\,C \), on the UploV triangle of C only.

The other triangle of C is neither read nor written. Fast paths: register-tiled for float and a tiled kernel for std::complex<float>, both needing an nd_item executor and a workspace; with an nd_item<3> and no fast path only tile-group (0, 0) computes.

Template Parameters

UploV	triangle of C updated
TransV	NoTrans: A is n x k; otherwise A is k x n
Group	sycl::group, sycl::sub_group or an nd_item; every work-item must call

Parameters

group	executor
a	single input matrix
operand	C (n x n), alpha, beta
workspace	local memory of syrk_workspace_elements() elements, or nullptr

Precondition: C holds finite values even when beta == 0; C does not alias A.

Note: Each C element has one writer; barrier before other work-items read C.

See also: herk()

Definition at line 485 of file group_blas_rankk.hh.

◆ syrk() [2/4]

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

void batchlas::device::syrk	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		RankKOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As syrk(), with an explicit DeviceBlasPolicy.

Definition at line 498 of file group_blas_rankk.hh.

◆ syrk() [3/4]

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

void batchlas::device::syrk	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As syrk(), with the operands passed separately instead of in an operand struct.

Definition at line 536 of file group_blas_rankk.hh.

◆ syrk() [4/4]

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

void batchlas::device::syrk	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As syrk(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 551 of file group_blas_rankk.hh.

◆ herk() [1/4]

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

void batchlas::device::herk	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		RankKOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

Hermitian rank-k update \( C := \alpha\,A\,A^{H} + \beta\,C \) (NoTrans) or \( C := \alpha\,A^{H} A + \beta\,C \), on the UploV triangle.

For real T this is exactly syrk(). alpha and beta are T, not real as in BLAS; the imaginary part of every diagonal element of C is set to zero, but a non-real alpha still gives a non-Hermitian off-diagonal.

Parameters

group	executor (see syrk())
a	single input matrix
operand	C, alpha, beta
workspace	local memory of herk_workspace_elements() elements, or nullptr

Definition at line 570 of file group_blas_rankk.hh.

◆ herk() [2/4]

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

void batchlas::device::herk	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		RankKOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As herk(), with an explicit DeviceBlasPolicy.

Definition at line 583 of file group_blas_rankk.hh.

◆ herk() [3/4]

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

void batchlas::device::herk	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As herk(), with the operands passed separately instead of in an operand struct.

Definition at line 621 of file group_blas_rankk.hh.

◆ herk() [4/4]

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

void batchlas::device::herk	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As herk(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 636 of file group_blas_rankk.hh.

◆ syr2k() [1/4]

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

void batchlas::device::syr2k	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

Symmetric rank-2k update \( C := \alpha\,(A B^{T} + B A^{T}) + \beta\,C \) (NoTrans) or \( C := \alpha\,(A^{T} B + B^{T} A) + \beta\,C \), on the UploV triangle.

Paths and executor rules as for syrk().

Template Parameters

UploV	triangle of C updated
TransV	NoTrans: A and B are n x k; otherwise k x n

Parameters

group	executor (see syrk())
a	first single input matrix
operand	B (same shape as A), C (n x n), alpha, beta
workspace	local memory of syr2k_workspace_elements() elements, or nullptr

Precondition: C holds finite values even when beta == 0; C does not alias A or B.

See also: her2k()

Definition at line 657 of file group_blas_rankk.hh.

◆ syr2k() [2/4]

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

void batchlas::device::syr2k	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As syr2k(), with an explicit DeviceBlasPolicy.

Definition at line 670 of file group_blas_rankk.hh.

◆ syr2k() [3/4]

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

void batchlas::device::syr2k	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As syr2k(), with the operands passed separately instead of in an operand struct.

Definition at line 708 of file group_blas_rankk.hh.

◆ syr2k() [4/4]

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

void batchlas::device::syr2k	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As syr2k(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 724 of file group_blas_rankk.hh.

◆ her2k() [1/4]

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

void batchlas::device::her2k	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

Hermitian rank-2k update \( C := \alpha\,A B^{H} + \bar{\alpha}\,B A^{H} + \beta\,C \) (NoTrans) or \( \alpha\,A^{H} B + \bar{\alpha}\,B^{H} A + \beta\,C \), on the UploV triangle.

For real T this is exactly syr2k(). The imaginary part of every diagonal element of C is set to zero; beta is T, not real as in BLAS.

Parameters

group	executor (see syrk())
a	first single input matrix
operand	B, C, alpha, beta
workspace	local memory of her2k_workspace_elements() elements, or nullptr

Definition at line 743 of file group_blas_rankk.hh.

◆ her2k() [2/4]

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

void batchlas::device::her2k	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As her2k(), with an explicit DeviceBlasPolicy.

Definition at line 756 of file group_blas_rankk.hh.

◆ her2k() [3/4]

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

void batchlas::device::her2k	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As her2k(), with the operands passed separately instead of in an operand struct.

Definition at line 794 of file group_blas_rankk.hh.

◆ her2k() [4/4]

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

void batchlas::device::her2k	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As her2k(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 810 of file group_blas_rankk.hh.

◆ symm() [1/4]

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

void batchlas::device::symm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

Symmetric matrix-matrix product \( C := \alpha\,A\,B + \beta\,C \) (Left) or \( C := \alpha\,B\,A + \beta\,C \) (Right), reading only the UploV triangle of A.

There is no device hemm. The fast paths are as for trmm(), except that B and C must not overlap.

Template Parameters

SideV	side A multiplies from
UploV	stored triangle of A
Group	sycl::group, sycl::sub_group or an nd_item; every work-item must call

Parameters

group	executor
a	single square matrix, order B.rows() (Left) or B.cols() (Right)
operand	B, C (same shape as B, not overlapping it), alpha, beta
workspace	local memory of symm_workspace_elements() elements, or nullptr

Precondition: C holds finite values even when beta == 0.

Note: On the generic path each C element is written by the group leader; barrier before other work-items read C.

Definition at line 307 of file group_blas_symm.hh.

◆ symm() [2/4]

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

void batchlas::device::symm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As symm(), with an explicit DeviceBlasPolicy.

Definition at line 320 of file group_blas_symm.hh.

◆ symm() [3/4]

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

void batchlas::device::symm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As symm(), with the operands passed separately instead of in an operand struct.

Definition at line 355 of file group_blas_symm.hh.

◆ symm() [4/4]

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

void batchlas::device::symm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As symm(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 371 of file group_blas_symm.hh.

◆ symv() [1/4]

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

void batchlas::device::symv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixVectorOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

Symmetric matrix-vector product \( y := \alpha\,A\,x + \beta\,y \), reading only the UploV triangle of A.

With a workspace, a non-Generic policy, at least 16 work-items and n >= 16, 16 x 16 blocks of A are staged through local memory and each off-diagonal block is used for both its own and its mirrored contribution.

Template Parameters

UploV	stored triangle of A
Group	sycl::group or sycl::sub_group (the tiled path barriers on it); every work-item must call

Parameters

group	executor
a	single n x n matrix
operand	x, y (each of length at least n), alpha, beta
workspace	local memory of symv_workspace_elements() elements, or nullptr

Precondition: y holds finite values even when beta == 0; y does not alias x.

See also: hemv(), Device-side group BLAS

Definition at line 268 of file group_blas_symv.hh.

◆ symv() [2/4]

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

void batchlas::device::symv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixVectorOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As symv(), with an explicit DeviceBlasPolicy.

Definition at line 277 of file group_blas_symv.hh.

◆ symv() [3/4]

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

void batchlas::device::symv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As symv(), with the operands passed separately instead of in an operand struct.

Definition at line 308 of file group_blas_symv.hh.

◆ symv() [4/4]

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

void batchlas::device::symv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As symv(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 320 of file group_blas_symv.hh.

◆ hemv() [1/4]

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

void batchlas::device::hemv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixVectorOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

Hermitian matrix-vector product \( y := \alpha\,A\,x + \beta\,y \); symv() with conjugated mirroring.

For real T this is exactly symv(). The diagonal is read as stored; its imaginary part is not forced to zero.

Parameters

group	executor (sycl::group or sycl::sub_group)
a	single n x n matrix, UploV triangle read
operand	x, y, alpha, beta
workspace	local memory of hemv_workspace_elements() elements, or nullptr

Definition at line 339 of file group_blas_symv.hh.

◆ hemv() [2/4]

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

void batchlas::device::hemv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixVectorOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As hemv(), with an explicit DeviceBlasPolicy.

Definition at line 348 of file group_blas_symv.hh.

◆ hemv() [3/4]

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

void batchlas::device::hemv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As hemv(), with the operands passed separately instead of in an operand struct.

Definition at line 379 of file group_blas_symv.hh.

◆ hemv() [4/4]

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

void batchlas::device::hemv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As hemv(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 391 of file group_blas_symv.hh.

◆ trmm() [1/4]

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

void batchlas::device::trmm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

Triangular matrix-matrix product \( C := \alpha\,\mathrm{op}(A)\,B + \beta\,C \) (Left) or \( C := \alpha\,B\,\mathrm{op}(A) + \beta\,C \) (Right).

Unlike BLAS trmm it has a separate output and a beta. B and C may overlap: the generic path orders its traversal so no entry is overwritten before it is read, and overlap disables both fast paths. The fast paths (float, an nd_item executor; the register-tiled one also an nd_item<3>, 256 work-items and a workspace) are otherwise as for gemm().

Template Parameters

SideV	side A multiplies from
UploV	stored triangle of A
TransV	op() applied to A
DiagV	Unit: the diagonal is taken as 1
Group	sycl::group, sycl::sub_group or an nd_item; every work-item must call

Parameters

group	executor
a	single square triangular matrix, order B.rows() (Left) or B.cols() (Right)
operand	B, C (same shape as B), alpha, beta
workspace	local memory of trmm_workspace_elements() elements, or nullptr

Precondition: C holds finite values even when beta == 0.

Note: On the generic path each C element is written by the group leader; barrier before other work-items read C.

Definition at line 371 of file group_blas_trmm.hh.

◆ trmm() [2/4]

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

void batchlas::device::trmm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixMatrixOperand< T >	operand,
		T *	workspace = nullptr )

inlineconstexpr

As trmm(), with an explicit DeviceBlasPolicy.

Definition at line 386 of file group_blas_trmm.hh.

◆ trmm() [3/4]

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

void batchlas::device::trmm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As trmm(), with the operands passed separately instead of in an operand struct.

Definition at line 444 of file group_blas_trmm.hh.

◆ trmm() [4/4]

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

void batchlas::device::trmm	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const KernelMatrixView< T, MatrixFormat::Dense > &	b,
		const KernelMatrixView< T, MatrixFormat::Dense > &	c,
		T	alpha = T(1),
		T	beta = T(0),
		T *	workspace = nullptr )

inlineconstexpr

As trmm(), with an explicit DeviceBlasPolicy and the operands passed separately.

Definition at line 462 of file group_blas_trmm.hh.

◆ trmv() [1/4]

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

void batchlas::device::trmv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixVectorOperand< T >	operand )

inlineconstexpr

Triangular matrix-vector product \( y := \alpha\,\mathrm{op}(A)\,x + \beta\,y \), A triangular.

Out of place, unlike BLAS trmv: A is read only on its UploV triangle (and not on the diagonal when DiagV is Unit). x and y may be the same vector: rows are visited in the order that never overwrites an entry a later row still reads.

Template Parameters

UploV	stored triangle of A
TransV	op() applied to A
DiagV	Unit: the diagonal is taken as 1
Group	sycl::group, sycl::sub_group or an nd_item; every work-item must call

Parameters

group	executor
a	single n x n matrix
operand	x, y (each of length at least n), alpha, beta

Precondition: y holds finite values even when beta == 0.

Note: Each y element is written by the group leader; barrier before other work-items read it.

Definition at line 70 of file group_blas_trmv.hh.

◆ trmv() [2/4]

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

void batchlas::device::trmv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		MatrixVectorOperand< T >	operand )

inlineconstexpr

As trmv(), with an explicit DeviceBlasPolicy (ignored: one implementation).

Definition at line 83 of file group_blas_trmv.hh.

◆ trmv() [3/4]

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

void batchlas::device::trmv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1),
		T	beta = T(0) )

inlineconstexpr

As trmv(), with the operands passed separately instead of in an operand struct.

Definition at line 96 of file group_blas_trmv.hh.

◆ trmv() [4/4]

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

void batchlas::device::trmv	(	const Group &	group,
		const KernelMatrixView< T, MatrixFormat::Dense > &	a,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1),
		T	beta = T(0) )

inlineconstexpr

As trmv(), with an explicit DeviceBlasPolicy and the operands passed separately (ignored: one implementation).

Definition at line 112 of file group_blas_trmv.hh.

◆ copy()

template<typename Group, typename T>

void batchlas::device::copy	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y )

inlineconstexpr

\( y := x \), cooperatively across group.

Template Parameters

Group sycl::group, sycl::sub_group or an nd_item; every work-item must call

Parameters

group	executor whose work-items share the elements
x	input vector
y	output vector

Precondition: x.size() == y.size(); both single vectors (batch_size() == 1). Checked by assert only.

Definition at line 34 of file group_blas_vector.hh.

◆ copyc()

template<typename Group, typename T>

void batchlas::device::copyc	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y )

inlineconstexpr

\( y := \bar{x} \) (a plain copy for real T).

Parameters

group	executor whose work-items share the elements
x	input vector
y	output vector

Precondition: x.size() == y.size().

Definition at line 52 of file group_blas_vector.hh.

◆ scal()

template<typename Group, typename T>

void batchlas::device::scal	(	const Group &	group,
		const VectorView< T > &	x,
		T	alpha )

inlineconstexpr

\( x := \alpha\,x \).

Parameters

group	executor whose work-items share the elements
x	vector scaled in place
alpha	scale

Definition at line 69 of file group_blas_vector.hh.

◆ axpy()

template<typename Group, typename T>

void batchlas::device::axpy	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		T	alpha = T(1) )

inlineconstexpr

\( y := \alpha\,x + y \).

Parameters

group	executor whose work-items share the elements
x	input vector
y	vector updated in place
alpha	scale of x

Precondition: x.size() == y.size().

Definition at line 88 of file group_blas_vector.hh.

◆ hadamard() [1/2]

template<typename Group, typename T>

void batchlas::device::hadamard	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y,
		const VectorView< T > &	z )

inlineconstexpr

Element-wise product \( z_i := x_i\,y_i \).

Parameters

group	executor whose work-items share the elements
x	first factor
y	second factor
z	output; may alias x or y

Precondition: all three vectors have the same size.

Definition at line 108 of file group_blas_vector.hh.

◆ hadamard() [2/2]

template<typename Group, typename T, typename Op, typename... Inputs>
requires (sizeof...(Inputs) > 0 && (detail::VectorOperandFor<T, Inputs> && ...))

void batchlas::device::hadamard	(	const Group &	group,
		const VectorView< T > &	z,
		Op	op,
		const Inputs &...	inputs )

inlineconstexpr

Element-wise n-ary map \( z_i := \mathrm{op}(u_i, v_i, \ldots) \).

Template Parameters

Op	callable taking one T per input and returning T; must be device-callable
Inputs	one or more VectorView<T>

Parameters

group	executor whose work-items share the elements
z	output vector
op	the element-wise function
inputs	input vectors, each of z.size()

Definition at line 124 of file group_blas_vector.hh.

◆ dotu()

template<typename Group, typename T>

T batchlas::device::dotu	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y )

inlineconstexpr

Unconjugated dot product \( \sum_i x_i\,y_i \).

Parameters

group	executor whose work-items share the elements
x	first vector
y	second vector

Returns: the group-wide sum, on every work-item of group

Precondition: x.size() == y.size().

Definition at line 139 of file group_blas_vector.hh.

◆ dotc()

template<typename Group, typename T>

T batchlas::device::dotc	(	const Group &	group,
		const VectorView< T > &	x,
		const VectorView< T > &	y )

inlineconstexpr

Conjugated dot product \( \sum_i \bar{x}_i\,y_i \) (dotu for real T).

Parameters

group	executor whose work-items share the elements
x	conjugated vector
y	second vector

Returns: the group-wide sum, on every work-item of group

Precondition: x.size() == y.size().

Definition at line 161 of file group_blas_vector.hh.

Enumerations

◆ DeviceBlasPolicy

enum class batchlas::device::DeviceBlasPolicy

strong

Which implementation family a device BLAS call may use.

Level-1 operations, ger and trmv have one implementation and ignore it.

See also: Device-side group BLAS

Enumerator
Auto	any eligible path; the sub-group fast paths accept sub-group size 16 or 32
Generic	the portable group-reduction path only; never stages through local memory
Subgroup16	fast paths only when the actual sub-group size is 16
Subgroup32	fast paths only when the actual sub-group size is 32

Definition at line 28 of file group_blas_common.hh.

◆ DeviceBlasLaunchKind

enum class batchlas::device::DeviceBlasLaunchKind

strong

Kind of launch a DeviceBlasLaunchInfo describes.

Enumerator
Group	the call receives a sycl::group (or sycl::sub_group)
NdItem1D	the call receives a sycl::nd_item<1>
NdItem3D	the call receives a sycl::nd_item<3> whose dimensions 1 and 2 index output tiles

Definition at line 36 of file group_blas_common.hh.

Topics
	Device-side group BLAS: lower-level API
	Launch-description and operand builders, and workspace sizes.

Files
file	include/batchlas/blas/device.hh
	Umbrella header for the in-kernel group BLAS (namespace batchlas::device).
file	include/batchlas/blas/device/group_blas.hh
	All device group BLAS operations; include <batchlas/blas/device.hh> instead.

Classes
struct	batchlas::device::DeviceBlasLaunchInfo
	Host-side description of a launch, for the *_workspace_elements queries. More...
struct	batchlas::device::MatrixVectorOperand< T >
	Operands of a matrix-vector operation: \( y := \alpha\,\mathrm{op}(A)\,x + \beta\,y \). More...
struct	batchlas::device::MatrixMatrixOperand< T >
	Operands of a matrix-matrix operation: \( C := \alpha\,(\ldots B \ldots) + \beta\,C \). More...
struct	batchlas::device::RankKOperand< T >
	Operands of a rank-k update: \( C := \alpha\,\mathrm{op}(A)\,\mathrm{op}(A)^{T\|H} + \beta\,C \). More...
struct	batchlas::device::Rank1UpdateOperand< T >
	Operands of a rank-1 update: \( A := A + \alpha\,x\,y^{T} \) (conjugations per call). More...
struct	batchlas::device::MatrixVectorTransform
	Runtime transpose mode of a matrix-vector operation. More...
struct	batchlas::device::MatrixVectorTransformTag< TransV >
	Compile-time transpose mode of a matrix-vector operation. More...
struct	batchlas::device::GeneralMatrixTransform
	Runtime transpose modes of a general matrix-matrix product. More...
struct	batchlas::device::GeneralMatrixTransformTag< TransAV, TransBV >
	Compile-time transpose modes of a general matrix-matrix product. More...
struct	batchlas::device::TriangularTransform
	Mode of a triangular product (trmv, trmm). More...
struct	batchlas::device::SymmetricTransform
	Mode of a symmetric or Hermitian product (symv, hemv, symm). More...
struct	batchlas::device::OuterProductTransform
	Conjugation mode of a rank-1 update (ger, geru, gerc). More...
struct	batchlas::device::OuterProductTransformTag< ConjugateXV, ConjugateYV >
	Compile-time conjugation mode of a rank-1 update. More...
struct	batchlas::device::SymmetricRank2kTransform
	Mode of a symmetric or Hermitian rank-2k update (syr2k, her2k). More...
struct	batchlas::device::SymmetricRankKTransform
	Mode of a symmetric or Hermitian rank-k update (syrk, herk). More...

Overview

Functions

◆ fill() [1/2]

◆ fill() [2/2]

◆ gemm() [1/4]

◆ gemm() [2/4]

◆ gemm() [3/4]

◆ gemm() [4/4]

◆ gemv() [1/4]

◆ gemv() [2/4]

◆ gemv() [3/4]

◆ gemv() [4/4]

◆ ger() [1/4]

◆ ger() [2/4]

◆ ger() [3/4]

◆ ger() [4/4]

◆ geru() [1/4]

◆ geru() [2/4]

◆ geru() [3/4]

◆ geru() [4/4]

◆ gerc() [1/4]

◆ gerc() [2/4]

◆ gerc() [3/4]

◆ gerc() [4/4]

◆ syrk() [1/4]

◆ syrk() [2/4]

◆ syrk() [3/4]

◆ syrk() [4/4]

◆ herk() [1/4]

◆ herk() [2/4]

◆ herk() [3/4]

◆ herk() [4/4]

◆ syr2k() [1/4]

◆ syr2k() [2/4]

◆ syr2k() [3/4]

◆ syr2k() [4/4]

◆ her2k() [1/4]

◆ her2k() [2/4]

◆ her2k() [3/4]

◆ her2k() [4/4]

◆ symm() [1/4]

◆ symm() [2/4]

◆ symm() [3/4]

◆ symm() [4/4]

◆ symv() [1/4]

◆ symv() [2/4]

◆ symv() [3/4]

◆ symv() [4/4]

◆ hemv() [1/4]

◆ hemv() [2/4]

◆ hemv() [3/4]

◆ hemv() [4/4]

◆ trmm() [1/4]

◆ trmm() [2/4]

◆ trmm() [3/4]

◆ trmm() [4/4]

◆ trmv() [1/4]

◆ trmv() [2/4]

◆ trmv() [3/4]

◆ trmv() [4/4]

◆ copy()

◆ copyc()

◆ scal()

◆ axpy()

◆ hadamard() [1/2]

◆ hadamard() [2/2]

◆ dotu()

◆ dotc()

Enumerations

◆ DeviceBlasPolicy

◆ DeviceBlasLaunchKind

Topics

Files

Classes