queue-dispatch.hh
Go to the documentation of this file.
1#pragma once
2
3#include <complex>
4#include <stdexcept>
5#include <string>
6#include <type_traits>
7#include <utility>
8
9#include <batchlas/backend_config.h>
11#include <batchlas/settings.hh>
16
17/// @file
18/// @brief Backend deduction from the Queue and the USM pointer check.
19/// @ingroup api_dispatch
20
21namespace batchlas {
22
23/**
24 * @addtogroup api_dispatch
25 * @details
26 * <b>Every entry point has two spellings.</b> `f<Backend::CUDA>(ctx, ...)` fixes the
27 * backend at compile time; `f(ctx, ...)` takes it from `ctx.backend()` through
28 * with_backend(). The deducing overload of a positional entry point is generated by
29 * `BATCHLAS_DISPATCH_ON_QUEUE(f)`, and the option-struct overloads in
30 * `blas/options.hh` deduce the same way. Library code templated on `Backend` must
31 * call the `f<B>` spelling, never the deducing one.
32 *
33 * <b>Every entry point also accepts owning arguments.</b> A `Matrix` is a `MatrixView`
34 * and a `Vector` is a `VectorView` (public inheritance), so an owning argument binds to
35 * a view parameter directly and deduces `T` and `MatrixFormat` through derived-to-base
36 * deduction, in any position and with any template arguments named.
37 *
38 * <b>Pointer checks.</b> The deducing spellings check that every non-empty pointer
39 * argument is reachable from the queue's device and throw `std::invalid_argument`
40 * naming the argument otherwise; the `f<Backend>` spellings do not check.
41 * `BATCHLAS_SKIP_POINTER_CHECKS` turns the check off (latched once per process).
42 *
43 * The generating macros are omitted from this reference; see @ref design_api_conventions
44 * for their design and traps.
45 * @{
46 */
47
48namespace detail {
49
50// Element type of a container argument (Matrix, MatrixView, Vector, VectorView, Span,
51// UnifiedVector); void for anything else. Scalars are deliberately not containers: a
52// float op may take a double knob (fill_graded_hermitian's condition number).
53template <class A> struct container_elem { using type = void; };
54// A one-parameter class template over the scalar (StedcParams<T>, GemmOptions<T>) names the
55// op's type as well: stedc_buffer_size(ctx, n, batch, job, StedcParams<double>) has no container.
56template <template <class> class P, class T> struct container_elem<P<T>> { using type = T; };
57template <class T> struct container_elem<std::complex<T>> { using type = void; };
58template <class T, MatrixFormat F> struct container_elem<Matrix<T, F>> { using type = T; };
59template <class T, MatrixFormat F> struct container_elem<MatrixView<T, F>> { using type = T; };
60template <class T> struct container_elem<Vector<T>> { using type = T; };
61template <class T> struct container_elem<VectorView<T>> { using type = T; };
62template <class T> struct container_elem<Span<T>> { using type = T; };
63template <class T> struct container_elem<UnifiedVector<T>> { using type = T; };
64
65template <class T>
66inline constexpr bool is_fp64_scalar_v =
67 std::is_same_v<std::remove_cv_t<T>, double> || std::is_same_v<std::remove_cv_t<T>, std::complex<double>>;
68
69// True when any container argument holds double or complex<double>: Backend::SYCL
70// instantiates those types only when BATCHLAS_HAS_DEVICE_FP64.
71template <class... A>
72inline constexpr bool needs_fp64_v =
74
75[[noreturn]] inline void throw_no_device_fp64(const Queue& ctx, const char* op) {
76 const std::string what = std::string("BatchLAS: ") + (op ? op : "this call") +
77 " was called with double or complex<double> data on a queue whose device (" +
78 ctx.device().get_name() + ")";
79 // A CPU queue pinned to Backend::SYCL has fp64; what it lacks is the build's SYCL double kernels.
80 if (ctx.device().supports_fp64()) {
82 what + " runs Backend::SYCL, and this build compiled no Backend::SYCL kernels for those "
83 "types. Use Backend::NETLIB (the default backend of a CPU Queue), or convert the data "
84 "to float or complex<float>.");
85 }
87 what + " has no fp64 support, so this build compiled no Backend::SYCL kernels for those "
88 "types. Run double work on a CPU Queue (Queue(Device(\"cpu\"))), whose backend is "
89 "NETLIB, or convert the data to float or complex<float>.");
90}
91
92/// Throws throw_no_device_fp64 for a double or complex<double> @p T on a GPU without fp64. For
93/// kernels that are not Backend-templated (or NETLIB ones that launch on the queue's device),
94/// whose double instantiations exist and would otherwise fail in the device JIT.
95template <class T>
96inline void require_device_fp64(const Queue& ctx, const char* op) {
97 if constexpr (is_fp64_scalar_v<T> && !BATCHLAS_HAS_DEVICE_FP64) {
98 if (ctx.device().type == DeviceType::GPU && !ctx.device().supports_fp64())
99 throw_no_device_fp64(ctx, op);
100 }
101}
102
103// with_backend with the fp64 gate: the SYCL branch is discarded (never instantiated) when
104// NeedsFp64 and the build has no device fp64, so no <SYCL, double> symbol is referenced.
105template <bool NeedsFp64, typename F>
106inline auto with_backend_gated(Queue& ctx, const char* op, F&& f) {
107 static_assert(BATCHLAS_HAS_CUDA_BACKEND || BATCHLAS_HAS_ROCM_BACKEND ||
108 BATCHLAS_HAS_SYCL_BACKEND || BATCHLAS_HAS_HOST_BACKEND,
109 "BatchLAS was built with no backends; nothing can be dispatched.");
110
111 switch (ctx.backend()) {
112#if BATCHLAS_HAS_CUDA_BACKEND
113 case Backend::CUDA:
114 return f(std::integral_constant<Backend, Backend::CUDA>{});
115#endif
116#if BATCHLAS_HAS_ROCM_BACKEND
117 case Backend::ROCM:
118 return f(std::integral_constant<Backend, Backend::ROCM>{});
119#endif
120#if BATCHLAS_HAS_SYCL_BACKEND
121 case Backend::SYCL:
122 if constexpr (NeedsFp64 && !BATCHLAS_HAS_DEVICE_FP64) {
123 throw_no_device_fp64(ctx, op);
124 } else {
125 return f(std::integral_constant<Backend, Backend::SYCL>{});
126 }
127#endif
128#if BATCHLAS_HAS_HOST_BACKEND
129 case Backend::NETLIB:
130 return f(std::integral_constant<Backend, Backend::NETLIB>{});
131#endif
132 default:
133 break;
134 }
136 std::string("BatchLAS: backend ") + std::string(to_string(ctx.backend())) +
137 " has no implementation in this build. "
138 "Check Queue::backend_available() before pinning a backend.");
139}
140
141} // namespace detail
142
143/// @brief Call @p f with the queue's backend as a compile-time constant.
144///
145/// Converts `ctx.backend()` into `std::integral_constant<Backend, B>` and invokes
146/// `f` with it, so the body can name `B.value` as a template argument:
147/// @code
148/// with_backend(ctx, [&](auto B) { return gemm<B.value, T>(ctx, ...); });
149/// @endcode
150/// This is a `switch` over existing instantiations, not a virtual call.
151/// @tparam F callable taking `std::integral_constant<Backend, B>`
152/// @param ctx queue whose backend is dispatched on
153/// @param f callable; every branch must return the same type
154/// @return whatever @p f returns
155/// @throws batchlas::unsupported if `ctx.backend()` has no implementation in this
156/// build (always the case for Backend::MAGMA)
157/// @note In a build without device fp64 (`BATCHLAS_HAS_DEVICE_FP64` 0), Backend::SYCL
158/// has no double or complex<double> instantiations: @p f must not name them for
159/// Backend::SYCL. The entry points' deducing overloads check this and throw
160/// batchlas::unsupported instead.
161/// @see Queue::backend_available
162template <typename F>
163inline auto with_backend(Queue& ctx, F&& f) {
164 return detail::with_backend_gated<false>(ctx, nullptr, std::forward<F>(f));
165}
166
167/// @brief with_backend() for a call on element type @p T: throws batchlas::unsupported
168/// before dispatch when @p T is double or complex<double> and the queue's backend is
169/// Backend::SYCL in a build without device fp64.
170template <typename T, typename F>
171inline auto with_backend_for(Queue& ctx, const char* op, F&& f) {
172 return detail::with_backend_gated<detail::is_fp64_scalar_v<T>>(ctx, op, std::forward<F>(f));
173}
174
175/// @}
176
177namespace detail {
178
179// USM contract, enforced: host memory on a GPU queue throws invalid_argument here
180// instead of aborting in the runtime. Latched once (it is on every call's path), so
181// a settings reload does not reach it. Any value of BATCHLAS_SKIP_POINTER_CHECKS not
182// starting with '0' disables it ("=off" too); do not tighten that to env_truthy.
183// evidence: docs/design/api-conventions.md#api-conventions-the-usm-pointer-check
185 static const bool enabled = !batchlas::settings().unsafe.skip_pointer_checks;
186 return enabled;
187}
188
189// Matrix/MatrixView spell it data_ptr(); Span/UnifiedVector spell it data().
190template <typename A>
191concept HasDataPtr = requires(const A& a) { { a.data_ptr() } -> std::convertible_to<const void*>; };
192template <typename A>
193concept HasData = requires(const A& a) {
194 { a.data() } -> std::convertible_to<const void*>;
195 { a.size() } -> std::convertible_to<size_t>;
196};
197
198// Extents: MatrixView rows()/cols()/batch_size(); VectorView size()/batch_size().
199template <typename A>
200concept HasMatrixExtent = requires(const A& a) {
201 { a.rows() } -> std::convertible_to<long long>;
202 { a.cols() } -> std::convertible_to<long long>;
203 { a.batch_size() } -> std::convertible_to<long long>;
204};
205template <typename A>
206concept HasVectorExtent = requires(const A& a) {
207 { a.size() } -> std::convertible_to<long long>;
208 { a.batch_size() } -> std::convertible_to<long long>;
209};
210
211// True when no kernel can dereference the argument, whatever its pointer.
212template <typename A>
213inline bool addresses_no_elements(const A& arg) {
214 using U = std::remove_cvref_t<A>;
215 if constexpr (HasMatrixExtent<U>) {
216 return arg.rows() == 0 || arg.cols() == 0 || arg.batch_size() == 0;
217 } else if constexpr (HasVectorExtent<U>) {
218 return arg.size() == 0 || arg.batch_size() == 0;
219 } else {
220 return false;
221 }
222}
223
224template <typename A>
225inline void require_arg_accessible(const Queue& ctx, const A& arg, const std::string& what) {
226 if constexpr (HasDataPtr<std::remove_cvref_t<A>>) {
227 // A default-constructed view means "optional matrix not in use"; it must pass.
228 if (addresses_no_elements(arg)) return;
229 ctx.require_device_accessible(static_cast<const void*>(arg.data_ptr()), what.c_str());
230 } else if constexpr (HasData<std::remove_cvref_t<A>>) {
231 // BumpAllocator sizing passes hand out empty spans by design.
232 if (arg.size() != 0) {
233 ctx.require_device_accessible(static_cast<const void*>(arg.data()), what.c_str());
234 }
235 }
236 // Anything else (option structs, scalars, enums) carries no pointer.
237}
238
239// Positional labelling: the dispatch macro forwards an unnamed pack.
240template <typename... Args>
241inline void require_pack_accessible(const Queue& ctx, const char* fn, const Args&... args) {
242 if (!pointer_checks_enabled()) return;
243 int pos = 0;
244 (void)std::initializer_list<int>{
245 (++pos,
246 require_arg_accessible(ctx, args,
247 std::string(fn) + ": argument " + std::to_string(pos)),
248 0)...};
249}
250
251// Any compiled-in backend; used only to probe "is the positional call well-formed?".
252inline constexpr Backend kProbeBackend =
253#if BATCHLAS_HAS_CUDA_BACKEND
255#elif BATCHLAS_HAS_ROCM_BACKEND
257#else
259#endif
260
261} // namespace detail
262
263} // namespace batchlas
264
265// Backend-deducing overload `NAME(ctx, args...)` of `NAME<Backend, T, ...>(ctx, ...)`.
266// The pack is forwarded, not restated, so the primary's defaults still apply. The
267// requires-clause must stay: unconstrained, this beats option-struct overloads and
268// fails inside its body. Even constrained, it beats `const MA&` for prvalues, so
269// checks must also live on the <Backend> option overloads.
270// evidence: docs/design/api-conventions.md#api-conventions-the-variadic-dispatch-overload-and-its-requires-clause
271#define BATCHLAS_DISPATCH_ON_QUEUE(NAME) \
272 template <typename... Args> \
273 requires requires(Queue& probe_ctx, Args&&... probe_args) { \
274 NAME<::batchlas::detail::kProbeBackend>(probe_ctx, \
275 std::forward<Args>(probe_args)...); \
276 } \
277 inline auto NAME(Queue& ctx, Args&&... args) { \
278 ::batchlas::detail::require_pack_accessible(ctx, #NAME, args...); \
279 return ::batchlas::detail::with_backend_gated< \
280 ::batchlas::detail::needs_fp64_v<Args...>>(ctx, #NAME, [&](auto Back) { \
281 return NAME<Back.value>(ctx, std::forward<Args>(args)...); \
282 }); \
283 }
Non-owning view of a batch of matrices: the argument type of every entry point.
Definition matrix.hh:485
Owning batch of batch_size() matrices of equal (capacity) shape.
Definition matrix.hh:1399
Non-owning view of a batch of strided vectors: the vector argument type of the entry points.
Definition matrix.hh:1754
No kernel, backend or vendor entry point in this build on this device serves the request.
Definition error.hh:122
Enumerations, scalar traits and scalar concepts shared by every public header.
const Settings & settings()
The process-wide settings.
auto with_backend_for(Queue &ctx, const char *op, F &&f)
with_backend() for a call on element type T: throws batchlas::unsupported before dispatch when T is d...
auto with_backend(Queue &ctx, F &&f)
Call f with the queue's backend as a compile-time constant.
constexpr std::string_view to_string(MatrixFormat v)
Spelling of v as in source.
Definition enums.hh:84
Backend
Library family an entry point dispatches to, used as a template argument.
Definition enums.hh:111
@ NETLIB
Host CBLAS / LAPACKE; the fallback for any device.
Definition enums.hh:119
@ SYCL
Native SYCL kernels on a GPU with no vendor library (AdaptiveCpp Metal); float and complex<float> onl...
Definition enums.hh:117
@ ROCM
AMD GPU: rocBLAS / rocSOLVER / rocSPARSE and the native SYCL kernels.
Definition enums.hh:114
@ CUDA
NVIDIA GPU: cuBLAS / cuSOLVER / cuSPARSE and the native SYCL kernels.
Definition enums.hh:113
The batched containers: Matrix (owning), MatrixView (non-owning), KernelMatrixView (device-side),...
void throw_no_device_fp64(const Queue &ctx, const char *op)
void require_arg_accessible(const Queue &ctx, const A &arg, const std::string &what)
void require_device_fp64(const Queue &ctx, const char *op)
Throws throw_no_device_fp64 for a double or complex<double> T on a GPU without fp64.
constexpr Backend kProbeBackend
constexpr bool needs_fp64_v
bool addresses_no_elements(const A &arg)
constexpr bool is_fp64_scalar_v
void require_pack_accessible(const Queue &ctx, const char *fn, const Args &... args)
auto with_backend_gated(Queue &ctx, const char *op, F &&f)
Every BATCHLAS_* environment variable the library reads, captured once into batchlas::Settings,...
DeviceType type
The device's type.
bool supports_fp64() const
True when the device runs double-precision kernels (sycl::aspect::fp64).
std::string get_name() const
The device's name as reported by SYCL.
Where BatchLAS work runs: a device, a backend, an ordered SYCL queue and a workspace arena.
batchlas::Backend backend() const
The backend calls dispatch to; never Backend::AUTO (an AUTO Queue resolves once, on first query,...
Device device() const
The device this Queue runs on.
void require_device_accessible(const void *ptr, const char *what) const
As is_device_accessible(), but throws.
UnsafeSettings unsafe
Overrides that remove a guarantee.
Definition settings.hh:491
Non-owning view of size() contiguous elements of type T.
Definition sycl-span.hh:39
Owning, growable array of T in USM shared memory (sycl::malloc_shared).
bool skip_pointer_checks
BATCHLAS_SKIP_POINTER_CHECKS: disables the per-argument USM reachability check that turns host memory...
Definition settings.hh:459
Owning batch of batch_size() vectors of length size() in USM shared memory.
Definition matrix.hh:1963
batchlas::Device, batchlas::Event and batchlas::Queue: where work runs and how it is ordered.
batchlas::Span, the non-owning pointer + length every entry point takes for 1-D data.
batchlas::UnifiedVector, an owning growable array in USM shared memory.