Utilities#

RAFT contains numerous utility functions and primitives that are easily usable. This page provides C++ API references for the publicly-exposed utility functions.

Memory Pool#

#include <raft/utils/memory_pool.cuh>

namespace raft

Warning

doxygengroup: Cannot find group “memory_pool” in doxygen xml output for project “RAFT” from directory: ../../cpp/doxygen/_xml/

Kernel Launch#

#include <raft/util/kernel_launch.hpp>

namespace raft

inline cudaLaunchAttribute cooperative()#

Launch attribute: the kernel synchronizes across the whole grid.

Such a launch fails unless the whole grid is resident on the device at once, so its grid size has to come from an occupancy query rather than from the problem size.

inline cudaLaunchAttribute shmem_carveout(unsigned percent)#

Launch attribute: preferred share of the combined L1/shared memory to use as shared memory, in percent.

Only a hint; the driver may pick a different split. Unrelated to the cap on dynamic shared memory, which is a property of the kernel rather than of the launch.

template<typename ...Args> void launch_kernel( launch_on where, dim3 grid, dim3 block, std::type_identity_t<void (*)(std::remove_cvref_t<Args>...)> kernel, Args&&... args )#

Launch kernel with args, which already have the kernel parameter types.

The launch arguments are checked against the kernel parameters at compile time, and a failed launch throws raft::cuda_error blaming the call site:

raft::launch_kernel(res, grid, block, my_kernel, arg0, arg1);

The function-pointer parameter type is a non-deduced context derived from args, so a partially specified function template (e.g. map_kernel<R, PassOffset>) can still convert to a unique __global__ pointer by deducing its remaining template parameters from that type. Overload sets that remain ambiguous after that conversion are not supported.

Parameters:

where – [in] stream to launch on, dynamic shared memory size, and the call site
grid – [in] grid dimensions
block – [in] block dimensions
kernel – [in] the __global__ function to launch
args – [in] arguments to pass to kernel

template<typename ...Params, typename ...Args> void launch_kernel( launch_on where, dim3 grid, dim3 block, void (*kernel)(Params...), Args&&... args )#

Launch kernel, converting args to the kernel parameter types.

Handles call sites where an argument merely converts to its parameter (e.g. T* to const T*), so they do not need casts. kernel must name a single specialization here, because its parameter types are what the arguments are converted to.

Parameters:

where – [in] stream to launch on, dynamic shared memory size, and the call site
grid – [in] grid dimensions
block – [in] block dimensions
kernel – [in] the __global__ function to launch
args – [in] arguments to convert and pass to kernel

template<typename ...Params, typename ...Args> void launch_kernel( launch_on where, dim3 grid, dim3 block, kernel_ref<void(Params...)> kernel, Args&&... args )#

Launch a kernel named by a runtime handle, converting args to its parameter types.

Behaves like the converting overload above, except that the kernel and its parameter types come from kernel rather than from a __global__ function pointer:

raft::launch_kernel({res, smem}, grid, block,
                    raft::kernel_ref<scan_kernel_t>{launcher->get_kernel()}, queries, n);

Unlike the two function-pointer overloads, this one accepts arguments that already have the parameter types too, because there is no exactly-matching overload for them to prefer.

Parameters:

where – [in] stream to launch on, dynamic shared memory size, attributes, and the call site
grid – [in] grid dimensions
block – [in] block dimensions
kernel – [in] handle to the loaded kernel, with the signature to launch it by
args – [in] arguments to convert and pass to kernel

struct launch_on#

#include <kernel_launch.hpp>

How and where a kernel is launched: the stream, the dynamic shared memory size, the launch attributes, and the call site to blame for launch errors.

Converts implicitly from raft resources or from a stream, so that a launch reads as a single call and the diagnostics of a failed launch point at the launch expression:

raft::launch_kernel(res, grid, block, my_kernel, arg0, arg1);
raft::launch_kernel({stream, smem}, grid, block, my_kernel, arg0, arg1);
raft::launch_kernel({res, smem, {raft::cooperative()}}, grid, block, my_kernel, arg0, arg1);

Launching on raft resources is dry run compliant: the kernel does not run when the handle has the dry run flag set, so such a launch needs no guard of its own. The stream overloads cannot know that, hence their kSkipExecution argument.

Copy and move are deleted and launch_kernel takes this by value, so the parameter can only be initialized from a prvalue: an instance stored in a variable can never be launched, and the captured location is therefore always the one of the launch expression. That is also what makes it safe for config to point at attrs, whose backing array lives until the end of that expression.

Public Functions

inline launch_on( resources const &res, std::size_t smem = 0, std::initializer_list<cudaLaunchAttribute> attrs = {}, std::source_location loc = std::source_location::current() )#

Launch on the stream owned by the resources.

The launch is skipped when the handle is in dry run mode; do not add a dry-run guard around it (see docs/source/dry_run_protocol.md).

Parameters:

res – [in] raft resources providing the stream to launch on
smem – [in] dynamic shared memory size in bytes
attrs – [in] launch attributes, e.g. raft::cooperative()
loc – [in] call site to blame for launch errors; leave at its default

inline launch_on( rmm::cuda_stream_view stream, std::size_t smem = 0, bool kSkipExecution = false, std::initializer_list<cudaLaunchAttribute> attrs = {}, std::source_location loc = std::source_location::current() )#

Launch on an explicit stream, which carries no dry-run state of its own.

In code reachable from an API taking raft::resources, either launch on the resources instead or pass kSkipExecution, otherwise the kernel runs in dry-run mode, which must not execute any CUDA work.

Parameters:

stream – [in] stream to launch on
smem – [in] dynamic shared memory size in bytes
kSkipExecution – [in] whether to skip the launch, e.g. a dry-run flag plumbed by the caller
attrs – [in] launch attributes, e.g. raft::cooperative()
loc – [in] call site to blame for launch errors; leave at its default

inline launch_on( cudaStream_t stream, std::size_t smem = 0, bool kSkipExecution = false, std::initializer_list<cudaLaunchAttribute> attrs = {}, std::source_location loc = std::source_location::current() )#

Launch on an explicit stream, which carries no dry-run state of its own.

Parameters:

stream – [in] stream to launch on
smem – [in] dynamic shared memory size in bytes
kSkipExecution – [in] whether to skip the launch, e.g. a dry-run flag plumbed by the caller
attrs – [in] launch attributes, e.g. raft::cooperative()
loc – [in] call site to blame for launch errors; leave at its default

Public Members

std::source_location location#: Call site to blame for launch errors.

cudaLaunchConfig_t config = {}#: Launch configuration; the grid and block dimensions are filled in by the launch.

detail::launch_flags flags = {}#: How to launch; derived from the resources rather than given at the call site.

template<typename Signature> struct kernel_ref#

#include <kernel_launch.hpp>

A kernel that exists only at run time, together with the signature it was compiled with.

A kernel loaded from a runtime-linked library (cudaLibraryGetKernel, e.g. after a JIT LTO link) has no __global__ function pointer for launch_kernel to read the parameter types from, so the signature is named explicitly:

using scan_kernel_t = void(float const*, std::uint32_t);
raft::launch_kernel({res, smem}, grid, block,
                    raft::kernel_ref<scan_kernel_t>{handle}, queries, n_queries);

Whether the handle really has that signature is on whoever loaded it. Given the signature, the launch converts each argument to its parameter type, so a call site does not need casts to make the argument types match the kernel exactly.

Template Parameters:: Signature – the kernel’s function type, e.g. void(float const*, std::uint32_t)

template<typename ...Params> struct kernel_ref<void(Params...)>#

#include <kernel_launch.hpp>

Public Functions

inline explicit kernel_ref(cudaKernel_t kernel)#

Parameters:: kernel – [in] handle to a loaded kernel whose signature is void(Params...)

Public Members

cudaKernel_t handle#: Handle to the loaded kernel.