Overview

Headers installed because public headers include them.

Not a stable interface.

Functions

◆ createBackendHandle()

template<typename T, MatrixFormat MType>

std::shared_ptr< BackendMatrixHandle< T, MType > > batchlas::createBackendHandle ( const MatrixView< T, MType > & view )

Creates the backend descriptor for a view (src/backends/matrix_handle_impl.cc).

◆ ormqr_blocked()

template<Backend B, typename T>

Event batchlas::ormqr_blocked	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	a,
		const MatrixView< T, MatrixFormat::Dense > &	c,
		Side	side,
		Transpose	trans,
		Span< T >	tau,
		Span< std::byte >	workspace,
		int32_t	block_size = tuning::ORMQR_BLOCK_SIZE_MEDIUM )

Blocked native arm of ormqr(): applies Q with the compact WY form.

Groups the geqrf() reflectors (forward, columnwise, unit-lower V) into panels of block_size and applies each as \( I - V T V^H \) (LAPACK larft/larfb) with level-3 GEMMs. Same result as ormqr(); batched through strided-batch views. Meant for medium sizes, where the CTA kernels do not apply.

Precondition: ctx is in-order (the pack / larft / GEMM steps are not separately ordered); workspace holds at least ormqr_blocked_buffer_size() bytes (V, T and W)

Exceptions

batchlas::invalid_argument	on mismatched batch or order, a short `tau`, or an out-of-order `ctx`
batchlas::unsupported	for Transpose::Trans with complex T (use ConjTrans)

◆ ormqr_blocked_buffer_size()

template<Backend B, typename T>

size_t batchlas::ormqr_blocked_buffer_size	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	a,
		const MatrixView< T, MatrixFormat::Dense > &	c,
		Side	side,
		Transpose	trans,
		Span< T >	tau,
		int32_t	block_size = tuning::ORMQR_BLOCK_SIZE_MEDIUM )

Workspace, in bytes, that ormqr_blocked() needs for this block_size.

◆ group_broadcast()

template<typename Group, typename T, typename... Id>
requires sycl::is_group_v<Group>

T batchlas::portable::group_broadcast	(	Group	g,
		T	x,
		Id...	id )

inline

sycl::group_broadcast for any trivially copyable T. Collective.

Definition at line 151 of file group-collectives.hh.

◆ select_from_group()

template<typename Group, typename T, typename Id>
requires sycl::is_group_v<Group>

T batchlas::portable::select_from_group	(	Group	g,
		T	x,
		Id	id )

inline

sycl::select_from_group for any trivially copyable T. Collective.

Definition at line 162 of file group-collectives.hh.

◆ shift_group_left()

template<typename Group, typename T, typename... Delta>
requires sycl::is_group_v<Group>

T batchlas::portable::shift_group_left	(	Group	g,
		T	x,
		Delta...	delta )

inline

sycl::shift_group_left for any trivially copyable T. Collective.

Definition at line 173 of file group-collectives.hh.

◆ shift_group_right()

template<typename Group, typename T, typename... Delta>
requires sycl::is_group_v<Group>

T batchlas::portable::shift_group_right	(	Group	g,
		T	x,
		Delta...	delta )

inline

sycl::shift_group_right for any trivially copyable T. Collective.

Definition at line 184 of file group-collectives.hh.

◆ permute_group_by_xor()

template<typename Group, typename T, typename Mask>
requires sycl::is_group_v<Group>

T batchlas::portable::permute_group_by_xor	(	Group	g,
		T	x,
		Mask	mask )

inline

sycl::permute_group_by_xor for any trivially copyable T. Collective.

Definition at line 195 of file group-collectives.hh.

◆ reduce_over_group()

template<typename Group, typename T, typename Op>
requires sycl::is_group_v<Group>

T batchlas::portable::reduce_over_group	(	Group	g,
		T	x,
		Op	op )

inline

sycl::reduce_over_group, also for a std::complex sum; a work-group floating result is exact (see detail::wg_reduce_floating).

Collective.

Definition at line 207 of file group-collectives.hh.

◆ joint_reduce() [1/2]

template<typename Group, typename Ptr, typename T, typename Op>
requires sycl::is_group_v<Group>

T batchlas::portable::joint_reduce	(	Group	g,
		Ptr	first,
		Ptr	last,
		T	init,
		Op	op )

inline

sycl::joint_reduce with an initial value, also for a std::complex sum.

Collective.

Precondition: first and last are raw pointers: AdaptiveCpp has no accessor-iterator overload.

Definition at line 232 of file group-collectives.hh.

◆ joint_reduce() [2/2]

template<typename Group, typename Ptr, typename Op>
requires sycl::is_group_v<Group>

auto batchlas::portable::joint_reduce	(	Group	g,
		Ptr	first,
		Ptr	last,
		Op	op )

inline

sycl::joint_reduce without an initial value (an empty range gives T{}).

Collective.

Precondition: first and last are raw pointers.

Definition at line 253 of file group-collectives.hh.

◆ joint_exclusive_scan()

template<typename Group, typename InPtr, typename OutPtr, typename T, typename Op>
requires sycl::is_group_v<Group>

OutPtr batchlas::portable::joint_exclusive_scan	(	Group	g,
		InPtr	first,
		InPtr	last,
		OutPtr	result,
		T	init,
		Op	op )

inline

sycl::joint_exclusive_scan with an initial value; result may equal first.

Collective.

Precondition: first, last and result are raw pointers; T is an integral type.

Definition at line 271 of file group-collectives.hh.

◆ invoke_one()

template<typename Group, typename Fn, typename... Args>

void batchlas::invoke_one	(	const Group &	group,
		Fn &&	fn,
		Args &&...	args )

inlineconstexpr

Calls fn(args...) on the group's leader work-item only.

No barrier and no broadcast: other work-items do not see side effects until the caller synchronises.

Template Parameters

Group any type with leader(): a SYCL group, sub-group, or sub-group partition

Definition at line 49 of file group-invoke.hh.

◆ broadcast_from_leader()

template<typename Group, typename T>

T batchlas::broadcast_from_leader	(	const Group &	group,
		T	value )

inlineconstexpr

Returns the leader's value on every work-item of group.

Collective.

Uses the group's sg_leader_broadcast when one is found by ADL (sub-group partitions), else batchlas::portable::group_broadcast (group-collectives.hh).

Precondition: T is trivially copyable.

Definition at line 61 of file group-invoke.hh.

◆ invoke_one_broadcast()

template<typename Group, typename Fn, typename... Args>

auto batchlas::invoke_one_broadcast	(	const Group &	group,
		Fn &&	fn,
		Args &&...	args ) -> std::invoke_result_t< Fn, Args... >

inlineconstexpr

Calls fn(args...) on the leader and returns its result on every work-item.

Collective.

Returns: the leader's result; non-leaders contribute a value-initialised placeholder that is discarded

Precondition: the result type is trivially copyable.

Definition at line 71 of file group-invoke.hh.

◆ compute_optimal_wg_size()

size_t batchlas::compute_optimal_wg_size	(	const Device &	device,
		KernelType	kernel_type,
		size_t	problem_size = 0,
		size_t	batch_size = 1,
		size_t	memory_per_problem = 0 )

inline

Compute a work-group size from the device type and kernel type.

Starts from a per-device-type base (GPU 256, CPU min(64, 2 x compute units), accelerator 128), caps it per kernel type, then by problem_size (when nonzero) and the device's maximum work-group size. Never returns 0.

Parameters

device	The target device
kernel_type	Type of kernel operation
problem_size	Characteristic problem size (e.g., matrix dimension); 0 = no cap
batch_size	Number of matrices in batch; currently unused
memory_per_problem	Currently unused

Returns: Work-group size

Trap: For REDUCTION and SCAN the "power of two" rounding is 1 << (31 - __builtin_clzl(x)), which assumes a 32-bit long. On LP64 the shift count is negative (undefined behaviour), so the result is NOT guaranteed to be a power of two.

Definition at line 65 of file kernel-heuristics.hh.

◆ compute_batched_nd_range_sizes()

BatchedNdRangeSizes batchlas::compute_batched_nd_range_sizes	(	size_t	total_work,
		const Device &	device,
		KernelType	kernel_type,
		size_t	batch_size,
		size_t	problem_size,
		size_t	preferred_wg_size = 0,
		size_t	footprint_per_problem = 0,
		size_t	max_wg_size_for_kernel = 0 )

inline

Compute nd_range sizes for a batched operation.

When total_work exceeds INT32_MAX / 2 the result grid-strides with four work-groups per compute unit. Otherwise, with at least as many matrices as compute units, each matrix gets ceil(problem_size / local_size) work-groups; with fewer, the compute units are spread over the matrices, capped at what one matrix needs. The global size is never below total_work rounded up.

Parameters

total_work	Total number of work items needed
device	Target device
kernel_type	Type of kernel operation
batch_size	Number of matrices in batch
problem_size	Elements (or work items) per matrix
preferred_wg_size	Optional preferred work-group size (0 = auto)
footprint_per_problem	Bytes per problem; currently has no effect (its estimate is overwritten)
max_wg_size_for_kernel	Currently unused

Returns: BatchedNdRangeSizes {global_size, local_size, use_grid_stride}

Definition at line 169 of file kernel-heuristics.hh.

◆ compute_batched_matrix_decomposition()

BatchedMatrixDecomposition batchlas::compute_batched_matrix_decomposition	(	size_t	batch_size,
		size_t	elements_per_matrix,
		const Device &	device,
		KernelType	kernel_type,
		size_t	preferred_wg_size = 0 )

inline

Compute batched matrix decomposition for work-group assignment.

With at least as many matrices as compute units, one work-group per matrix, and the kernel grid-strides over the rest of the matrix. Otherwise work_groups_per_matrix = ceil(CUs / batch_size), never more than ceil(elements_per_matrix / local_size). The global size is then capped to fit a signed 32-bit int, keeping at least one work-group per matrix; the kernel's grid-stride picks up the slack.

Parameters

batch_size	Number of matrices in batch
elements_per_matrix	Number of elements per matrix
device	Target device
kernel_type	Type of kernel operation
preferred_wg_size	Optional preferred work-group size (0 = auto)

Returns: BatchedMatrixDecomposition {global_size, local_size, work_groups_per_matrix}

Definition at line 259 of file kernel-heuristics.hh.

◆ compute_nd_range_sizes()

std::pair< size_t, size_t > batchlas::compute_nd_range_sizes	(	size_t	total_work,
		const Device &	device,
		KernelType	kernel_type,
		size_t	preferred_wg_size = 0 )

inline

Compute optimal global and local work sizes for nd_range kernels with overflow protection.

Parameters

total_work	Total number of work items needed
device	Target device
kernel_type	Type of kernel operation
preferred_wg_size	Optional preferred work-group size (0 = auto)

Returns: std::pair<size_t, size_t> {global_size, local_size}

Definition at line 307 of file kernel-heuristics.hh.

◆ MiniAccMain()

int miniacc::MiniAccMain	(	int	argc,
		char **	argv )

inline

Parses the command line and runs every registered case that matches; the body of MINI_ACC_MAIN().

Returns: 0, including after --help

Definition at line 719 of file miniacc.hh.

◆ get_raw_ptr()

template<typename T, int Dims = 1>

T * batchlas::util::get_raw_ptr ( const sycl::local_accessor< T, Dims > & accessor )

inline

Helper function to get a raw pointer from a SYCL local accessor.

This function provides a convenient wrapper around the verbose local_accessor::get_multi_ptr<sycl::access::decorated::no>().get() syntax, avoiding the deprecation warning for get_pointer().

Template Parameters

T	The element type of the local accessor
Dims	The dimensionality of the accessor

Parameters

accessor The local accessor to get a pointer from

Returns: T* Raw pointer to the local accessor's storage

Definition at line 31 of file sycl-local-accessor-helpers.hh.

Enumerations

◆ KernelType

enum class batchlas::KernelType

strong

Kernel types for heuristic selection.

Enumerator
ELEMENTWISE	Element-wise operations (fill, copy, etc.).
REDUCTION	Reduction operations (sum, max, etc.).
SCAN	Prefix scan operations.
MEMORY_BOUND	Memory bandwidth limited.
COMPUTE_BOUND	Compute intensive.
SPARSE	Sparse matrix operations.
GEMM	General matrix multiply.
SMALL_MATRIX	Small matrix operations.
TASK_BASED	Task-based parallelism, e.g. no fine-grained parallelism, each thread solves one problem.

Definition at line 22 of file kernel-heuristics.hh.

Files
file	include/batchlas/internal/ormqr_blocked.hh
	Blocked (WY) application of Q from geqrf reflectors.
file	include/batchlas/util/group-collectives.hh
	batchlas::portable:: group broadcast, shuffles and sums for any trivially copyable or complex T.
file	include/batchlas/util/group-invoke.hh
	Leader-only invocation and leader broadcast for SYCL groups and sub-group partitions.
file	include/batchlas/util/kernel-heuristics.hh
	Launch-size heuristics used by the library's own kernels.
file	include/batchlas/util/miniacc.hh
	miniacc: a header-only accuracy-benchmark harness (register cases, sample, summarise, CSV).
file	include/batchlas/util/sycl-local-accessor-helpers.hh
	util::get_raw_ptr() for sycl::local_accessor, used to pass device-BLAS workspaces.

Classes
class	batchlas::BackendMatrixHandle< T, MType >
	Backend-library descriptor for a matrix (cuSPARSE / rocSPARSE handles); defined in src/. More...
struct	batchlas::BatchedNdRangeSizes
	Result of compute_batched_nd_range_sizes(). More...
struct	batchlas::BatchedMatrixDecomposition
	Result of compute_batched_matrix_decomposition(). More...
struct	miniacc::Config
	Run-wide settings parsed from the command line. More...
class	miniacc::State
	Per-case context handed to a benchmark: its arguments, the run settings, and the sample sink. More...
struct	miniacc::Benchmark
	A registered case: name, function, and its argument lists (Args, ArgRange, ArgsProduct). More...

Overview

Functions

◆ createBackendHandle()

◆ ormqr_blocked()

◆ ormqr_blocked_buffer_size()

◆ group_broadcast()

◆ select_from_group()

◆ shift_group_left()

◆ shift_group_right()

◆ permute_group_by_xor()

◆ reduce_over_group()

◆ joint_reduce() [1/2]

◆ joint_reduce() [2/2]

◆ joint_exclusive_scan()

◆ invoke_one()

◆ broadcast_from_leader()

◆ invoke_one_broadcast()

◆ compute_optimal_wg_size()

◆ compute_batched_nd_range_sizes()

◆ compute_batched_matrix_decomposition()

◆ compute_nd_range_sizes()

◆ MiniAccMain()

◆ get_raw_ptr()

Enumerations

◆ KernelType

Files

Classes