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