Skip to content

Commit 53569d4

Browse files
ktfktf
committed
GPU: Metal branches in the common definition macros
Defines the GPUd()/GPUshared()/GPU*ref() family for __METAL__ and __METAL_HOST__, and teaches GPUCommonDef.h about the Metal host and device compilation passes. The Metal backend targets MSL 4.1 and later only. That is what lets GPUdDefault() expand to nothing: up to MSL 4.0 a member function's implicit this is thread, which is wrong for objects living in device memory, and pinning defaulted constructors to device made the same type unusable in thread or threadgroup. MSL 4.1 makes an unannotated this generic, which is the C++ semantics this codebase already assumes. The *ref() macros stay explicit regardless: they are correct from the OpenCL port, explicit is never slower than generic, and constant is not covered by generic pointers at all. Inert unless __METAL__ or __METAL_HOST__ is defined.
1 parent 55db997 commit 53569d4

2 files changed

Lines changed: 97 additions & 3 deletions

File tree

‎GPU/Common/GPUCommonDef.h‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -31,12 +31,12 @@
3131
#include "GPUCommonDefSettings.h"
3232

3333
#if !defined(__CLING__) && !defined(G__ROOT) // No GPU code for ROOT
34-
#if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__)
34+
#if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__) || defined(__METAL_HOST__)
3535
#define GPUCA_GPUCODE // Compiled by GPU compiler
3636
#endif
3737

3838
#if defined(GPUCA_GPUCODE)
39-
#if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__)
39+
#if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__) || defined(__METAL_VERSION__)
4040
#define GPUCA_GPUCODE_DEVICE // Executed on device
4141
#endif
4242
#if defined(__CUDACC__)
@@ -45,6 +45,8 @@
4545
#define GPUCA_GPUTYPE HIP
4646
#elif defined(__OPENCL__) || defined(__OPENCL_HOST__)
4747
#define GPUCA_GPUTYPE OCL
48+
#elif defined(__METAL__) || defined(__METAL_HOST__)
49+
#define GPUCA_GPUTYPE METAL
4850
#endif
4951
#endif
5052
#endif

‎GPU/Common/GPUCommonDefAPI.h‎

Lines changed: 93 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -124,6 +124,98 @@
124124
#if (!defined(__OPENCL__) || !defined(GPUCA_NO_CONSTANT_MEMORY))
125125
#define GPUconstantref() GPUconstant()
126126
#endif
127+
#elif defined(__METAL_HOST__)
128+
#define GPUd() // device function
129+
#define GPUdDefault() // default (constructor / operator) device function
130+
#define GPUhdDefault() // default (constructor / operator) host device function
131+
#define GPUdi() inline // to-be-inlined device function
132+
#define GPUdii() // Only on GPU to-be-inlined (forced) device function
133+
#define GPUdni() // Device function, not-to-be-inlined
134+
#define GPUdnii() inline // Device function, not-to-be-inlined on device, inlined on host
135+
#define GPUh() // Host-only function
136+
// NOTE: All GPUd*() functions are also compiled on the host during host compilation.
137+
// The GPUh*() macros are for the rare cases of functions that you want to compile for the host during GPU compilation.
138+
// Usually, you do not need the GPUh*() versions. If in doubt, use GPUd*()!
139+
#define GPUhi() inline // to-be-inlined host-only function
140+
#define GPUhd() // Host and device function, inlined during GPU compilation to avoid symbol clashes in host code
141+
#define GPUhdi() inline // Host and device function, to-be-inlined on host and device
142+
#define GPUhdni() // Host and device function, not to-be-inlined automatically
143+
#define GPUg() INVALID_TRIGGER_ERROR_NO_GPU_CODE // GPU kernel
144+
#define GPUshared() // shared memory variable declaration
145+
#define GPUglobal() // global memory variable declaration (only used for kernel input pointers)
146+
#define GPUconstant() // constant memory variable declaraion
147+
#define GPUconstexpr() static constexpr // constexpr on GPU that needs to be instantiated for dynamic access (e.g. arrays), becomes __constant on GPU
148+
#define GPUprivate() // private memory variable declaration
149+
#define GPUgeneric() // reference / ptr to generic address space
150+
#define GPUbarrier() // synchronize all GPU threads in block
151+
#define GPUbarrierWarp() // synchronize threads inside warp
152+
#define GPUAtomic(type) type // atomic variable type
153+
#define GPUsharedref() // reference / ptr to shared memory
154+
#define GPUglobalref() // reference / ptr to global memory
155+
#define GPUconstantref() // reference / ptr to constant memory
156+
#define GPUconstexprref() // reference / ptr to variable declared as GPUconstexpr()
157+
158+
#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
159+
struct float4 { float x, y, z, w; };
160+
struct float3 { float x, y, z; };
161+
struct float2 { float x; float y; };
162+
struct uchar2 { uint8_t x, y; };
163+
struct short2 { int16_t x, y; };
164+
struct ushort2 { uint16_t x, y; };
165+
struct int2 { int32_t x, y; };
166+
struct int3 { int32_t x, y, z; };
167+
struct int4 { int32_t x, y, z, w; };
168+
struct uint1 { uint32_t x; };
169+
struct uint2 { uint32_t x, y; };
170+
struct uint3 { uint32_t x, y, z; };
171+
struct uint4 { uint32_t x, y, z, w; };
172+
struct dim3 { uint32_t x, y, z; };
173+
#endif
174+
#elif defined(__METAL__) //Defines for Metal Shading Language
175+
// ADDRESS SPACES. This backend targets MSL 4.1 (macOS 27) and later only --
176+
// see -std=metal4.1 in the CMakeLists, which fails the build on anything
177+
// older rather than miscompiling quietly.
178+
//
179+
// That version is what makes the port tractable: up to MSL 4.0 a member
180+
// function's implicit `this` is `thread`, which is wrong for us, since most
181+
// objects the kernels touch live in `device` memory. Pinning defaulted
182+
// constructors and operators to `device` was the 4.0 workaround, and it made
183+
// the same type unusable in `thread` or `threadgroup`. In 4.1 an unannotated
184+
// `this` is GENERIC and resolves to whichever address space the object is in
185+
// -- the C++ semantics this codebase already assumes -- so GPUdDefault()
186+
// needs nothing at all. The compiler resolves it statically in almost every
187+
// case; it only falls back to a runtime branch where it cannot see through,
188+
// such as argument buffers or dynamic libraries.
189+
//
190+
// The *ref() macros below stay explicit even so. They are already correct
191+
// from the OpenCL port, an explicit annotation is never slower than a generic
192+
// one, and `constant` is not covered by generic pointers at all.
193+
#define GPUdDefault() // generic `this` (MSL 4.1+)
194+
#define GPUd()
195+
#define GPUhdDefault()
196+
#define GPUdi() inline
197+
#define GPUdii() inline
198+
#define GPUdni()
199+
#define GPUdnii()
200+
#define GPUh() inline
201+
#define GPUhi() inline
202+
#define GPUhd() inline
203+
#define GPUhdi() inline
204+
#define GPUhdni()
205+
#define GPUg() kernel
206+
#define GPUshared() threadgroup
207+
#define GPUglobal() __global
208+
#define GPUconstant() constant // TODO: possibly add const __restrict where possible later!
209+
#define GPUconstexpr() constant
210+
#define GPUprivate() private
211+
#define GPUgeneric()
212+
#define GPUglobalref() device
213+
#define GPUsharedref() threadgroup
214+
#define GPUprivateref() thread
215+
#define GPUconstantref() constant
216+
#define GPUconstexprref() GPUconstexpr()
217+
#define GPUdouble() float
218+
#define GPUAtomic(type) atomic<type> // atomic variable type
127219
#elif defined(__HIPCC__) //Defines for HIP
128220
#define GPUd() __device__
129221
#define GPUdDefault() __device__
@@ -230,5 +322,5 @@
230322
#define get_group_id(dim) iBlock
231323
#endif
232324

233-
// clang-format on
325+
// clang-format on
234326
#endif

0 commit comments

Comments
 (0)