macros.hpp File Reference

Toolchain detection and portability macros. More...

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_FORCEINLINE   inline
	__forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS_HOST_DEVICE for why), plain inline elsewhere.
#define	TDLS_RESTRICT   __restrict__
	Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC).
#define	TDLS_UNROLL_FORCE
	Full-unroll pragma in the local compiler dialect.

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_FORCEINLINE=inline) without editing headers.

Macro Definition Documentation

TDLS_FORCEINLINE

#define TDLS_FORCEINLINE   inline

__forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS_HOST_DEVICE for why), 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.

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