macros.hpp File Reference

Toolchain detection and portability macros. More...

#include <cstddef>
#include <tdls/core/target.hpp>

Include dependency graph for macros.hpp:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Macros
#define	TDLS_HOST_DEVICE
	__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
#define	TDLS_FORCE_INLINE_GPU   inline
	Inlining forced on GPU only: __forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS_HOST_DEVICE for why) and in the device pass of SYCL (__SYCL_DEVICE_ONLY__: icpx/DPC++ for Intel GPUs, and its NVIDIA and AMD targets), plain inline elsewhere.
#define	TDLS_FORCE_INLINE_ALL   inline
	Inlining forced on every backend, CPU included: TDLS_FORCE_INLINE_GPU on GPU, __forceinline under MSVC, always_inline under the GNU dialect, plain inline elsewhere.
#define	TDLS_NO_INLINE_GPU   inline
	Keeps a function out of line on GPU only: noinline under CUDA, HIP and in the device pass of SYCL, plain inline elsewhere; the pair of TDLS_FORCE_INLINE_GPU.
#define	TDLS_NO_INLINE
	Keeps a function out of line on every backend.
#define	TDLS_RESTRICT   __restrict__
	Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC; nothing in the code that nvc++ generates).
#define	TDLS_STRIDED(i, s)
	Offset of entry i of a strided object, i times stride s, in 64-bit arithmetic (std::size_t); i and s are nonnegative ints.
#define	TDLS_STRIDED_WIDE(i, s)
	Offset of entry i of a strided object, i times stride s, for a 64-bit stride s (std::size_t); i is a nonnegative int.
#define	TDLS_UNROLL_FORCE
	Full-unroll pragma in the local compiler dialect.
#define	TDLS_EXEC_CHECK_DISABLE
	Lifts the execution-space check of nvcc on the next function template; empty elsewhere.

Detailed Description

Toolchain detection and portability macros.

Author

Tristan Chenaille

Copyright

Copyright (C) 2026 CEA. All rights reserved. This project is publicly released under the BSD 3-Clause License (see the LICENSE file). CEA may also distribute it under specific licensing conditions.

Targeted toolchains: plain CPU (gcc, clang, MSVC), CUDA (nvcc, clang), HIP (hipcc/amdclang++), SYCL (icpx, AdaptiveCpp), stdpar (nvc++), OpenMP host/target, Kokkos (through its backend compiler). Only CUDA and HIP require function-space decorations; single-source models (SYCL, stdpar, OpenMP target) decide device callability when kernels are outlined, so no decoration is emitted there. OpenMP-target / OpenACC declare target / routine seq wrappers belong to the caller's dispatch layer, never to the solvers.

Every macro is #ifndef-guarded: each can be overridden from the command line (e.g. -DTDLS_FORCE_INLINE_GPU=inline) without editing headers.

Macro Definition Documentation

TDLS_EXEC_CHECK_DISABLE

#define TDLS_EXEC_CHECK_DISABLE

Lifts the execution-space check of nvcc on the next function template; empty elsewhere.

The CooperativeLUpp solvers call the barrier of the caller, a device function on GPU, from host-device function templates. nvcc checks the execution space of that call in the template itself and rejects it when the instantiation starts from a non-template kernel, although every instantiation actually used runs on the device. The pragma lifts the check for that one function, the practice of Thrust and CUB for user functors. clang, for CUDA and HIP alike, only checks the functions it emits and needs nothing.

TDLS_FORCE_INLINE_ALL

#define TDLS_FORCE_INLINE_ALL   inline

Inlining forced on every backend, CPU included: TDLS_FORCE_INLINE_GPU on GPU, __forceinline under MSVC, always_inline under the GNU dialect, plain inline elsewhere.

Reserved for the few helpers whose body was factored out of a solver routine: inlined, they leave its generated code as it was, while a call the CPU inliner may decline costs time (see TiledLUppSolverStatic::permute_cycles).

TDLS_FORCE_INLINE_GPU

#define TDLS_FORCE_INLINE_GPU   inline

Inlining forced on GPU only: __forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS_HOST_DEVICE for why) and in the device pass of SYCL (__SYCL_DEVICE_ONLY__: icpx/DPC++ for Intel GPUs, and its NVIDIA and AMD targets), plain inline elsewhere.

On GPU backends, forced inlining is the guard that keeps register tiles alive: an out-of-line call would demote them to slow local memory. On CPU a register spill to the stack is cheap and the heuristic inliner makes better calls than a blanket always_inline (which also bloats debug builds and compile times), so a plain inline hint is kept there. Its pair TDLS_FORCE_INLINE_ALL forces the inlining on CPU too.

TDLS_HOST_DEVICE

#define TDLS_HOST_DEVICE

__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.

Marks every solver function as callable from host code and from inside device kernels. Under CUDA the __host__ / __device__ macros are always predefined by the toolchain; under HIP they only exist once hip/hip_runtime.h has been included, which this library does not require, so the HIP branch spells the attributes those macros expand to and works whatever the include order.

TDLS_NO_INLINE

#define TDLS_NO_INLINE

Keeps a function out of line on every backend.

For the rare paths of a solver: an inlined rare path would share the registers of the common one and could push it to local memory on GPU (see the second pass of ArrowheadLUppSolverStatic).

TDLS_NO_INLINE_GPU

#define TDLS_NO_INLINE_GPU   inline

Keeps a function out of line on GPU only: noinline under CUDA, HIP and in the device pass of SYCL, plain inline elsewhere; the pair of TDLS_FORCE_INLINE_GPU.

For a function that the GPU inliner would merge into a large caller, where the merged function gets a worse register allocation than the two apart (see ArrowheadLUppConfig::streamed_out_of_line). On CPU the inliner is left free.

TDLS_RESTRICT

#define TDLS_RESTRICT   __restrict__

Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC; nothing in the code that nvc++ generates).

From -O2 on, nvc++ 25.11 fuses loops that depend on each other. It draws that right from the qualifier of the pointers that the solvers pass to one another, once inlined. Its host code then returns wrong solutions: the physical row interchanges of TiledLUpp, the blocks of ArrowheadLUpp eliminated in registers, ArrowheadLUpp in its stored and streamed solves. It also stops on an internal error on some TiledLUpp configurations. Without the qualifier, its code is right. Measured on the stored solve of ArrowheadLUpp: the same speed through -stdpar=gpu on an RTX 5080 and through -stdpar=multicore on 12 CPU threads, 0 to 5 % slower on one CPU core. The device pass of nvcc keeps the qualifier, whatever its host compiler.

TDLS_STRIDED

#define TDLS_STRIDED	(		i,
			s )

Value:

    (static_cast<std::size_t>(static_cast<unsigned>(i)) * static_cast<unsigned>(s))

Offset of entry i of a strided object, i times stride s, in 64-bit arithmetic (std::size_t); i and s are nonnegative ints.

The addressing macros of the solvers use it for every element of an object in remote memory: the offset of the last element of a batch in SoA, its index times the number of systems, may exceed 2^32 on a large GPU, where a 32-bit product would wrap around silently. Measured against the 32-bit product: same results bit for bit, no perceptible cost (RTX 5080, CPU).

In the code that nvc++ generates, i is extended with its sign rather than with zeros. From -O2 on, nvc++ 25.11 derives the addresses of some loops from an index that the loop never forms, a negative one, extended as this macro extends i. Extended with zeros, every address of such a loop lies 2^32 strides beyond the right one: the runtime solvers of TiledLUpp and ArrowheadLUpp read outside their operands and fault. Extended with its sign, every address is right. Both extensions give the same offset for every index the solvers form.

TDLS_STRIDED_WIDE

#define TDLS_STRIDED_WIDE	(		i,
			s )

Value:

(static_cast<std::size_t>(static_cast<unsigned>(i)) * (s))

Offset of entry i of a strided object, i times stride s, for a 64-bit stride s (std::size_t); i is a nonnegative int.

The stride between the columns of a multi right-hand-side block: a column of a batch in SoA lies the dimension times the number of systems away from the previous one, a product that may exceed 2^31 where the strides of the entries do not. In the code that nvc++ generates, i is extended with its sign, as in TDLS_STRIDED.

TDLS_UNROLL_FORCE

#define TDLS_UNROLL_FORCE

Full-unroll pragma in the local compiler dialect.

Dialect only: which loops carry it is a compile-time Config knob of each solver, applied through a two-branch if constexpr whose no-unroll branch carries no pragma at all. The GCC dialect needs an explicit unrolling factor. 64 fully unrolls the loops bounded by the tile size and those bounded by a system dimension up to 64; beyond, GCC unrolls them partially.