Skip to content

Commit 25061d5

Browse files
f3schf3sch
committed
GPUCommonDefAPI: add GPUhdii()
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent 3f159b8 commit 25061d5

1 file changed

Lines changed: 33 additions & 28 deletions

File tree

‎GPU/Common/GPUCommonDefAPI.h‎

Lines changed: 33 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -28,37 +28,38 @@
2828
//All host-functions in GPU code are automatically inlined, to avoid duplicate symbols.
2929
//For non-inline host only functions, use no keyword at all!
3030
#if !defined(GPUCA_GPUCODE) || defined(__OPENCL_HOST__) || defined(__METAL_HOST__) || defined(__CPUDEVICE__) // For host / ROOT dictionary, and for the CPUDevice backend whose device code runs on the host
31-
#define GPUd() // device function
32-
#define GPUdDefault() // default (constructor / operator) device function
33-
#define GPUhdDefault() // default (constructor / operator) host device function
34-
#define GPUdi() inline // to-be-inlined device function
35-
#define GPUdii() // Only on GPU to-be-inlined (forced) device function
36-
#define GPUdni() // Device function, not-to-be-inlined
37-
#define GPUdnii() inline // Device function, not-to-be-inlined on device, inlined on host
38-
#define GPUh() // Host-only function
31+
#define GPUd() // device function
32+
#define GPUdDefault() // default (constructor / operator) device function
33+
#define GPUhdDefault() // default (constructor / operator) host device function
34+
#define GPUdi() inline // to-be-inlined device function
35+
#define GPUdii() // Only on GPU to-be-inlined (forced) device function
36+
#define GPUdni() // Device function, not-to-be-inlined
37+
#define GPUdnii() inline // Device function, not-to-be-inlined on device, inlined on host
38+
#define GPUh() // Host-only function
3939
// NOTE: All GPUd*() functions are also compiled on the host during host compilation.
4040
// The GPUh*() macros are for the rare cases of functions that you want to compile for the host during GPU compilation.
4141
// Usually, you do not need the GPUh*() versions. If in doubt, use GPUd*()!
42-
#define GPUhi() inline // to-be-inlined host-only function
43-
#define GPUhd() // Host and device function, inlined during GPU compilation to avoid symbol clashes in host code
44-
#define GPUhdi() inline // Host and device function, to-be-inlined on host and device
45-
#define GPUhdni() // Host and device function, not to-be-inlined automatically
46-
#define GPUg() INVALID_TRIGGER_ERROR_NO_GPU_CODE // GPU kernel
47-
#define GPUshared() // shared memory variable declaration
48-
#define GPUglobal() // global memory variable declaration (only used for kernel input pointers)
49-
#define GPUconstant() // constant memory variable declaraion
50-
#define GPUconstexpr() static constexpr // constexpr on GPU that needs to be instantiated for dynamic access (e.g. arrays), becomes __constant on GPU
51-
#define GPUglobalconstexpr() constexpr // constexpr variable at program scope, needs the constant address space in MSL
52-
#define GPUnoexcept() noexcept // noexcept where the backend supports it
53-
#define GPUprivate() // private memory variable declaration
54-
#define GPUgeneric() // reference / ptr to generic address space
55-
#define GPUbarrier() // synchronize all GPU threads in block
56-
#define GPUbarrierWarp() // synchronize threads inside warp
57-
#define GPUAtomic(type) type // atomic variable type
58-
#define GPUsharedref() // reference / ptr to shared memory
59-
#define GPUglobalref() // reference / ptr to global memory
60-
#define GPUconstantref() // reference / ptr to constant memory
61-
#define GPUconstexprref() // reference / ptr to variable declared as GPUconstexpr()
42+
#define GPUhi() inline // to-be-inlined host-only function
43+
#define GPUhd() // Host and device function, inlined during GPU compilation to avoid symbol clashes in host code
44+
#define GPUhdi() inline // Host and device function, to-be-inlined on host and device
45+
#define GPUhdii() inline __attribute__((always_inline)) // Host and device function, to-be-inlined (forced) on host and device
46+
#define GPUhdni() // Host and device function, not to-be-inlined automatically
47+
#define GPUg() INVALID_TRIGGER_ERROR_NO_GPU_CODE // GPU kernel
48+
#define GPUshared() // shared memory variable declaration
49+
#define GPUglobal() // global memory variable declaration (only used for kernel input pointers)
50+
#define GPUconstant() // constant memory variable declaraion
51+
#define GPUconstexpr() static constexpr // constexpr on GPU that needs to be instantiated for dynamic access (e.g. arrays), becomes __constant on GPU
52+
#define GPUglobalconstexpr() constexpr // constexpr variable at program scope, needs the constant address space in MSL
53+
#define GPUnoexcept() noexcept // noexcept where the backend supports it
54+
#define GPUprivate() // private memory variable declaration
55+
#define GPUgeneric() // reference / ptr to generic address space
56+
#define GPUbarrier() // synchronize all GPU threads in block
57+
#define GPUbarrierWarp() // synchronize threads inside warp
58+
#define GPUAtomic(type) type // atomic variable type
59+
#define GPUsharedref() // reference / ptr to shared memory
60+
#define GPUglobalref() // reference / ptr to global memory
61+
#define GPUconstantref() // reference / ptr to constant memory
62+
#define GPUconstexprref() // reference / ptr to variable declared as GPUconstexpr()
6263

6364
#ifndef __VECTOR_TYPES_H__ // FIXME: ROOT will pull in these CUDA definitions if built against CUDA, so we have to add an ugly protection here
6465
struct float4 { float x, y, z, w; };
@@ -88,6 +89,7 @@
8889
#define GPUhi() INVALID_TRIGGER_ERROR_NO_HOST_CODE
8990
#define GPUhd() inline
9091
#define GPUhdi() inline
92+
#define GPUhdii() __attribute__((always_inline)) inline
9193
#define GPUhdni()
9294
#define GPUg() __kernel
9395
#define GPUshared() __local
@@ -156,6 +158,7 @@
156158
#define GPUhi() inline
157159
#define GPUhd() inline
158160
#define GPUhdi() inline
161+
#define GPUhdii() inline __attribute__((always_inline))
159162
#define GPUhdni()
160163
#define GPUg() kernel
161164
#define GPUshared() threadgroup
@@ -186,6 +189,7 @@
186189
#define GPUhi() __host__ inline
187190
#define GPUhd() __host__ __device__ inline
188191
#define GPUhdi() __host__ __device__ inline
192+
#define GPUhdii() __host__ __device__ __forceinline__
189193
#define GPUhdni() __host__ __device__
190194
#define GPUg() __global__
191195
#define GPUshared() __shared__
@@ -221,6 +225,7 @@
221225
#define GPUhi() __host__ inline
222226
#define GPUhd() __host__ __device__ inline
223227
#define GPUhdi() __host__ __device__ inline
228+
#define GPUhdii() __host__ __device__ __forceinline__
224229
#define GPUhdni() __host__ __device__
225230
#define GPUg() __global__
226231
#define GPUshared() __shared__

0 commit comments

Comments
 (0)