From 53569d4c602b70957876a12c3981dc7bebc1c25f Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Sun, 6 Sep 2026 09:47:30 +0200 Subject: [PATCH] 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. --- GPU/Common/GPUCommonDef.h | 6 ++- GPU/Common/GPUCommonDefAPI.h | 94 +++++++++++++++++++++++++++++++++++- 2 files changed, 97 insertions(+), 3 deletions(-) diff --git a/GPU/Common/GPUCommonDef.h b/GPU/Common/GPUCommonDef.h index ffe5551f02f1b..28768b57c43de 100644 --- a/GPU/Common/GPUCommonDef.h +++ b/GPU/Common/GPUCommonDef.h @@ -31,12 +31,12 @@ #include "GPUCommonDefSettings.h" #if !defined(__CLING__) && !defined(G__ROOT) // No GPU code for ROOT - #if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__) + #if defined(__CUDACC__) || defined(__OPENCL__) || defined(__HIPCC__) || defined(__OPENCL_HOST__) || defined(__METAL_HOST__) #define GPUCA_GPUCODE // Compiled by GPU compiler #endif #if defined(GPUCA_GPUCODE) - #if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__) + #if defined(__CUDA_ARCH__) || defined(__OPENCL__) || defined(__HIP_DEVICE_COMPILE__) || defined(__METAL_VERSION__) #define GPUCA_GPUCODE_DEVICE // Executed on device #endif #if defined(__CUDACC__) @@ -45,6 +45,8 @@ #define GPUCA_GPUTYPE HIP #elif defined(__OPENCL__) || defined(__OPENCL_HOST__) #define GPUCA_GPUTYPE OCL + #elif defined(__METAL__) || defined(__METAL_HOST__) + #define GPUCA_GPUTYPE METAL #endif #endif #endif diff --git a/GPU/Common/GPUCommonDefAPI.h b/GPU/Common/GPUCommonDefAPI.h index 4d4e04f10b2fa..6dd1b386b1648 100644 --- a/GPU/Common/GPUCommonDefAPI.h +++ b/GPU/Common/GPUCommonDefAPI.h @@ -124,6 +124,98 @@ #if (!defined(__OPENCL__) || !defined(GPUCA_NO_CONSTANT_MEMORY)) #define GPUconstantref() GPUconstant() #endif +#elif defined(__METAL_HOST__) + #define GPUd() // device function + #define GPUdDefault() // default (constructor / operator) device function + #define GPUhdDefault() // default (constructor / operator) host device function + #define GPUdi() inline // to-be-inlined device function + #define GPUdii() // Only on GPU to-be-inlined (forced) device function + #define GPUdni() // Device function, not-to-be-inlined + #define GPUdnii() inline // Device function, not-to-be-inlined on device, inlined on host + #define GPUh() // Host-only function + // NOTE: All GPUd*() functions are also compiled on the host during host compilation. + // The GPUh*() macros are for the rare cases of functions that you want to compile for the host during GPU compilation. + // Usually, you do not need the GPUh*() versions. If in doubt, use GPUd*()! + #define GPUhi() inline // to-be-inlined host-only function + #define GPUhd() // Host and device function, inlined during GPU compilation to avoid symbol clashes in host code + #define GPUhdi() inline // Host and device function, to-be-inlined on host and device + #define GPUhdni() // Host and device function, not to-be-inlined automatically + #define GPUg() INVALID_TRIGGER_ERROR_NO_GPU_CODE // GPU kernel + #define GPUshared() // shared memory variable declaration + #define GPUglobal() // global memory variable declaration (only used for kernel input pointers) + #define GPUconstant() // constant memory variable declaraion + #define GPUconstexpr() static constexpr // constexpr on GPU that needs to be instantiated for dynamic access (e.g. arrays), becomes __constant on GPU + #define GPUprivate() // private memory variable declaration + #define GPUgeneric() // reference / ptr to generic address space + #define GPUbarrier() // synchronize all GPU threads in block + #define GPUbarrierWarp() // synchronize threads inside warp + #define GPUAtomic(type) type // atomic variable type + #define GPUsharedref() // reference / ptr to shared memory + #define GPUglobalref() // reference / ptr to global memory + #define GPUconstantref() // reference / ptr to constant memory + #define GPUconstexprref() // reference / ptr to variable declared as GPUconstexpr() + + #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 + struct float4 { float x, y, z, w; }; + struct float3 { float x, y, z; }; + struct float2 { float x; float y; }; + struct uchar2 { uint8_t x, y; }; + struct short2 { int16_t x, y; }; + struct ushort2 { uint16_t x, y; }; + struct int2 { int32_t x, y; }; + struct int3 { int32_t x, y, z; }; + struct int4 { int32_t x, y, z, w; }; + struct uint1 { uint32_t x; }; + struct uint2 { uint32_t x, y; }; + struct uint3 { uint32_t x, y, z; }; + struct uint4 { uint32_t x, y, z, w; }; + struct dim3 { uint32_t x, y, z; }; + #endif +#elif defined(__METAL__) //Defines for Metal Shading Language + // ADDRESS SPACES. This backend targets MSL 4.1 (macOS 27) and later only -- + // see -std=metal4.1 in the CMakeLists, which fails the build on anything + // older rather than miscompiling quietly. + // + // That version is what makes the port tractable: up to MSL 4.0 a member + // function's implicit `this` is `thread`, which is wrong for us, since most + // objects the kernels touch live in `device` memory. Pinning defaulted + // constructors and operators to `device` was the 4.0 workaround, and it made + // the same type unusable in `thread` or `threadgroup`. In 4.1 an unannotated + // `this` is GENERIC and resolves to whichever address space the object is in + // -- the C++ semantics this codebase already assumes -- so GPUdDefault() + // needs nothing at all. The compiler resolves it statically in almost every + // case; it only falls back to a runtime branch where it cannot see through, + // such as argument buffers or dynamic libraries. + // + // The *ref() macros below stay explicit even so. They are already correct + // from the OpenCL port, an explicit annotation is never slower than a generic + // one, and `constant` is not covered by generic pointers at all. + #define GPUdDefault() // generic `this` (MSL 4.1+) + #define GPUd() + #define GPUhdDefault() + #define GPUdi() inline + #define GPUdii() inline + #define GPUdni() + #define GPUdnii() + #define GPUh() inline + #define GPUhi() inline + #define GPUhd() inline + #define GPUhdi() inline + #define GPUhdni() + #define GPUg() kernel + #define GPUshared() threadgroup + #define GPUglobal() __global + #define GPUconstant() constant // TODO: possibly add const __restrict where possible later! + #define GPUconstexpr() constant + #define GPUprivate() private + #define GPUgeneric() + #define GPUglobalref() device + #define GPUsharedref() threadgroup + #define GPUprivateref() thread + #define GPUconstantref() constant + #define GPUconstexprref() GPUconstexpr() + #define GPUdouble() float + #define GPUAtomic(type) atomic // atomic variable type #elif defined(__HIPCC__) //Defines for HIP #define GPUd() __device__ #define GPUdDefault() __device__ @@ -230,5 +322,5 @@ #define get_group_id(dim) iBlock #endif - // clang-format on +// clang-format on #endif