namespace

`SushiBLAS::Launch`

Declared in: include/SushiBLAS/support/launch_1d.hpp

Contains

SushiBLAS::Launch::Shape1D

Variables

template <typename T>
constexpr int K_VECTOR_WIDTH = (Core::is_complex_v<T> || sizeof(T) >= Device::K_VECTOR_TRANSACTION_BYTES) ? 1 : static_cast<int>(Device::K_VECTOR_TRANSACTION_BYTES / sizeof(T))

Holds the sycl::vec width the library considers for element type T; 1 for std::complex.

See also

include/SushiBLAS/support/README.md

Functions

bool is_vector_aligned(const void *p) noexcept

Returns whether p can be the base address of a 128-bit access.

See also

include/SushiBLAS/support/README.md

template <typename... Ptrs>
bool vector_aligned(const Ptrs *... ps) noexcept

Whether every pointer given can be the base address of a 128-bit access.

Template parameters

Ptrs: The pointee types, deduced.

Parameters

ps: The pointers to test.

Returns

True only if all of ps satisfy is_vector_aligned.

template <typename T, std::size_t N>
std::array< T *, N > typed_ptrs(const std::array< void *, N > &raw) noexcept

Reinterprets a dtype-erased pointer group as a typed one.

See also

include/SushiBLAS/support/README.md

Shape1D plan_for(const Device::Profile &p, std::size_t n, std::size_t width) noexcept

Turns a device profile and an access width into a launch geometry.

Parameters

width: Elements per access: 1 for scalar, the sycl::vec width for a 128-bit path.

Returns

The geometry; Shape1D::vector_steps is 0 when width is 1.

See also

include/SushiBLAS/support/README.md

template <typename T>
Shape1D plan(const sycl::queue &q, std::size_t n, bool aligned)

Decides the launch geometry for n elements of type T.

Parameters

aligned: Whether every operand satisfies is_vector_aligned.

Returns

The geometry; Shape1D::vector_steps is 0 when the wide path is not taken.

Shape1D plan_scalar(const sycl::queue &q, std::size_t n)

Decides the launch geometry for n indices, with no wide path.

Returns

The geometry, always with Shape1D::vector_steps == 0.

See also

include/SushiBLAS/support/README.md

template <typename Func, typename T, typename... Ins>
void map(sycl::handler &h, const Shape1D &shape, std::size_t n, Func op, T *out, const Ins *... ins)

Launches out[i] = op(in0[i], ..., inM[i]) over n elements in a given shape.

Precondition

shape.vector_steps * K_VECTOR_WIDTH<T> <= n, and out and every input satisfy is_vector_aligned when shape.vector_steps is above zero.

Parameters

op: Device-copyable, captures by value, holds no host pointer, allocates and throws nothing.
out: May alias an input at the same index only.

See also

include/SushiBLAS/support/README.md

template <typename Func, typename T, typename... Ins>
void map(sycl::handler &h, const sycl::queue &q, std::size_t n, bool aligned, Func op, T *out, const Ins *... ins)

Launches out[i] = op(in0[i], ..., inM[i]) over n elements in the shape the queue's device plans.

Parameters

aligned: Whether out and every input satisfy is_vector_aligned.

See also

include/SushiBLAS/support/README.md

template <typename T, std::size_t N, typename Func>
void zip_inplace(sycl::handler &h, const Shape1D &shape, std::size_t n, const std::array< T *, N > &ptrs, Func op)

Launches op(t0[i], ..., tN-1[i]) over n elements in a given shape, every slot writable.

Precondition

The preconditions of the shape form of map, for every array in ptrs.

Parameters

ptrs: The arrays to update, in the order op expects; they must not overlap each other.
op: Callable (auto&...) -> void under the same device-kernel constraints as map.

See also

include/SushiBLAS/support/README.md

template <typename T, std::size_t N, typename Func>
void zip_inplace(sycl::handler &h, const sycl::queue &q, std::size_t n, bool aligned, const std::array< T *, N > &ptrs, Func op)

Launches op(t0[i], ..., tN-1[i]) over n elements in the shape the queue's device plans.

Parameters

aligned: Whether every array in ptrs satisfies is_vector_aligned.

See also

include/SushiBLAS/support/README.md

template <typename Func>
void indexed(sycl::handler &h, const Shape1D &shape, std::size_t n, Func op)

Launches op(i) once for each index in [0, n) in a given shape, with no wide path.

Parameters

op: Callable (std::size_t i) -> void under the same device-kernel constraints as map.

See also

include/SushiBLAS/support/README.md

template <typename Func>
void indexed(sycl::handler &h, const sycl::queue &q, std::size_t n, Func op)

Launches op(i) once for each index in [0, n) in the shape the queue's device plans.

See also

include/SushiBLAS/support/README.md