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

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

struct launch_on#
#include <kernel_launch.hpp>

Where a kernel is launched: the stream, the dynamic shared memory size, 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);

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.

Public Functions

inline launch_on(
resources const &res,
std::size_t smem = 0,
std::source_location loc = std::source_location::current()
)#
Parameters:

  • res – [in] raft resources providing the stream to launch on

  • smem – [in] dynamic shared memory size in bytes

  • 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,
std::source_location loc = std::source_location::current()
)#
Parameters:

  • stream – [in] stream to launch on

  • smem – [in] dynamic shared memory size in bytes

  • loc – [in] call site to blame for launch errors; leave at its default

inline launch_on(
cudaStream_t stream,
std::size_t smem = 0,
std::source_location loc = std::source_location::current()
)#
Parameters:

  • stream – [in] stream to launch on

  • smem – [in] dynamic shared memory size in bytes

  • 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.