TNL::Backend Namespace Reference

Internal namespace for CUDA/HIP backend support. More...

Classes
struct	LaunchConfiguration
	Holds the parameters necessary to launch a CUDA or HIP kernel (i.e. schedule it for execution on some stream of some device). More...
struct	SharedMemory
struct	SharedMemory< T, 16 >
struct	SharedMemory< T, 32 >
struct	SharedMemory< T, 64 >
struct	SharedMemory< T, 8 >
class	Stream
class	StreamPool
struct	Uninitialized
	Storage wrapper for type-punning through an array of device words. More...

Typedefs
using	error_t = cudaError_t
using	stream_t = cudaStream_t

Enumerations
enum	{ StreamDefault = cudaStreamDefault , StreamNonBlocking = cudaStreamNonBlocking }
enum	FuncCache { FuncCachePreferNone = cudaFuncCachePreferNone , FuncCachePreferShared = cudaFuncCachePreferShared , FuncCachePreferL1 = cudaFuncCachePreferL1 , FuncCachePreferEqual = cudaFuncCachePreferEqual }
enum	MemcpyKind {   MemcpyHostToHost = cudaMemcpyHostToHost , MemcpyHostToDevice = cudaMemcpyHostToDevice , MemcpyDeviceToHost = cudaMemcpyDeviceToHost , MemcpyDeviceToDevice = cudaMemcpyDeviceToDevice ,   MemcpyDefault = cudaMemcpyDefault }

Functions
template<typename Element, typename FillBuffer, typename PushBuffer>
void	bufferedTransfer (std::size_t size, FillBuffer &fill, PushBuffer &push)
template<typename Element, typename FillBuffer>
void	bufferedTransferToDevice (Element *destination, std::size_t size, FillBuffer &fill)
template<typename Element, typename PushBuffer>
void	bufferedTransferToHost (const Element *source, std::size_t size, PushBuffer &push)
void	checkErrorCode (const char *file_name, int line, error_t error)
void	deviceSynchronize ()
template<class T>
void	funcSetCacheConfig (T *func, enum FuncCache cacheConfig)
int	getArchitectureMajor (int deviceNum)
int	getArchitectureMinor (int deviceNum)
int	getDevice ()
	Returns the ID of the active device.
int	getDeviceCores (int deviceNum)
int	getDeviceCoresPerMultiprocessors (int deviceNum)
int	getDeviceCount ()
	Returns the number of devices available in the system.
int	getDeviceMultiprocessors (int deviceNum)
std::string	getDeviceName (int deviceNum)
bool	getECCEnabled (int deviceNum)
std::size_t	getFreeGlobalMemory ()
__device__ std::size_t	getGlobalBlockIdx_x (const dim3 &gridIdx)
__device__ std::size_t	getGlobalBlockIdx_y (const dim3 &gridIdx)
__device__ std::size_t	getGlobalBlockIdx_z (const dim3 &gridIdx)
std::size_t	getGlobalMemorySize (int deviceNum)
__device__ std::size_t	getGlobalThreadIdx_x (const dim3 &gridIdx)
__device__ std::size_t	getGlobalThreadIdx_y (const dim3 &gridIdx)
__device__ std::size_t	getGlobalThreadIdx_z (const dim3 &gridIdx)
template<typename Index>
__device__ Index	getInterleaving (const Index index)
constexpr int	getMaxBlockXSize ()
constexpr int	getMaxBlockYSize ()
constexpr int	getMaxBlockZSize ()
constexpr std::size_t	getMaxGridXSize ()
constexpr std::size_t	getMaxGridYSize ()
constexpr std::size_t	getMaxGridZSize ()
constexpr int	getMaxThreadsPerBlock ()
constexpr std::size_t	getMaxThreadsPerGrid ()
constexpr int	getMaxWarpSize ()
	Returns the maximum possible warp/wavefront size for the current build.
constexpr int	getMinWarpSize ()
	Returns the minimum possible warp/wavefront size across all platforms.
template<typename Index>
Index	getNumberOfBlocks (const Index threads, const unsigned int blockSize)
template<typename Index>
Index	getNumberOfGrids (const Index blocks, const unsigned int gridSize)
constexpr int	getNumberOfSharedMemoryBanks ()
int	getRegistersPerMultiprocessor (int deviceNum)
template<typename T>
__device__ T *	getSharedMemory ()
std::size_t	getSharedMemoryPerBlock (int deviceNum)
constexpr std::size_t	getTransferBufferSize ()
constexpr auto	getWarpFullMask ()
	Returns the full mask for warp shuffle operations.
constexpr int	getWarpSize ()
int	getWarpSize (int deviceId)
	Runtime query for the warp/wavefront size of a specific device.
template<typename RawKernel, typename... KernelParameters>
void	launchKernel (RawKernel kernel_function, LaunchConfiguration launch_configuration, KernelParameters &&... parameters)
template<typename RawKernel, typename... KernelParameters>
void	launchKernelAsync (RawKernel kernel_function, LaunchConfiguration launch_configuration, KernelParameters &&... parameters)
template<typename RawKernel, typename... KernelParameters>
void	launchKernelSync (RawKernel kernel_function, LaunchConfiguration launch_configuration, KernelParameters &&... parameters)
template<class T>
__device__ T	ldg (const T &value)
	Loads data from a global memory using the __ldg() intrinsic.
void	memcpy (void *dst, const void *src, std::size_t sizeBytes, MemcpyKind kind)
std::ostream &	operator<< (std::ostream &str, const dim3 &d)
void	printThreadsSetup (const dim3 &blockSize, const dim3 &blocksCount, const dim3 &gridSize, const dim3 &gridsCount, std::ostream &str=std::cout)
void	setDevice (int device)
	Sets the active device.
void	setupGrid (const dim3 &blocksCount, const dim3 &gridsCount, const dim3 &gridIdx, dim3 &gridSize)
void	setupThreads (const dim3 &blockSize, dim3 &blocksCount, dim3 &gridsCount, long long int xThreads, long long int yThreads=0, long long int zThreads=0)
stream_t	streamCreateWithPriority (unsigned int flags, int priority)
void	streamDestroy (stream_t stream)
void	streamSynchronize (stream_t stream)
template<typename T>
__device__ T	warp_shuffle_down (T val, int delta)
	Generic warp shuffle down for any bit-copyable type.
template<typename T>
__device__ T	warp_shuffle_up (T val, int delta)
	Generic warp shuffle up for any bit-copyable type.
template<typename T>
__device__ T	warp_shuffle_xor (T val, int laneMask)
	Generic warp shuffle xor for any bit-copyable type.

Variables
template<typename T>
constexpr bool	is_warp_shuffle_native_v
	Variable template listing the types with native __shfl_*_sync overloads.

Detailed Description

Internal namespace for CUDA/HIP backend support.

Function Documentation

bufferedTransferToDevice()

template<typename Element, typename FillBuffer>

void TNL::Backend::bufferedTransferToDevice	(	Element *	destination,
		std::size_t	size,
		FillBuffer &	fill )

This function creates a buffer on the host, the fill handler fills it with data and this function transfers data from the buffer to the destination, which is a pointer to device memory.

bufferedTransferToHost()

template<typename Element, typename PushBuffer>

void TNL::Backend::bufferedTransferToHost	(	const Element *	source,
		std::size_t	size,
		PushBuffer &	push )

This function creates a buffer on the host, fills it with data transferred from source, which is a pointer to device memory, and the push handler processes the data in the buffer.

getMaxWarpSize()

int TNL::Backend::getMaxWarpSize ( )

constexpr

Returns the maximum possible warp/wavefront size for the current build.

Unlike getWarpSize(), this is consistent between host and device compilation:

• CUDA: always 32
• HIP: always 64 (the maximum across all AMD architectures)
• Host-only: 32

Use this for compile-time guards that must be consistent on the host side, e.g. deciding whether to instantiate TPS=64 kernel variants.

getMinWarpSize()

int TNL::Backend::getMinWarpSize ( )

constexpr

Returns the minimum possible warp/wavefront size across all platforms.

Always returns 32. Use this for compile-time checks that must allow configs valid on any architecture, e.g. column-major SlicedEllpack launch configs where TPS * SliceSize >= getMinWarpSize() ensures the group covers at least one warp on the smallest-warp-size device.

getWarpFullMask()

auto TNL::Backend::getWarpFullMask ( )

nodiscardconstexpr

Returns the full mask for warp shuffle operations.

HIP shfl intrinsics require a 64-bit mask regardless of the wavefront size (unused upper bits are zero for wave32). CUDA uses a 32-bit mask.

Warning: this function relies on __GFX8__/__GFX9__ arch macros which are defined only during device compilation. On the host side, it falls through to the wavefront=32 mask. Do not call from host code — use only in __device__ functions.

getWarpSize() [1/2]

int TNL::Backend::getWarpSize ( )

constexpr

The warpSize variable is of type int and contains the warp size (in threads) for the target device. This should be used only from device code in order to develop portable wave-aware code.

Note that NVIDIA devices return 32; AMD devices return 64 for gfx8/gfx9 and 32 for gfx10 and above.

Warning: the returned value may be inconsistent when used from a host-code context on HIP, because __GFX8__/__GFX9__ are defined only during device compilation. On the host side, the function falls through to return 32. Use getMaxWarpSize() for compile-time host-side guards, or the runtime overload getWarpSize(int deviceId) for host-side dispatch decisions. See https://clang.llvm.org/docs/HIPSupport.html#predefined-macros for details.

https://rocm.docs.amd.com/projects/HIP/en/latest/reference/kernel_language.html#warpsize

getWarpSize() [2/2]

int TNL::Backend::getWarpSize ( int deviceId )

inlinenodiscard

Runtime query for the warp/wavefront size of a specific device.

Use this for host-side launch configuration decisions (TPS selection, etc.). For CUDA, always returns 32. For HIP, queries the runtime API.

Parameters

deviceId

The device ID to query (default: current device).

warp_shuffle_down()

template<typename T>

__device__ T TNL::Backend::warp_shuffle_down	(	T	val,
		int	delta )

Generic warp shuffle down for any bit-copyable type.

See warp_shuffle_xor for details.

warp_shuffle_up()

template<typename T>

__device__ T TNL::Backend::warp_shuffle_up	(	T	val,
		int	delta )

Generic warp shuffle up for any bit-copyable type.

See warp_shuffle_xor for details.

warp_shuffle_xor()

template<typename T>

__device__ T TNL::Backend::warp_shuffle_xor	(	T	val,
		int	laneMask )

Generic warp shuffle xor for any bit-copyable type.

CUDA's __shfl_xor_sync only has overloads for a fixed set of primitive types (see is_warp_shuffle_native_v). This wrapper splits any other type into an array of words and shuffles each word separately — the same pattern used by NVIDIA's CUB and libcu++ libraries. Works for TNL::Arithmetics::Complex and any other trivially copyable struct with sizeof <= 32.

Variable Documentation

is_warp_shuffle_native_v

template<typename T>

bool TNL::Backend::is_warp_shuffle_native_v

inlineconstexpr

Initial value:

                                            =
std::is_same_v< T, int > || std::is_same_v< T, unsigned int > || std::is_same_v< T, long >
|| std::is_same_v< T, unsigned long > || std::is_same_v< T, long long > || std::is_same_v< T, unsigned long long >
|| std::is_same_v< T, float > || std::is_same_v< T, double >

Variable template listing the types with native __shfl_*_sync overloads.

CUDA/HIP __shfl_*_sync functions only accept int, unsigned int, long, unsigned long, long long, unsigned long long, float, and double. Other arithmetic types like bool, char, and short do not have overloads.