Template Numerical Library: TNL::Backend Namespace Reference

Searching...

No Matches

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)
	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.
template<typename Element, typename PushBuffer>
void	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.
void	checkErrorCode (const char *file_name, int line, error_t error)
void	checkLaunchConfiguration (const LaunchConfiguration &launch_configuration)
	Checks that the launch configuration satisfies the limits of the device.
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)
unsigned int	getGridSize (std::size_t totalBlocksCount, std::size_t gridIdx, std::size_t maxGridSize)
	Returns the number of blocks of the grid with index gridIdx when totalBlocksCount blocks are split into grids of at most maxGridSize blocks.
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 ()
	Returns the warp size for the target device.
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 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 &, const dim3 &gridIdx, dim3 &gridSize)
	Sets the size of the grid with index gridIdx when blocksCount blocks are split into grids of at most getMaxGrid{X,Y,Z}Size() blocks.
void	setupGrid (const dim3 &blocksCount, const dim3 &gridsCount, std::size_t gridIdx, dim3 &gridSize)
	Sets the size of the grid with index gridIdx in the x-dimension.
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

◆ checkLaunchConfiguration()

void TNL::Backend::checkLaunchConfiguration ( const LaunchConfiguration & launch_configuration )

inline

Checks that the launch configuration satisfies the limits of the device.

Some violations of the limits are not reported by the runtime, e.g. a grid with too many threads silently runs fewer threads on AMD GPUs.

Exceptions

std::invalid_argument when the configuration violates some limit.

◆ getGridSize()

unsigned int TNL::Backend::getGridSize	(	std::size_t	totalBlocksCount,
		std::size_t	gridIdx,
		std::size_t	maxGridSize )

inline

Returns the number of blocks of the grid with index gridIdx when totalBlocksCount blocks are split into grids of at most maxGridSize blocks.

All grids except the last one have maxGridSize blocks, the last one has the remaining blocks (which is maxGridSize when totalBlocksCount is its multiple).

gridIdx must be smaller than the number of grids, i.e. getNumberOfGrids( totalBlocksCount, maxGridSize ).

◆ 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

Returns the warp size for the target device.

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 to develop portable wave-aware code.

Note: 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 for all devices on the first call and returns the cached value afterwards.

Parameters

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

◆ setupGrid() [1/2]

void TNL::Backend::setupGrid	(	const dim3 &	blocksCount,
		const dim3 &	,
		const dim3 &	gridIdx,
		dim3 &	gridSize )

inline

Sets the size of the grid with index gridIdx when blocksCount blocks are split into grids of at most getMaxGrid{X,Y,Z}Size() blocks.

Each component of gridIdx must be smaller than the number of grids in the same dimension.

◆ setupGrid() [2/2]

void TNL::Backend::setupGrid	(	const dim3 &	blocksCount,
		const dim3 &	gridsCount,
		std::size_t	gridIdx,
		dim3 &	gridSize )

inline

Sets the size of the grid with index gridIdx in the x-dimension.

This overload is selected when gridIdx is a scalar. Without it, the scalar would be converted to dim3( gridIdx ), whose y and z components are 1 and thus refer to a nonexistent second grid in these dimensions.

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

Classes

Typedefs

Enumerations

Functions

Variables

Detailed Description

Function Documentation

◆ checkLaunchConfiguration()

◆ getGridSize()

◆ getMaxWarpSize()

◆ getMinWarpSize()

◆ getWarpFullMask()

◆ getWarpSize() [1/2]

◆ getWarpSize() [2/2]

◆ setupGrid() [1/2]

◆ setupGrid() [2/2]

◆ warp_shuffle_down()

◆ warp_shuffle_up()

◆ warp_shuffle_xor()

Variable Documentation

◆ is_warp_shuffle_native_v