GPU portability layer: host/device decoration macros and backend detection. More...

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

Macros
#define	EBGEOMETRY_HOST_DEVICE
	Mark a function as callable from both host and device code (no-op on host compilers).

#define	EBGEOMETRY_HOST
	Mark a function as callable from host code only (no-op on host compilers).

#define	EBGEOMETRY_GLOBAL
	Mark a function as a device kernel entry point (no-op on host compilers).

#define	EBGEOMETRY_DEVICE_NOINLINE
	Keep a function out of line in the device pass (CUDA/HIP); nothing on the host.

#define	EBGEOMETRY_HOST_ALWAYS_INLINE
	Always inline a function in a host pass (GCC and Clang); nothing in a device pass.

#define	EBGEOMETRY_HOST_NOINLINE
	Keep a function out of line in a host pass (GCC and Clang); nothing in a device pass.

#define	EBGEOMETRY_GRID_CONSTANT
	Let a kernel read a const parameter in place, through its address, rather than copy it (CUDA).

Detailed Description

GPU portability layer: host/device decoration macros and backend detection.

This header is the single home for the thin backend-abstraction macros used by the EBGeometry GPU port. It performs no allocation and pulls in no backend runtime headers; it only detects which offload compiler (if any) is translating the current translation unit and expands the decoration macros accordingly. On an ordinary host C++17 compiler every decoration macro expands to nothing, so annotated code is byte-for-byte identical to unannotated code.

It also defines the inlining and launch macros: EBGEOMETRY_DEVICE_NOINLINE (out of line in a device pass), EBGEOMETRY_HOST_ALWAYS_INLINE and EBGEOMETRY_HOST_NOINLINE (GCC/Clang attributes in a host pass, so these two are not empty on a host compiler), and EBGEOMETRY_GRID_CONSTANT (a kernel parameter read in place, CUDA only).

Decoration macros

EBGEOMETRY_HOST_DEVICE — callable from host and device (value types, the tape interpreter).
EBGEOMETRY_HOST — host-only. Unannotated functions are already implicitly host-only under CUDA/HIP; this macro is applied to host-only public API (file parsers, mesh readers, host-side builders) to make that intent explicit and to let the device compiler reject accidental device use.
EBGEOMETRY_GLOBAL — device kernel entry point.

Supported backends

CUDA (__CUDACC__): decorations expand to __host__ / __device__ / __global__.
HIP (__HIPCC__): same expansions as CUDA (HIP uses the identical attributes).

Any other compiler, including a SYCL or OpenACC compiler, sees host code: the macros expand to nothing, and nothing in EBGeometry is callable inside a SYCL kernel or an OpenACC compute region.

Author: Robert Marskar

Macros

Detailed Description