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..45cc987453830 100644 --- a/GPU/Common/GPUCommonDefAPI.h +++ b/GPU/Common/GPUCommonDefAPI.h @@ -27,7 +27,7 @@ //Define macros for GPU keywords. i-version defines inline functions. //All host-functions in GPU code are automatically inlined, to avoid duplicate symbols. //For non-inline host only functions, use no keyword at all! -#if !defined(GPUCA_GPUCODE) || defined(__OPENCL_HOST__) // For host / ROOT dictionary +#if !defined(GPUCA_GPUCODE) || defined(__OPENCL_HOST__) || defined(__METAL_HOST__) // For host / ROOT dictionary #define GPUd() // device function #define GPUdDefault() // default (constructor / operator) device function #define GPUhdDefault() // default (constructor / operator) host device function @@ -124,6 +124,53 @@ #if (!defined(__OPENCL__) || !defined(GPUCA_NO_CONSTANT_MEMORY)) #define GPUconstantref() GPUconstant() #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() device + #define GPUconstant() constant // TODO: possibly add const __restrict where possible later! + #define GPUconstexpr() constant + #define GPUprivate() thread + #define GPUgeneric() + #define GPUglobalref() device + #define GPUsharedref() threadgroup + #define GPUprivateref() thread + #define GPUconstantref() constant + #define GPUconstexprref() GPUconstexpr() + #define GPUdouble() float + #define GPUbarrier() threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup) + #define GPUbarrierWarp() simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup) + #define GPUAtomic(type) atomic // atomic variable type #elif defined(__HIPCC__) //Defines for HIP #define GPUd() __device__ #define GPUdDefault() __device__ @@ -230,5 +277,5 @@ #define get_group_id(dim) iBlock #endif - // clang-format on +// clang-format on #endif