diff --git a/Common/MathUtils/include/MathUtils/detail/StatAccumulator.h b/Common/MathUtils/include/MathUtils/detail/StatAccumulator.h index abb8a716cc5ee..d35327086eeee 100644 --- a/Common/MathUtils/include/MathUtils/detail/StatAccumulator.h +++ b/Common/MathUtils/include/MathUtils/detail/StatAccumulator.h @@ -42,7 +42,7 @@ struct StatAccumulator { wsum += w; n++; } - double getMean() const { return wsum > 0. ? sum / wsum : 0.; } + double getMean() const { return wsum > 0. ? sum / wsum : double{}; } #ifndef GPUCA_GPUCODE_DEVICE template diff --git a/Common/MathUtils/include/MathUtils/detail/trigonometric.h b/Common/MathUtils/include/MathUtils/detail/trigonometric.h index e13d965663dc9..afdfbb57b0f28 100644 --- a/Common/MathUtils/include/MathUtils/detail/trigonometric.h +++ b/Common/MathUtils/include/MathUtils/detail/trigonometric.h @@ -48,7 +48,7 @@ GPUhdi() T to02Pi(T phi) template GPUhdi() void bringTo02Pi(T& phi) { - phi = to02Pi(phi); + phi = to02Pi(phi); } template @@ -68,7 +68,7 @@ inline T to02PiGen(T phi) template inline void bringTo02PiGen(T& phi) { - phi = to02PiGen(phi); + phi = to02PiGen(phi); } template @@ -87,7 +87,7 @@ GPUhdi() T toPMPi(T phi) template GPUhdi() void bringToPMPi(T& phi) { - phi = toPMPi(phi); + phi = toPMPi(phi); } template @@ -107,10 +107,10 @@ inline T toPMPiGen(T phi) template inline void bringToPMPiGen(T& phi) { - phi = toPMPiGen(phi); + phi = toPMPiGen(phi); } -#ifdef __OPENCL__ // TODO: get rid of that stupid workaround for OpenCL template address spaces +#if defined(__OPENCL__) || defined(__METAL__) // TODO: get rid of that stupid workaround for OpenCL template address spaces template GPUhdi() void sincos(T ang, S& s, U& c) { @@ -275,12 +275,11 @@ GPUhdi() constexpr T fastATan2(T y, T x) // Average inaccuracy: 0.00048 // Max inaccuracy: 0.00084 // Speed: 6.2 times faster than atan2f() - constexpr T Pi = 3.1415926535897932384626433832795; - auto atan = [](T a) -> T { // returns the arctan for the angular range [-Pi/4, Pi/4] // the polynomial coefficients are taken from: // https://stackoverflow.com/questions/42537957/fast-accurate-atan-arctan-approximation-algorithm + constexpr T Pi = 3.1415926535897932384626433832795; constexpr T A = 0.0776509570923569; constexpr T B = -0.287434475393028; constexpr T C = ((Pi / 4) - A - B); @@ -290,6 +289,7 @@ GPUhdi() constexpr T fastATan2(T y, T x) auto atan2P = [atan](T yy, T xx) -> T { // fast atan2(yy,xx) for the angular range [0,+Pi] + constexpr T Pi = 3.1415926535897932384626433832795; constexpr T Pi025 = 1 * Pi / 4; constexpr T Pi075 = 3 * Pi / 4; const T x1 = xx + yy; // point p1 (x1,y1) = (xx,yy) - Pi/4 diff --git a/DataFormats/Reconstruction/src/TrackParametrizationWithError.cxx b/DataFormats/Reconstruction/src/TrackParametrizationWithError.cxx index 748cb47094d26..abecef528f7bb 100644 --- a/DataFormats/Reconstruction/src/TrackParametrizationWithError.cxx +++ b/DataFormats/Reconstruction/src/TrackParametrizationWithError.cxx @@ -180,10 +180,10 @@ GPUd() bool TrackParametrizationWithError::propagateTo(value_t xk, Trac } value_t kb = bz * constants::math::B2C; // evaluate in double prec. - double snpRef0 = linRef0.getSnp(), cspRef0 = gpu::CAMath::Sqrt((1 - snpRef0) * (1 + snpRef0)); - double snpRef1 = linRef1.getSnp(), cspRef1 = gpu::CAMath::Sqrt((1 - snpRef1) * (1 + snpRef1)); - double cspRef0Inv = 1 / cspRef0, cspRef1Inv = 1 / cspRef1, cc = cspRef0 + cspRef1, ccInv = 1 / cc, dy2dx = (snpRef0 + snpRef1) * ccInv; - double dxccInv = dx * ccInv, hh = dxccInv * cspRef1Inv * (1 + cspRef0 * cspRef1 + snpRef0 * snpRef1), jj = dx * (dy2dx - snpRef1 * cspRef1Inv); + double snpRef0 = linRef0.getSnp(), cspRef0 = gpu::CAMath::Sqrt((1.f - snpRef0) * (1.f + snpRef0)); + double snpRef1 = linRef1.getSnp(), cspRef1 = gpu::CAMath::Sqrt((1.f - snpRef1) * (1.f + snpRef1)); + double cspRef0Inv = 1.f / cspRef0, cspRef1Inv = 1.f / cspRef1, cc = cspRef0 + cspRef1, ccInv = 1.f / cc, dy2dx = (snpRef0 + snpRef1) * ccInv; + double dxccInv = dx * ccInv, hh = dxccInv * cspRef1Inv * (1.f + cspRef0 * cspRef1 + snpRef0 * snpRef1), jj = dx * (dy2dx - snpRef1 * cspRef1Inv); double f02 = hh * cspRef0Inv; double f04 = hh * dxccInv * kb; @@ -638,7 +638,7 @@ GPUd() bool TrackParametrizationWithError::propagateTo(value_t xk, cons } double r1pr2Inv = 1. / (r1 + r2), r2inv = 1. / r2, r1inv = 1. / r1; double dy2dx = (f1 + f2) * r1pr2Inv, dx2r1pr2 = dx * r1pr2Inv; - value_t step = (gpu::CAMath::Abs(x2r) < 0.05f) ? dx * gpu::CAMath::Abs(r2 + f2 * dy2dx) // chord + value_t step = (gpu::CAMath::Abs(x2r) < 0.05f) ? value_t(dx * gpu::CAMath::Abs(r2 + f2 * dy2dx)) // chord : 2.f * gpu::CAMath::ASin(0.5f * dx * gpu::CAMath::Sqrt(1.f + dy2dx * dy2dx) * crv) / crv; // arc step *= gpu::CAMath::Sqrt(1.f + this->getTgl() * this->getTgl()); // @@ -1086,7 +1086,7 @@ GPUd() auto TrackParametrizationWithError::getPredictedChi2(const value auto chi2 = (d * (szz * d - sdz * z) + z * (sdd * z - d * sdz)) / det; if (chi2 < 0.) { #ifndef GPUCA_ALIGPUCODE - LOGP(warning, "Negative chi2={}, Cluster: {} {} {} Dy:{} Dz:{} | sdd:{} sdz:{} szz:{} det:{}", chi2, cov[0], cov[1], cov[2], d, z, sdd, sdz, szz, det); + LOGP(warning, "Negative chi2={}, Cluster: {} {} {} Dy:{} Dz:{} | sdd:{} sdz:{} szz:{} det:{}", double(chi2), cov[0], cov[1], cov[2], d, z, double(sdd), double(sdz), double(szz), double(det)); LOGP(warning, "Track: {}", asString()); #endif } @@ -1286,7 +1286,7 @@ GPUd() bool TrackParametrizationWithError::update(const TrackParametriz } // updated covariance: Cov0 = Cov0 - K*Cov0 - matK *= o2::math_utils::SMatrix>(matC0); + matK *= MatrixD5(matC0); mC[kSigY2] -= matK(kY, kY); mC[kSigZY] -= matK(kZ, kY); mC[kSigZ2] -= matK(kZ, kZ); diff --git a/GPU/Common/GPUCommonAlgorithm.h b/GPU/Common/GPUCommonAlgorithm.h index 338a6fb4c5ad4..2e87f08580cb2 100644 --- a/GPU/Common/GPUCommonAlgorithm.h +++ b/GPU/Common/GPUCommonAlgorithm.h @@ -32,13 +32,13 @@ class GPUCommonAlgorithm template GPUd() static void sort(T* begin, T* end); template - GPUd() static void sortInBlock(T* begin, T* end); + GPUd() static void sortInBlock(int32_t nThreads, int32_t iThread, T* begin, T* end); template GPUd() static void sortDeviceDynamic(T* begin, T* end); template GPUd() static void sort(T* begin, T* end, const S& comp); template - GPUd() static void sortInBlock(T* begin, T* end, const S& comp); + GPUd() static void sortInBlock(int32_t nThreads, int32_t iThread, T* begin, T* end, const S& comp); template GPUd() static void sortDeviceDynamic(T* begin, T* end, const S& comp); #if __cplusplus >= 202002L // sortOnDevice takes an auto parameter @@ -268,29 +268,29 @@ GPUdi() void GPUCommonAlgorithm::sort(T* begin, T* end, const S& comp) } template -GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end) +GPUdi() void GPUCommonAlgorithm::sortInBlock(int32_t nThreads, int32_t iThread, T* begin, T* end) { #ifndef GPUCA_GPUCODE GPUCommonAlgorithm::sort(begin, end); #else - GPUCommonAlgorithm::sortInBlock(begin, end, [](auto&& x, auto&& y) { return x < y; }); + GPUCommonAlgorithm::sortInBlock(nThreads, iThread, begin, end, [](auto&& x, auto&& y) { return x < y; }); #endif } template -GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp) +GPUdi() void GPUCommonAlgorithm::sortInBlock(int32_t nThreads, int32_t iThread, T* begin, T* end, const S& comp) { #ifndef GPUCA_GPUCODE GPUCommonAlgorithm::sort(begin, end, comp); #elif defined(GPUCA_DETERMINISTIC_MODE) // Not using GPUCA_DETERMINISTIC_CODE, which is enforced in TPC compression - if (get_local_id(0) == 0) { + if (iThread == 0) { GPUCommonAlgorithm::sort(begin, end, comp); } GPUbarrier(); #else int32_t n = end - begin; for (int32_t i = 0; i < n; i++) { - for (int32_t tIdx = get_local_id(0); tIdx < n; tIdx += get_local_size(0)) { + for (int32_t tIdx = iThread; tIdx < n; tIdx += nThreads) { int32_t offset = i % 2; int32_t curPos = 2 * tIdx + offset; int32_t nextPos = curPos + 1; diff --git a/GPU/Common/GPUCommonDef.h b/GPU/Common/GPUCommonDef.h index 90746019d9a99..3b15e16314eef 100644 --- a/GPU/Common/GPUCommonDef.h +++ b/GPU/Common/GPUCommonDef.h @@ -81,6 +81,14 @@ #define GPUCA_RTC_CONSTEXPR #endif +#if defined(GPUCA_DETERMINISTIC_MODE) && defined(__METAL__) + // The deterministic paths compute the transcendentals in double (see + // GPUCommonMath::SinCos). The emulated double is bit-exact for + - * / but + // its sin and cos are within 2 ulp of libm, not identical to it, so the + // results could not match the other backends. + #error "GPUCA_DETERMINISTIC_MODE is not supported on Metal" +#endif + #ifndef GPUCA_DETERMINISTIC_CODE #ifdef GPUCA_DETERMINISTIC_MODE #define GPUCA_DETERMINISTIC_CODE(det, indet) det // In deterministic mode, take deterministic code path diff --git a/GPU/Common/GPUCommonDefAPI.h b/GPU/Common/GPUCommonDefAPI.h index 36e8cd403e083..d38dcf86cf9ff 100644 --- a/GPU/Common/GPUCommonDefAPI.h +++ b/GPU/Common/GPUCommonDefAPI.h @@ -166,14 +166,14 @@ #define GPUnoexcept() #define GPUprivate() thread #define GPUgeneric() - #define GPUglobalref() device + #define GPUglobalref() #define GPUsharedref() threadgroup #define GPUprivateref() thread #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 + #define GPUAtomic(type) type // atomic variable type #elif defined(__HIPCC__) //Defines for HIP #define GPUd() __device__ #define GPUdDefault() __device__ @@ -277,6 +277,15 @@ #define get_group_id(dim) (blockIdx.x) #elif defined(__OPENCL__) // Using OpenCL defaults +#elif defined(__METAL__) + // MSL has no work-item builtins; these come in as kernel attributes, declared + // by GPUCA_KRNL_GRID_ARGS on every entry point. + #define get_global_id(dim) (_metalTgIg * _metalTPerTg + _metalTiTg) + #define get_global_size(dim) (_metalTPerTg * _metalTgPerG) + #define get_num_groups(dim) (_metalTgPerG) + #define get_local_id(dim) (_metalTiTg) + #define get_local_size(dim) (_metalTPerTg) + #define get_group_id(dim) (_metalTgIg) #else #define get_global_id(dim) iBlock #define get_global_size(dim) nBlocks diff --git a/GPU/Common/GPUCommonDouble.h b/GPU/Common/GPUCommonDouble.h new file mode 100644 index 0000000000000..2f7108b409a7e --- /dev/null +++ b/GPU/Common/GPUCommonDouble.h @@ -0,0 +1,49 @@ +// Copyright 2019-2025 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +/// \file GPUCommonDouble.h +/// \brief Storage for double precision members shared with a device that has none + +/// \brief What MSL needs once the double keyword names the emulated type + +#ifndef GPUCOMMONDOUBLE_H +#define GPUCOMMONDOUBLE_H + +#include "GPUCommonDef.h" +#include "GPUCommonMath.h" + +#ifdef __METAL__ + +namespace o2::gpu +{ + +static_assert(sizeof(double) == 8, "the emulated double must match the size of a real one"); +static_assert(alignof(double) == 8, "the emulated double must match the alignment of a real one"); + +// CAMath::Abs deduces its parameter rather than taking a float, so a call on a +// double picks the primary template, which has no definition. The rest of CAMath +// takes float and is reached through the implicit conversion. +template <> +GPUhdi() constexpr double GPUCommonMath::Abs(double x) +{ + return double::fromBits(x.bits() & ~GPUCA_B64_SIGN); +} + +// metal::fabs is not constant-evaluable, so this also fails to compile if the +// specialisation above is ever dropped and the call falls back to it in float +static_assert(GPUCommonMath::Abs(GPUdoubleBinary64::fromBits(0xBFF0000000000001ULL)).bits() == 0x3FF0000000000001ULL, + "Abs on the emulated double must clear the sign bit and keep every other one"); + +} // namespace o2::gpu + +#endif // __METAL__ + +#endif // GPUCOMMONDOUBLE_H diff --git a/GPU/Common/GPUCommonDoubleBinary64.h b/GPU/Common/GPUCommonDoubleBinary64.h new file mode 100644 index 0000000000000..572e17797c0a2 --- /dev/null +++ b/GPU/Common/GPUCommonDoubleBinary64.h @@ -0,0 +1,589 @@ +// Copyright 2019-2026 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +/// \file GPUCommonDoubleBinary64.h +/// \brief IEEE-754 binary64 in software, for a device that has none + +#ifndef GPUCOMMONDOUBLEBINARY64_H +#define GPUCOMMONDOUBLEBINARY64_H + +#if !defined(__METAL__) && !defined(GPUCA_B64_HOST_REFERENCE) +#error "the emulated binary64 is for Metal, which has no double of its own; every other backend has a real one. Define GPUCA_B64_HOST_REFERENCE to build it on the host as a test reference." +#endif + +#include "GPUCommonDef.h" +#ifndef GPUCA_GPUCODE_DEVICE +#include +#endif + +// IEEE-754 binary64 in software, for a device that has no double at all. Round to +// nearest even only, with subnormals, infinities and NaNs; NaN propagation follows +// the ARM64 order, so an Apple host is a bit-exact reference down to the payload. +// There is no fused multiply-add, no square root and no other rounding mode. +// +// Addition, subtraction, multiplication, division and the conversions to and from +// float and the 32-bit integers are exact, so a Metal build reproduces the CPU +// result bit for bit wherever it uses only those. sin and cos come from fdlibm and +// land within 2 ulp of libm rather than matching it. +// +// It costs of the order of a hundred times plain float on an M-series GPU, which +// the tracking can afford because double is a small fraction of its floating point +// work. + +namespace o2::gpu +{ + +namespace binary64_detail +{ + +#ifdef __METAL__ +typedef ulong u64; +typedef long i64; +typedef uint u32; +#else +typedef uint64_t u64; +typedef int64_t i64; +typedef uint32_t u32; +#endif + +#define GPUCA_B64_ALWAYS inline __attribute__((always_inline)) +#ifdef __METAL__ +// The arithmetic has to stay out of line: inlined into a kernel it drops the +// occupancy (maxTotalThreadsPerThreadgroup 832 -> 384) and runs two to five +// times slower than the call. +#define GPUCA_B64_OP __attribute__((noinline)) +#else +#define GPUCA_B64_OP inline +#endif + +#ifdef __METAL__ +// metal::clz is not constant-evaluable, which would keep the whole class out of +// constant expressions; the builtin folds to the same ctlz at run time. +GPUCA_B64_ALWAYS constexpr int32_t clz64(u64 x) { return x ? __builtin_clzl(x) : 64; } +GPUCA_B64_ALWAYS constexpr int32_t clz32(u32 x) { return x ? __builtin_clz(x) : 32; } +GPUCA_B64_ALWAYS constexpr u32 asu32(float f) { return as_type(f); } +GPUCA_B64_ALWAYS constexpr float asf32(u32 u) { return as_type(u); } +#else +GPUCA_B64_ALWAYS constexpr int32_t clz64(u64 x) { return x ? __builtin_clzll(x) : 64; } +GPUCA_B64_ALWAYS constexpr int32_t clz32(u32 x) { return x ? __builtin_clz(x) : 32; } + +GPUCA_B64_ALWAYS constexpr u32 asu32(float f) { return __builtin_bit_cast(uint32_t, f); } +GPUCA_B64_ALWAYS constexpr float asf32(u32 u) { return __builtin_bit_cast(float, u); } +#endif + +// The high half of a 64x64 product, from 32-bit partial products. Neither +// metal::mulhi nor __int128 can appear in a constant expression, and the +// program-scope constants are constexpr. +GPUCA_B64_ALWAYS constexpr u64 mulhiu(u64 a, u64 b) +{ + const u64 al = a & 0xffffffffu, ah = a >> 32, bl = b & 0xffffffffu, bh = b >> 32; + const u64 ll = al * bl, lh = al * bh, hl = ah * bl, hh = ah * bh; + const u64 mid = (ll >> 32) + (lh & 0xffffffffu) + (hl & 0xffffffffu); + return hh + (lh >> 32) + (hl >> 32) + (mid >> 32); +} +GPUCA_B64_ALWAYS constexpr i64 mulhis(i64 a, i64 b) +{ + u64 h = mulhiu((u64)a, (u64)b); + h -= (a < 0) ? (u64)b : 0ULL; + h -= (b < 0) ? (u64)a : 0ULL; + return (i64)h; +} + +#define GPUCA_B64_SIGN 0x8000000000000000ULL +#define GPUCA_B64_FRAC 0x000fffffffffffffULL +#define GPUCA_B64_IMPL 0x0010000000000000ULL +#define GPUCA_B64_QUIET 0x0008000000000000ULL +#define GPUCA_B64_INF 0x7ff0000000000000ULL +#define GPUCA_B64_DNAN 0x7ff8000000000000ULL // ARM default NaN + +// right shift keeping a sticky bit; any count >= 0, no shift count ever reaches 64 +GPUCA_B64_ALWAYS constexpr u64 shrJam(u64 m, int32_t s) +{ + const int32_t sc = s > 63 ? 63 : s; + const u64 r = m >> sc; + const u64 lost = (m << (63 - sc)) << 1; + const bool big = s > 63; + const u64 rr = big ? 0ULL : r; + const u64 st = big ? m : lost; + return rr | (st != 0 ? 1ULL : 0ULL); +} + +// m: leading bit at 62 for a normal result, 10 guard bits below the 53-bit significand, +// value = m * 2^(e - 1085) with e the biased exponent. Packing (e-1)<<52 + m lets a +// rounding carry ripple into the exponent. +GPUCA_B64_ALWAYS constexpr u64 roundPack(u64 sign, int32_t e, u64 m) +{ + if (e >= 0x7ff) { + return sign | GPUCA_B64_INF; + } + if (e <= 0) { + m = shrJam(m, 1 - e); + e = 1; + } + const u64 r = m & 0x3ffULL; + m >>= 10; + if (r > 0x200ULL || (r == 0x200ULL && (m & 1ULL))) { + ++m; + } + return sign | (((u64)(e - 1) << 52) + m); +} + +// an sNaN operand wins over a qNaN one, among equals the first wins, the result is quietened +GPUCA_B64_ALWAYS constexpr u64 propNaN(u64 a, u64 b, bool an, bool bn) +{ + const bool as = an && !(a & GPUCA_B64_QUIET), bs = bn && !(b & GPUCA_B64_QUIET); + if (as) { + return a | GPUCA_B64_QUIET; + } + if (bs) { + return b | GPUCA_B64_QUIET; + } + return an ? a : b; +} + +GPUCA_B64_OP constexpr u64 addsub(u64 a, u64 b0, bool neg) +{ + const u64 b = neg ? (b0 ^ GPUCA_B64_SIGN) : b0; + const bool sw = (a & ~GPUCA_B64_SIGN) < (b & ~GPUCA_B64_SIGN); + const u64 x = sw ? b : a, y = sw ? a : b; // |x| >= |y| + int32_t ex = (int32_t)((x >> 52) & 0x7ff), ey = (int32_t)((y >> 52) & 0x7ff); + u64 mx = x & GPUCA_B64_FRAC, my = y & GPUCA_B64_FRAC; + if (ex == 0x7ff) { // y can only be inf/NaN if x is too + const bool an = ((a >> 52) & 0x7ff) == 0x7ff && (a & GPUCA_B64_FRAC) != 0, bn = ((b0 >> 52) & 0x7ff) == 0x7ff && (b0 & GPUCA_B64_FRAC) != 0; + if (an || bn) { + return propNaN(a, b0, an, bn); + } + if (ey == 0x7ff) { + return ((x ^ y) & GPUCA_B64_SIGN) ? GPUCA_B64_DNAN : x; + } + return x; + } + const u64 sx = x & GPUCA_B64_SIGN; + const bool sub = ((x ^ y) >> 63) != 0; + mx = (ex ? (mx | GPUCA_B64_IMPL) : mx) << 10; + ex = ex ? ex : 1; + my = (ey ? (my | GPUCA_B64_IMPL) : my) << 10; + ey = ey ? ey : 1; + my = shrJam(my, ex - ey); + const u64 s = sub ? mx - my : mx + my; + const bool carry = (s >> 63) != 0; + const int32_t lz = clz64(s) - 1; // -1 on carry, 63 on zero + const u64 sN = carry ? ((s >> 1) | (s & 1ULL)) : (s << (lz & 63)); + const int32_t eN = carry ? ex + 1 : ex - lz; + const bool zero = s == 0; + return roundPack((zero && sub) ? 0ULL : sx, zero ? -1 : eN, sN); +} + +GPUCA_B64_OP constexpr u64 mul(u64 a, u64 b) +{ + const u64 sign = (a ^ b) & GPUCA_B64_SIGN; + int32_t ea = (int32_t)((a >> 52) & 0x7ff), eb = (int32_t)((b >> 52) & 0x7ff); + u64 ma = a & GPUCA_B64_FRAC, mb = b & GPUCA_B64_FRAC; + if (ea == 0x7ff || eb == 0x7ff) { + const bool an = ea == 0x7ff && ma != 0, bn = eb == 0x7ff && mb != 0; + if (an || bn) { + return propNaN(a, b, an, bn); + } + if ((ea == 0 && ma == 0) || (eb == 0 && mb == 0)) { + return GPUCA_B64_DNAN; // inf * 0 + } + return sign | GPUCA_B64_INF; + } + if (ea == 0) { + if (ma == 0) { + return sign; + } + const int32_t lz = clz64(ma) - 11; + ma <<= lz; + ea = 1 - lz; + } else { + ma |= GPUCA_B64_IMPL; + } + if (eb == 0) { + if (mb == 0) { + return sign; + } + const int32_t lz = clz64(mb) - 11; + mb <<= lz; + eb = 1 - lz; + } else { + mb |= GPUCA_B64_IMPL; + } + const u64 lo = ma * mb, hi = mulhiu(ma, mb); // 106-bit product in [2^104, 2^106) + const bool top = (hi & (1ULL << 41)) != 0; + const u64 m1 = (hi << 21) | (lo >> 43), m0 = (hi << 22) | (lo >> 42); + const u64 st = top ? (lo & ((1ULL << 43) - 1)) : (lo & ((1ULL << 42) - 1)); + const u64 m = (top ? m1 : m0) | (st != 0 ? 1ULL : 0ULL); + return roundPack(sign, ea + eb - 1023 + (top ? 1 : 0), m); +} + +GPUCA_B64_OP constexpr u64 div(u64 a, u64 b) +{ + const u64 sign = (a ^ b) & GPUCA_B64_SIGN; + int32_t ea = (int32_t)((a >> 52) & 0x7ff), eb = (int32_t)((b >> 52) & 0x7ff); + u64 ma = a & GPUCA_B64_FRAC, mb = b & GPUCA_B64_FRAC; + if (ea == 0x7ff || eb == 0x7ff) { + const bool an = ea == 0x7ff && ma != 0, bn = eb == 0x7ff && mb != 0; + if (an || bn) { + return propNaN(a, b, an, bn); + } + if (ea == 0x7ff && eb == 0x7ff) { + return GPUCA_B64_DNAN; // inf / inf + } + return ea == 0x7ff ? (sign | GPUCA_B64_INF) : sign; // inf / x, x / inf + } + if (eb == 0 && mb == 0) { + return (ea == 0 && ma == 0) ? GPUCA_B64_DNAN : (sign | GPUCA_B64_INF); // x / 0 + } + if (ea == 0) { + if (ma == 0) { + return sign; + } + const int32_t lz = clz64(ma) - 11; + ma <<= lz; + ea = 1 - lz; + } else { + ma |= GPUCA_B64_IMPL; + } + if (eb == 0) { + const int32_t lz = clz64(mb) - 11; + mb <<= lz; + eb = 1 - lz; + } else { + mb |= GPUCA_B64_IMPL; + } + const bool lt = ma < mb; + const int32_t e = ea - eb + 1023 - (lt ? 1 : 0); + const u64 A2 = lt ? (ma << 1) : ma; // A2 / mb in [1, 2) + // reciprocal R ~ 2^114 / mb in (2^61, 2^62], seeded from a float division on the top 24 bits + const u32 rb = asu32(1.0f / (float)(u32)(mb >> 29)); + u64 R = (u64)((rb & 0x7fffffu) | 0x800000u) << ((int32_t)((rb >> 23) & 0xff) - 127 + 62); + const u64 Bn = mb << 11; + for (int32_t it = 0; it < 2; ++it) { + const i64 E = (i64)(1ULL << 61) - (i64)mulhiu(Bn, R); + R = (u64)((i64)R + mulhis((i64)R, E * 8)); // not E << 3: shifting a negative value is not a constant expression + } + R = R > (1ULL << 62) ? (1ULL << 62) : R; + // Q ~ A2 * 2^62 / mb, then the exact remainder, which fits in a signed 64-bit + // word because Q is within a few units + u64 Q = mulhiu(A2 << 10, (R << 2) - 1); + i64 rem = (i64)((A2 << 62) - Q * mb); + const i64 adj = mulhis(rem, (i64)R) >> 50; + Q = (u64)((i64)Q + adj); + rem -= adj * (i64)mb; + // after adj the remainder is within one divisor of [0, mb): one predicated step each way + const bool ng = rem < 0; + Q = ng ? Q - 1 : Q; + rem = ng ? rem + (i64)mb : rem; + const bool bg = rem >= (i64)mb; + Q = bg ? Q + 1 : Q; + rem = bg ? rem - (i64)mb : rem; + return roundPack(sign, e, Q | (rem != 0 ? 1ULL : 0ULL)); +} + +// an int32 or a uint32 always fits the 53-bit significand, so these are exact +GPUCA_B64_ALWAYS constexpr u64 fromU32(u32 x) +{ + if (x == 0) { + return 0ULL; + } + const int32_t lz = clz32(x); + return ((u64)(31 - lz + 1023) << 52) | (((u64)x << (21 + lz)) & GPUCA_B64_FRAC); +} + +GPUCA_B64_ALWAYS constexpr u64 fromI32(int32_t x) +{ + return fromU32(x < 0 ? (u32)(-(i64)x) : (u32)x) | (x < 0 ? GPUCA_B64_SIGN : 0ULL); +} + +GPUCA_B64_ALWAYS constexpr u64 fromFloat(float f) +{ + const u32 u = asu32(f); + const u64 sign = (u64)(u & 0x80000000u) << 32; + int32_t e = (int32_t)((u >> 23) & 0xff); + u32 m = u & 0x7fffffu; + if (e == 0xff) { + return sign | GPUCA_B64_INF | ((u64)m << 29) | (m ? GPUCA_B64_QUIET : 0ULL); + } + if (e == 0) { + if (m == 0) { + return sign; + } + const int32_t lz = clz32(m) - 8; + m <<= lz; + e = 1 - lz; + } + return sign | ((u64)(e - 127 + 1023) << 52) | ((u64)(m & 0x7fffffu) << 29); +} + +// binary64 -> binary32, round to nearest even, subnormals, inf, NaN (payload kept, quietened) +GPUCA_B64_OP constexpr float toFloat(u64 d) +{ + const u32 sign = (u32)(d >> 32) & 0x80000000u; + const int32_t be = (int32_t)((d >> 52) & 0x7ff); + const u32 man = (u32)((d & GPUCA_B64_FRAC) >> 29); + const u32 drop = (u32)d & 0x1fffffffu; + u32 bits = 0; + if (be == 0x7ff) { + bits = sign | 0x7f800000u | man | ((d & GPUCA_B64_FRAC) ? 0x400000u : 0u); + } else if (be == 0) { + bits = sign; + } else { + const int32_t e = be - 1023 + 127; + if (e >= 0xff) { + bits = sign | 0x7f800000u; + } else if (e > 0) { + bits = sign | ((u32)e << 23) | man; + if ((drop & 0x10000000u) && ((drop & 0x0fffffffu) || (man & 1u))) { + bits += 1u; + } + } else if (e > -24) { + const u32 full = man | 0x800000u; + const u32 sh = (u32)(1 - e); + const u32 lost = full & ((1u << sh) - 1u); + const u32 halfb = 1u << (sh - 1); + u32 sub = full >> sh; + if (lost > halfb || (lost == halfb && ((sub & 1u) || drop))) { + sub += 1u; + } + bits = sign | sub; + } else { + bits = sign; + } + } + return asf32(bits); +} + +} // namespace binary64_detail + +class GPUdoubleBinary64 +{ + public: + GPUdDefault() GPUdoubleBinary64() = default; + GPUdi() constexpr GPUdoubleBinary64(float v) : mBits(binary64_detail::fromFloat(v)) {} + GPUdi() constexpr operator float() const { return binary64_detail::toFloat(mBits); } + + GPUdi() static constexpr GPUdoubleBinary64 fromBits(binary64_detail::u64 b) { return GPUdoubleBinary64(b, FromBits{}); } + GPUdi() constexpr binary64_detail::u64 bits() const { return mBits; } + + GPUdi() constexpr GPUdoubleBinary64 operator-() const { return fromBits(mBits ^ GPUCA_B64_SIGN); } + GPUdi() constexpr GPUdoubleBinary64 operator+(GPUdoubleBinary64 b) const { return fromBits(binary64_detail::addsub(mBits, b.mBits, false)); } + GPUdi() constexpr GPUdoubleBinary64 operator-(GPUdoubleBinary64 b) const { return fromBits(binary64_detail::addsub(mBits, b.mBits, true)); } + GPUdi() constexpr GPUdoubleBinary64 operator*(GPUdoubleBinary64 b) const { return fromBits(binary64_detail::mul(mBits, b.mBits)); } + GPUdi() constexpr GPUdoubleBinary64 operator/(GPUdoubleBinary64 b) const { return fromBits(binary64_detail::div(mBits, b.mBits)); } + + GPUdi() constexpr GPUdoubleBinary64 operator+(float b) const { return *this + GPUdoubleBinary64(b); } + GPUdi() constexpr GPUdoubleBinary64 operator-(float b) const { return *this - GPUdoubleBinary64(b); } + GPUdi() constexpr GPUdoubleBinary64 operator*(float b) const { return *this * GPUdoubleBinary64(b); } + GPUdi() constexpr GPUdoubleBinary64 operator/(float b) const { return *this / GPUdoubleBinary64(b); } +#ifndef __METAL__ + GPUdi() constexpr GPUdoubleBinary64 operator+(double b) const { return *this + (float)b; } + GPUdi() constexpr GPUdoubleBinary64 operator-(double b) const { return *this - (float)b; } + GPUdi() constexpr GPUdoubleBinary64 operator*(double b) const { return *this * (float)b; } + GPUdi() constexpr GPUdoubleBinary64 operator/(double b) const { return *this / (float)b; } +#endif + + // integral operands: an exact match, so `2 * x` does not sit ambiguously between + // converting the int up and converting *this down + GPUdi() constexpr GPUdoubleBinary64 operator+(int32_t b) const { return *this + fromBits(binary64_detail::fromI32(b)); } + GPUdi() constexpr GPUdoubleBinary64 operator-(int32_t b) const { return *this - fromBits(binary64_detail::fromI32(b)); } + GPUdi() constexpr GPUdoubleBinary64 operator*(int32_t b) const { return *this * fromBits(binary64_detail::fromI32(b)); } + GPUdi() constexpr GPUdoubleBinary64 operator/(int32_t b) const { return *this / fromBits(binary64_detail::fromI32(b)); } + GPUdi() constexpr GPUdoubleBinary64 operator+(uint32_t b) const { return *this + fromBits(binary64_detail::fromU32(b)); } + GPUdi() constexpr GPUdoubleBinary64 operator-(uint32_t b) const { return *this - fromBits(binary64_detail::fromU32(b)); } + GPUdi() constexpr GPUdoubleBinary64 operator*(uint32_t b) const { return *this * fromBits(binary64_detail::fromU32(b)); } + GPUdi() constexpr GPUdoubleBinary64 operator/(uint32_t b) const { return *this / fromBits(binary64_detail::fromU32(b)); } +#ifdef __METAL__ + // the same surface for an object that lives in the constant address space, which + // a generic `this` does not reach + GPUdi() constexpr GPUdoubleBinary64(float v) constant : mBits(binary64_detail::fromFloat(v)) {} + GPUdi() constexpr operator float() constant { return binary64_detail::toFloat(mBits); } + GPUdi() constexpr GPUdoubleBinary64 operator+(GPUdoubleBinary64 b) constant { return fromBits(binary64_detail::addsub(mBits, b.mBits, false)); } + GPUdi() constexpr GPUdoubleBinary64 operator-(GPUdoubleBinary64 b) constant { return fromBits(binary64_detail::addsub(mBits, b.mBits, true)); } + GPUdi() constexpr GPUdoubleBinary64 operator*(GPUdoubleBinary64 b) constant { return fromBits(binary64_detail::mul(mBits, b.mBits)); } + GPUdi() constexpr GPUdoubleBinary64 operator/(GPUdoubleBinary64 b) constant { return fromBits(binary64_detail::div(mBits, b.mBits)); } + GPUdi() constexpr GPUdoubleBinary64 operator+(float b) constant { return *this + GPUdoubleBinary64(b); } + GPUdi() constexpr GPUdoubleBinary64 operator-(float b) constant { return *this - GPUdoubleBinary64(b); } + GPUdi() constexpr GPUdoubleBinary64 operator*(float b) constant { return *this * GPUdoubleBinary64(b); } + GPUdi() constexpr GPUdoubleBinary64 operator/(float b) constant { return *this / GPUdoubleBinary64(b); } + GPUdi() constexpr GPUdoubleBinary64 operator*(int32_t b) constant { return *this * fromBits(binary64_detail::fromI32(b)); } + GPUdi() constexpr GPUdoubleBinary64 operator/(int32_t b) constant { return *this / fromBits(binary64_detail::fromI32(b)); } +#endif + + GPUdi() GPUdoubleBinary64& operator+=(GPUdoubleBinary64 b) { return *this = *this + b; } + GPUdi() GPUdoubleBinary64& operator-=(GPUdoubleBinary64 b) { return *this = *this - b; } + GPUdi() GPUdoubleBinary64& operator*=(GPUdoubleBinary64 b) { return *this = *this * b; } + GPUdi() GPUdoubleBinary64& operator/=(GPUdoubleBinary64 b) { return *this = *this / b; } + + private: + struct FromBits { + }; + GPUdi() constexpr GPUdoubleBinary64(binary64_detail::u64 b, FromBits) : mBits(b) {} + + binary64_detail::u64 mBits; +}; + +GPUdi() constexpr GPUdoubleBinary64 operator+(int32_t a, GPUdoubleBinary64 b) { return GPUdoubleBinary64::fromBits(binary64_detail::fromI32(a)) + b; } +GPUdi() constexpr GPUdoubleBinary64 operator-(int32_t a, GPUdoubleBinary64 b) { return GPUdoubleBinary64::fromBits(binary64_detail::fromI32(a)) - b; } +GPUdi() constexpr GPUdoubleBinary64 operator*(int32_t a, GPUdoubleBinary64 b) { return GPUdoubleBinary64::fromBits(binary64_detail::fromI32(a)) * b; } +GPUdi() constexpr GPUdoubleBinary64 operator/(int32_t a, GPUdoubleBinary64 b) { return GPUdoubleBinary64::fromBits(binary64_detail::fromI32(a)) / b; } +GPUdi() constexpr GPUdoubleBinary64 operator+(uint32_t a, GPUdoubleBinary64 b) { return GPUdoubleBinary64::fromBits(binary64_detail::fromU32(a)) + b; } +GPUdi() constexpr GPUdoubleBinary64 operator-(uint32_t a, GPUdoubleBinary64 b) { return GPUdoubleBinary64::fromBits(binary64_detail::fromU32(a)) - b; } +GPUdi() constexpr GPUdoubleBinary64 operator*(uint32_t a, GPUdoubleBinary64 b) { return GPUdoubleBinary64::fromBits(binary64_detail::fromU32(a)) * b; } +GPUdi() constexpr GPUdoubleBinary64 operator/(uint32_t a, GPUdoubleBinary64 b) { return GPUdoubleBinary64::fromBits(binary64_detail::fromU32(a)) / b; } +GPUdi() constexpr GPUdoubleBinary64 operator+(float a, GPUdoubleBinary64 b) { return GPUdoubleBinary64(a) + b; } +GPUdi() constexpr GPUdoubleBinary64 operator-(float a, GPUdoubleBinary64 b) { return GPUdoubleBinary64(a) - b; } +GPUdi() constexpr GPUdoubleBinary64 operator*(float a, GPUdoubleBinary64 b) { return GPUdoubleBinary64(a) * b; } +GPUdi() constexpr GPUdoubleBinary64 operator/(float a, GPUdoubleBinary64 b) { return GPUdoubleBinary64(a) / b; } +#ifndef __METAL__ +GPUdi() constexpr GPUdoubleBinary64 operator+(double a, GPUdoubleBinary64 b) { return GPUdoubleBinary64((float)a) + b; } +GPUdi() constexpr GPUdoubleBinary64 operator-(double a, GPUdoubleBinary64 b) { return GPUdoubleBinary64((float)a) - b; } +GPUdi() constexpr GPUdoubleBinary64 operator*(double a, GPUdoubleBinary64 b) { return GPUdoubleBinary64((float)a) * b; } +GPUdi() constexpr GPUdoubleBinary64 operator/(double a, GPUdoubleBinary64 b) { return GPUdoubleBinary64((float)a) / b; } +#endif + +// rounds once, as `someFloat += someDouble` does on the host +#ifdef __METAL__ +GPUdi() thread float& operator+=(thread float& a, GPUdoubleBinary64 b) { return a = (float)(GPUdoubleBinary64(a) + b); } +GPUdi() device float& operator+=(device float& a, GPUdoubleBinary64 b) { return a = (float)(GPUdoubleBinary64(a) + b); } +GPUdi() threadgroup float& operator+=(threadgroup float& a, GPUdoubleBinary64 b) { return a = (float)(GPUdoubleBinary64(a) + b); } +#else +GPUdi() float& operator+=(float& a, GPUdoubleBinary64 b) { return a = (float)(GPUdoubleBinary64(a) + b); } +#endif + +namespace binary64_detail +{ +// sin and cos are fdlibm's __kernel_sin / __kernel_cos and the medium-range +// branch of __ieee754_rem_pio2. The coefficients are spelled as bit patterns +// because MSL has no double literals. The argument reduction is exact for +// |x| <= 2^19 * pi/2; beyond that the accuracy degrades gracefully. +GPUCA_B64_OP GPUdoubleBinary64 kernelSin(GPUdoubleBinary64 x, GPUdoubleBinary64 y, bool iy) +{ + typedef GPUdoubleBinary64 b64; + const b64 S1 = b64::fromBits(0xBFC5555555555549ULL); + const b64 S2 = b64::fromBits(0x3F8111111110F8A6ULL); + const b64 S3 = b64::fromBits(0xBF2A01A019C161D5ULL); + const b64 S4 = b64::fromBits(0x3EC71DE357B1FE7DULL); + const b64 S5 = b64::fromBits(0xBE5AE5E68A2B9CEBULL); + const b64 S6 = b64::fromBits(0x3DE5D93A5ACFD57CULL); + const b64 z = x * x; + const b64 v = z * x; + const b64 r = S2 + z * (S3 + z * (S4 + z * (S5 + z * S6))); + if (!iy) { + return x + v * (S1 + z * r); + } + return x - ((z * (b64::fromBits(0x3FE0000000000000ULL) * y - v * r) - y) - v * S1); +} + +GPUCA_B64_OP GPUdoubleBinary64 kernelCos(GPUdoubleBinary64 x, GPUdoubleBinary64 y) +{ + typedef GPUdoubleBinary64 b64; + const b64 C1 = b64::fromBits(0x3FA555555555554CULL); + const b64 C2 = b64::fromBits(0xBF56C16C16C15177ULL); + const b64 C3 = b64::fromBits(0x3EFA01A019CB1590ULL); + const b64 C4 = b64::fromBits(0xBE927E4F809C52ADULL); + const b64 C5 = b64::fromBits(0x3E21EE9EBDB4B1C4ULL); + const b64 C6 = b64::fromBits(0xBDA8FAE9BE8838D4ULL); + const b64 one = b64::fromBits(0x3FF0000000000000ULL); + const b64 oneHalf = b64::fromBits(0x3FE0000000000000ULL); + const u32 ix = (u32)(x.bits() >> 32) & 0x7fffffffu; + const b64 z = x * x; + const b64 r = z * (C1 + z * (C2 + z * (C3 + z * (C4 + z * (C5 + z * C6))))); + if (ix < 0x3FD33333u) { // |x| < 0.3 + return one - (oneHalf * z - (z * r - x * y)); + } + const b64 qx = (ix > 0x3FE90000u) ? b64::fromBits(0x3FD2000000000000ULL) : b64::fromBits((u64)(ix - 0x00200000u) << 32); // 0.28125, else |x| / 4 + return (one - qx) - ((oneHalf * z - qx) - (z * r - x * y)); +} + +GPUCA_B64_ALWAYS int32_t truncToInt32(GPUdoubleBinary64 x) +{ + const u64 b = x.bits(); + const int32_t e = (int32_t)((b >> 52) & 0x7ff) - 1023; + if (e < 0) { + return 0; + } + const u64 m = (b & GPUCA_B64_FRAC) | GPUCA_B64_IMPL; + const int32_t v = (int32_t)(e >= 52 ? (m << (e - 52)) : (m >> (52 - e))); + return (b & GPUCA_B64_SIGN) ? -v : v; +} + +struct SinCosPair { + GPUdoubleBinary64 s, c; +}; + +GPUCA_B64_OP SinCosPair sincos(GPUdoubleBinary64 x) +{ + typedef GPUdoubleBinary64 b64; + const u32 ix = (u32)(x.bits() >> 32) & 0x7fffffffu; + SinCosPair out; + if (ix >= 0x7FF00000u) { // inf or NaN + out.s = out.c = b64::fromBits(GPUCA_B64_DNAN); + return out; + } + + if (ix < 0x3E400000u) { // |x| < 2^-27, where the sign of a zero x has to survive + out.s = x; + out.c = b64::fromBits(0x3FF0000000000000ULL); + return out; + } + + b64 y0 = x, y1 = b64::fromBits(0ULL); + int32_t n = 0; + const bool reduced = ix > 0x3FE921FBu; // |x| > pi/4 + if (reduced) { + const b64 t = b64::fromBits(x.bits() & ~GPUCA_B64_SIGN); + n = truncToInt32(t * b64::fromBits(0x3FE45F306DC9C883ULL) + b64::fromBits(0x3FE0000000000000ULL)); + const b64 fn = b64((float)n); + // Cody-Waite with pi/2 split over three terms, good to 151 bits + b64 r = t - fn * b64::fromBits(0x3FF921FB54400000ULL); + b64 s = r; + b64 w = fn * b64::fromBits(0x3DD0B4611A600000ULL); + r = s - w; + w = fn * b64::fromBits(0x3BA3198A2E037073ULL) - ((s - r) - w); + s = r; + w = fn * b64::fromBits(0x3BA3198A2E000000ULL); + r = s - w; + w = fn * b64::fromBits(0x397B839A252049C1ULL) - ((s - r) - w); + y0 = r - w; + y1 = (r - y0) - w; + if (x.bits() & GPUCA_B64_SIGN) { + y0 = -y0; + y1 = -y1; + n = -n; + } + } + + switch (n & 3) { + case 0: + out.s = kernelSin(y0, y1, reduced); + out.c = kernelCos(y0, y1); + break; + case 1: + out.s = kernelCos(y0, y1); + out.c = -kernelSin(y0, y1, reduced); + break; + case 2: + out.s = -kernelSin(y0, y1, reduced); + out.c = -kernelCos(y0, y1); + break; + default: + out.s = -kernelCos(y0, y1); + out.c = kernelSin(y0, y1, reduced); + break; + } + return out; +} +} // namespace binary64_detail + +} // namespace o2::gpu + +#endif // GPUCOMMONDOUBLEBINARY64_H diff --git a/GPU/Common/GPUCommonMath.h b/GPU/Common/GPUCommonMath.h index 7a78a5881dcfa..8ab250f4cfd98 100644 --- a/GPU/Common/GPUCommonMath.h +++ b/GPU/Common/GPUCommonMath.h @@ -105,7 +105,11 @@ class GPUCommonMath GPUd() constexpr static bool Finite(float x); GPUd() constexpr static bool IsNaN(float x); #ifndef __FAST_MATH__ +#ifdef __METAL__ // MSL has no nan(uint) + GPUd() constexpr static float QuietNaN() { return __builtin_nanf(""); } +#else GPUd() constexpr static float QuietNaN() { return GPUCA_CHOICE(std::numeric_limits::quiet_NaN(), __builtin_nanf(""), nan(0u)); } +#endif #endif GPUd() constexpr static uint32_t Clz(uint32_t val); GPUd() constexpr static uint32_t Ctz(uint32_t val); @@ -245,7 +249,11 @@ GPUdi() float2 GPUCommonMath::MakeFloat2(float x, float y) } GPUdi() constexpr float GPUCommonMath::Modf(float x, float y) { return GPUCA_CHOICE(fmodf(x, y), fmodf(x, y), fmod(x, y)); } +#ifdef __METAL__ // MSL has no remainder(); this is its definition +GPUhdi() float GPUCommonMath::Remainderf(float x, float y) { return x - y * rint(x / y); } +#else GPUhdi() float GPUCommonMath::Remainderf(float x, float y) { return GPUCA_CHOICE(std::remainderf(x, y), remainderf(x, y), remainder(x, y)); } +#endif GPUdi() uint32_t GPUCommonMath::Float2UIntReint(const float& x) { @@ -302,12 +310,20 @@ GPUhdi() void GPUCommonMath::SinCos(float x, float& s, float& c) __sincosf(x, &s, &c); #elif !defined(GPUCA_GPUCODE_DEVICE) && (defined(__GNU_SOURCE__) || defined(_GNU_SOURCE) || defined(GPUCA_GPUCODE)) sincosf(x, &s, &c); +#else +#ifdef __METAL__ // MSL's sincos returns sin and takes cos by thread reference, + // so it cannot write straight through a generic one + float metalCos; + s = sincos(x, metalCos); + c = metalCos; #else GPUCA_CHOICE((void)((s = sinf(x)) + (c = cosf(x))), sincosf(x, &s, &c), s = sincos(x, &c)); +#endif #endif ) // clang-format on } +#ifndef __METAL__ GPUhdi() void GPUCommonMath::SinCosd(double x, double& s, double& c) { #if !defined(GPUCA_GPUCODE_DEVICE) && defined(__APPLE__) @@ -318,6 +334,7 @@ GPUhdi() void GPUCommonMath::SinCosd(double x, double& s, double& c) GPUCA_CHOICE((void)((s = sin(x)) + (c = cos(x))), sincos(x, &s, &c), s = sincos(x, &c)); #endif } +#endif GPUdi() constexpr uint32_t GPUCommonMath::Clz(uint32_t x) { @@ -444,11 +461,13 @@ GPUhdi() constexpr float GPUCommonMath::Abs(float x) return GPUCA_CHOICE(fabsf(x), fabsf(x), fabs(x)); } +#ifndef __METAL__ template <> GPUhdi() constexpr double GPUCommonMath::Abs(double x) { return GPUCA_CHOICE(fabs(x), fabs(x), fabs(x)); } +#endif template <> GPUhdi() constexpr int32_t GPUCommonMath::Abs(int32_t x) @@ -456,6 +475,17 @@ GPUhdi() constexpr int32_t GPUCommonMath::Abs(int32_t x) return GPUCA_CHOICE(abs(x), abs(x), abs(x)); } +#ifdef __METAL__ +// The counters these operate on are plain integers in the transferred structures, +// as they are for CUDA and HIP. MSL's atomic operations want metal::atomic, which +// has the same size and alignment; the overloads keep the address space, which a +// generic pointer would not carry into atomic_*_explicit. +template +GPUdi() threadgroup metal::atomic* GPUCommonMathMetalAtomic(threadgroup T* p) { return reinterpret_cast*>(p); } +template +GPUdi() device metal::atomic* GPUCommonMathMetalAtomic(T* p) { return (device metal::atomic*)p; } +#endif + template GPUdi() uint32_t GPUCommonMath::AtomicExchInternal(S* addr, T val) { @@ -466,7 +496,7 @@ GPUdi() uint32_t GPUCommonMath::AtomicExchInternal(S* addr, T val) #elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__)) return ::atomicExch(addr, val); #elif defined(GPUCA_GPUCODE) && defined(__METAL__) - return atomic_exchange_explicit(addr, val, memory_order_relaxed); + return atomic_exchange_explicit(GPUCommonMathMetalAtomic(addr), val, memory_order_relaxed); #elif defined(WITH_OPENMP) uint32_t old; __atomic_exchange(addr, &val, &old, __ATOMIC_SEQ_CST); @@ -486,7 +516,7 @@ GPUdi() bool GPUCommonMath::AtomicCASInternal(S* addr, T cmp, T val) #elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__)) return ::atomicCAS(addr, cmp, val) == cmp; #elif defined(GPUCA_GPUCODE) && defined(__METAL__) - return atomic_compare_exchange_weak_explicit(addr, &cmp, val, memory_order_relaxed, memory_order_relaxed); + return atomic_compare_exchange_weak_explicit(GPUCommonMathMetalAtomic(addr), &cmp, val, memory_order_relaxed, memory_order_relaxed); #elif defined(WITH_OPENMP) return __atomic_compare_exchange(addr, &cmp, &val, true, __ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST); #else @@ -504,7 +534,7 @@ GPUdi() uint32_t GPUCommonMath::AtomicAddInternal(S* addr, T val) #elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__)) return ::atomicAdd(addr, val); #elif defined(GPUCA_GPUCODE) && defined(__METAL__) - return atomic_fetch_add_explicit(addr, val, memory_order_relaxed); + return atomic_fetch_add_explicit(GPUCommonMathMetalAtomic(addr), val, memory_order_relaxed); #elif defined(WITH_OPENMP) return __atomic_add_fetch(addr, val, __ATOMIC_SEQ_CST) - val; #else @@ -522,7 +552,7 @@ GPUdi() void GPUCommonMath::AtomicMaxInternal(S* addr, T val) #elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__)) ::atomicMax(addr, val); #elif defined(GPUCA_GPUCODE) && defined(__METAL__) - atomic_fetch_max_explicit(addr, val, memory_order_relaxed); + atomic_fetch_max_explicit(GPUCommonMathMetalAtomic(addr), val, memory_order_relaxed); #else S current; while ((current = *(volatile S*)addr) < val && !AtomicCASInternal(addr, current, val)) { @@ -540,7 +570,7 @@ GPUdi() void GPUCommonMath::AtomicMinInternal(S* addr, T val) #elif defined(GPUCA_GPUCODE) && (defined(__CUDACC__) || defined(__HIPCC__)) ::atomicMin(addr, val); #elif defined(GPUCA_GPUCODE) && defined(__METAL__) - atomic_fetch_min_explicit(addr, val, memory_order_relaxed); + atomic_fetch_min_explicit(GPUCommonMathMetalAtomic(addr), val, memory_order_relaxed); #else S current; while ((current = *(volatile S*)addr) > val && !AtomicCASInternal(addr, current, val)) { diff --git a/GPU/Common/GPUCommonTypeTraits.h b/GPU/Common/GPUCommonTypeTraits.h index 3f83e151b0f33..6d72f164771bc 100644 --- a/GPU/Common/GPUCommonTypeTraits.h +++ b/GPU/Common/GPUCommonTypeTraits.h @@ -114,7 +114,13 @@ struct is_pointer_t { }; template struct is_pointer { +#ifdef __METAL__ + // A bare T* partial specialization does not match a pointer type deduced from + // an argument, which carries its address space; metal::is_pointer does. + enum { value = metal::is_pointer::value }; +#else enum { value = is_pointer_t::type>::value }; +#endif }; template diff --git a/GPU/Common/test/testGPUsortCUDA.cu b/GPU/Common/test/testGPUsortCUDA.cu index b19235f9e8c6b..95464dcc96c0d 100644 --- a/GPU/Common/test/testGPUsortCUDA.cu +++ b/GPU/Common/test/testGPUsortCUDA.cu @@ -96,12 +96,12 @@ __global__ void sortInThreadWithOperator(float* data, size_t dataLength) __global__ void sortInBlock(float* data, size_t dataLength) { - o2::gpu::CAAlgo::sortInBlock(data, data + dataLength); + o2::gpu::CAAlgo::sortInBlock(blockDim.x, threadIdx.x, data, data + dataLength); } __global__ void sortInBlockWithOperator(float* data, size_t dataLength) { - o2::gpu::CAAlgo::sortInBlock(data, data + dataLength, [](float a, float b) { return a < b; }); + o2::gpu::CAAlgo::sortInBlock(blockDim.x, threadIdx.x, data, data + dataLength, [](float a, float b) { return a < b; }); } /////////////////////////////////////////////////////////////// diff --git a/GPU/GPUTracking/Base/GPUGeneralKernels.cxx b/GPU/GPUTracking/Base/GPUGeneralKernels.cxx index e1a3ce69dd8df..d27b778701ea1 100644 --- a/GPU/GPUTracking/Base/GPUGeneralKernels.cxx +++ b/GPU/GPUTracking/Base/GPUGeneralKernels.cxx @@ -19,12 +19,12 @@ using namespace o2::gpu; template <> GPUdii() void GPUMemClean16::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors, GPUglobalref() void* ptr, uint64_t size) { - const uint64_t stride = get_global_size(0); + const uint64_t stride = (nBlocks * nThreads); int4 i0; i0.x = i0.y = i0.z = i0.w = 0; int4* ptra = (int4*)ptr; uint64_t len = (size + sizeof(int4) - 1) / sizeof(int4); - for (uint64_t i = get_global_id(0); i < len; i += stride) { + for (uint64_t i = (iBlock * nThreads + iThread); i < len; i += stride) { ptra[i] = i0; } } @@ -32,8 +32,8 @@ GPUdii() void GPUMemClean16::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_ template <> GPUdii() void GPUitoa::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors, GPUglobalref() int32_t* ptr, uint64_t size) { - const uint64_t stride = get_global_size(0); - for (uint64_t i = get_global_id(0); i < size; i += stride) { + const uint64_t stride = (nBlocks * nThreads); + for (uint64_t i = (iBlock * nThreads + iThread); i < size; i += stride) { ptr[i] = i; } } diff --git a/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h b/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h index cc1c62bed507d..0887cadd7338d 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h +++ b/GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h @@ -63,8 +63,17 @@ #define GPUCA_ATTRRES(...) GPUCA_M_EXPAND(GPUCA_M_CAT(GPUCA_ATTRRES_, GPUCA_M_FIRST(__VA_ARGS__)))(__VA_ARGS__) // GPU Kernel entry point +// MSL requires every kernel parameter to carry an attribute, and supplies the +// grid dimensions the same way, so the backend gets to shape both ends of the +// parameter list. +#ifndef GPUCA_KRNL_SECTOR_ARG +#define GPUCA_KRNL_SECTOR_ARG int32_t _iSector_internal +#endif +#ifndef GPUCA_KRNL_GRID_ARGS +#define GPUCA_KRNL_GRID_ARGS +#endif #define GPUCA_KRNLGPU_DEF(x_class, x_attributes, x_arguments, ...) \ - GPUg() void GPUCA_ATTRRES(GPUCA_M_STRIP(x_attributes)) GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))(GPUCA_CONSMEM_PTR int32_t _iSector_internal GPUCA_M_STRIP(x_arguments)) + GPUg() void GPUCA_ATTRRES(GPUCA_M_STRIP(x_attributes)) GPUCA_M_CAT(krnl_, GPUCA_M_KRNL_NAME(x_class))(GPUCA_CONSMEM_PTR GPUCA_KRNL_SECTOR_ARG GPUCA_M_STRIP(x_arguments) GPUCA_KRNL_GRID_ARGS) #ifdef GPUCA_KRNL_DEFONLY #define GPUCA_KRNLGPU(...) GPUCA_KRNLGPU_DEF(__VA_ARGS__); diff --git a/GPU/GPUTracking/Base/GPUReconstructionThreading.h b/GPU/GPUTracking/Base/GPUReconstructionThreading.h index 374c7545e65da..f03ac3c06102a 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionThreading.h +++ b/GPU/GPUTracking/Base/GPUReconstructionThreading.h @@ -35,25 +35,25 @@ struct GPUReconstructionThreading { #endif -#define GPUCA_TBB_KERNEL_LOOP_HOST(rec, vartype, varname, iEnd, code) \ - for (vartype varname = get_global_id(0); varname < iEnd; varname += get_global_size(0)) { \ - code \ +#define GPUCA_TBB_KERNEL_LOOP_HOST(rec, nBlocks, nThreads, iBlock, iThread, vartype, varname, iEnd, code) \ + for (vartype varname = (iBlock) * (nThreads) + (iThread); varname < iEnd; varname += (nBlocks) * (nThreads)) { \ + code \ } #ifdef GPUCA_GPUCODE #define GPUCA_TBB_KERNEL_LOOP GPUCA_TBB_KERNEL_LOOP_HOST #else -#define GPUCA_TBB_KERNEL_LOOP(rec, vartype, varname, iEnd, code) \ - if (!rec.GetProcessingSettings().inKernelParallel) { \ - rec.mThreading->activeThreads->execute([&] { \ - tbb::parallel_for(tbb::blocked_range(get_global_id(0), iEnd, get_global_size(0)), [&](const tbb::blocked_range& _r_internal) { \ - for (vartype varname = _r_internal.begin(); varname < _r_internal.end(); varname += get_global_size(0)) { \ - code \ - } \ - }); \ - }); \ - } else { \ - GPUCA_TBB_KERNEL_LOOP_HOST(rec, vartype, varname, iEnd, code) \ +#define GPUCA_TBB_KERNEL_LOOP(rec, nBlocks, nThreads, iBlock, iThread, vartype, varname, iEnd, code) \ + if (!rec.GetProcessingSettings().inKernelParallel) { \ + rec.mThreading->activeThreads->execute([&] { \ + tbb::parallel_for(tbb::blocked_range((iBlock) * (nThreads) + (iThread), iEnd, (nBlocks) * (nThreads)), [&](const tbb::blocked_range& _r_internal) { \ + for (vartype varname = _r_internal.begin(); varname < _r_internal.end(); varname += (nBlocks) * (nThreads)) { \ + code \ + } \ + }); \ + }); \ + } else { \ + GPUCA_TBB_KERNEL_LOOP_HOST(rec, nBlocks, nThreads, iBlock, iThread, vartype, varname, iEnd, code) \ } #endif diff --git a/GPU/GPUTracking/Base/hip/test/testGPUsortHIP.hip b/GPU/GPUTracking/Base/hip/test/testGPUsortHIP.hip index ed13124ef65df..5758faaebebda 100644 --- a/GPU/GPUTracking/Base/hip/test/testGPUsortHIP.hip +++ b/GPU/GPUTracking/Base/hip/test/testGPUsortHIP.hip @@ -104,12 +104,12 @@ __global__ void sortInThreadWithOperator(float* data, size_t dataLength) __global__ void sortInBlock(float* data, size_t dataLength) { - o2::gpu::CAAlgo::sortInBlock(data, data + dataLength); + o2::gpu::CAAlgo::sortInBlock(blockDim.x, threadIdx.x, data, data + dataLength); } __global__ void sortInBlockWithOperator(float* data, size_t dataLength) { - o2::gpu::CAAlgo::sortInBlock(data, data + dataLength, [](float a, float b) { return a < b; }); + o2::gpu::CAAlgo::sortInBlock(blockDim.x, threadIdx.x, data, data + dataLength, [](float a, float b) { return a < b; }); } /////////////////////////////////////////////////////////////// diff --git a/GPU/GPUTracking/Base/metal/CMakeLists.txt b/GPU/GPUTracking/Base/metal/CMakeLists.txt new file mode 100644 index 0000000000000..9c9b7e4fd2b15 --- /dev/null +++ b/GPU/GPUTracking/Base/metal/CMakeLists.txt @@ -0,0 +1,110 @@ +# Copyright 2019-2020 CERN and copyright holders of ALICE O2. +# See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +# All rights not expressly granted are reserved. +# +# This software is distributed under the terms of the GNU General Public +# License v3 (GPL Version 3), copied verbatim in the file "COPYING". +# +# In applying this license CERN does not waive the privileges and immunities +# granted to it by virtue of its status as an Intergovernmental Organization +# or submit itself to any jurisdiction. + +set(MODULE GPUTrackingMETAL) +enable_language(ASM) + +message(STATUS "Building GPUTracking with Metal support") + +# convenience variables +if(ALIGPU_BUILD_TYPE STREQUAL "Standalone") + set(GPUDIR ${CMAKE_SOURCE_DIR}/../) +else() + set(GPUDIR ${CMAKE_SOURCE_DIR}/GPU/GPUTracking) +endif() +set(METAL_SRC ${GPUDIR}/Base/metal/GPUReconstructionMETAL.metal) +set(METAL_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionMetalCode) + +# MSL 4.1 is the first version with a generic address space: earlier versions +# reject an unannotated pointer or `this` outright, which GPUCommonDefAPI.h +# relies on for GPUgeneric() and GPUdDefault(). +set(METAL_FLAGS -std=metal4.1 ${GPUCA_METAL_DENORMALS_FLAGS}) +if(GPUCA_DETERMINISTIC_MODE GREATER_EQUAL ${GPUCA_DETERMINISTIC_MODE_MAP_NO_FAST_MATH}) + set(METAL_FLAGS ${METAL_FLAGS} ${GPUCA_METAL_NO_FAST_MATH_FLAGS}) +endif() +set(METAL_DEFINES "-D$,$-D>" + "-I$,EXCLUDE,^/usr/include/?>,$-I>" + -I${CMAKE_SOURCE_DIR}/Detectors/TRD/base/src + -I${CMAKE_SOURCE_DIR}/Detectors/Base/src + -I${CMAKE_SOURCE_DIR}/DataFormats/Reconstruction/src +) + +set(SRCS GPUReconstructionMetal.mm GPUReconstructionMetalKernels.mm) +set(HDRS GPUReconstructionMetal.h GPUReconstructionMetalIncludesHost.h) + +if(ALIGPU_BUILD_TYPE STREQUAL "O2") + o2_add_library(${MODULE} + SOURCES ${SRCS} + PUBLIC_LINK_LIBRARIES O2::GPUTracking + TARGETVARNAME targetName) + + target_link_libraries(${targetName} PUBLIC ${METAL_FRAMEWORKS}) + + target_compile_definitions(${targetName} PRIVATE $) + # the compile_defitions are not propagated automatically on purpose (they are + # declared PRIVATE) so we are not leaking them outside of the GPU** + # directories +endif() + +if(ALIGPU_BUILD_TYPE STREQUAL "Standalone") + add_library(${MODULE} SHARED ${SRCS}) + target_link_libraries(${MODULE} GPUTracking) + install(TARGETS ${MODULE}) + set(targetName ${MODULE}) +endif() + +if(METAL_ENABLED) # BUILD Metal source code for runtime compilation target + + # executes clang to preprocess + add_custom_command( + OUTPUT ${METAL_BIN}.metal + COMMAND xcrun -sdk macosx metal + -Wno-unused-command-line-argument + ${METAL_FLAGS} + ${METAL_DEFINES} + -MD -MT ${METAL_BIN}.src -MF ${METAL_BIN}.src.d + -E -P ${METAL_SRC} > ${METAL_BIN}.metal + DEPENDS ${METAL_SRC} + DEPFILE ${METAL_BIN}.src.d + COMMAND_EXPAND_LISTS + COMMENT "Preparing Metal source file for run time compilation ${METAL_BIN}.metal") + + # Create the ir + add_custom_command( + OUTPUT ${METAL_BIN}.ir + COMMAND xcrun -sdk macosx metal + -Wno-unused-command-line-argument + -Wno-c++17-extensions + -ferror-limit=10000 + ${METAL_FLAGS} + ${METAL_DEFINES} + ${METAL_BIN}.metal + -o ${METAL_BIN}.ir + DEPENDS ${METAL_BIN}.metal + COMMAND_EXPAND_LISTS + COMMENT "Preparing Metal intermediate representation for run time compilation ${METAL_BIN}.ir") + + add_custom_target(metal_preprocessed_code ALL DEPENDS ${METAL_BIN}.metal COMMENT "Needed to inject dependency on its creation") + add_custom_target(metal_intermediate_representation ALL DEPENDS ${METAL_BIN}.ir COMMENT "Needed to inject dependency on its creation") + + # Pack the compiled library into __DATA,__gpu_resource during final link. This + # way we do not need to create an intermediate object. Compiling the source at + # run time is not an option: the driver's compiler service dies on it. + target_link_options(${targetName} + PRIVATE + "-Wl,-sectcreate,__DATA,__gpu_resource,${METAL_BIN}.ir") + add_dependencies(${targetName} metal_preprocessed_code) + add_dependencies(${targetName} metal_intermediate_representation) +endif() + +install(FILES ${HDRS} DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/GPU) + +target_compile_definitions(${targetName} PRIVATE GPUCA_METAL_BUILD_FLAGS=$ ) diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal b/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal new file mode 100644 index 0000000000000..62a7430a1570a --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal @@ -0,0 +1,87 @@ +// Copyright 2019-2025 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +/// \file GPUReconstructionMETAL.metal + +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wgnu-zero-variadic-macro-arguments" +// clang-format off + +// --- Backend selection ------------------------------------------------------- +#define GPUCA_GPUTYPE_METAL 1 + +// --- Metal stdlib ------------------------------------------------------------ +#include +// MSL rejects derived classes outside of this pragma, and the kernels are +// class-based throughout. metal_stdlib itself uses it in 46 paired places. +#pragma METAL internals : enable +using namespace metal; + +// --- OpenCL compatibility shims --------------------------------------------- + +// Address space aliases (match OpenCL vernacular used by the project); constant +// is spelled the same in MSL +#define global device +#define local threadgroup + +#ifndef M_PI +#define M_PI 3.1415926535f +#endif + +// Disable assertions inside GPU code (same as OpenCL variant) +#ifdef assert +# undef assert +#endif +#define assert(param) + +// --- double ------------------------------------------------------------------ +// MSL has no double. GPUdoubleBinary64 is IEEE-754 binary64 in software, with the +// same eight bytes in the same order, so the keyword can simply name it and the +// shared code needs no separate spelling. Must come after metal_stdlib, which +// uses the token itself. +#include "GPUCommonDoubleBinary64.h" +#define double o2::gpu::GPUdoubleBinary64 +#include "GPUCommonDouble.h" + +// --- Project headers --------------------------------------------------------- +#include "GPUCommonDef.h" +#include "GPUCommonTypeTraits.h" +#include "GPUCommonArray.h" + +#include "GPUConstantMem.h" +#include "GPUReconstructionIncludesDeviceAll.h" + +// --- Kernel list expansion --------------------------------------------------- +#define GPUCA_KRNL(...) GPUCA_KRNLGPU(__VA_ARGS__) + +// --- Constant memory + global heap plumbing --------------------------------- +// The heap and the constant memory arrive as buffer(0) and buffer(1). The latter +// is untyped because a buffer of GPUConstantMem, which has base classes, is not +// a valid kernel argument type. +#define GPUCA_CONSMEM_PTR \ + device char* gpu_mem [[buffer(0)]], \ + device char* pConstantRaw [[buffer(1)]], +#define GPUCA_CONSMEM (*(device GPUConstantMem*)pConstantRaw) + +// Every kernel parameter needs an attribute, so the sector index arrives as a +// buffer rather than by value, and the grid dimensions come in at the end, where +// GPUCommonDefAPI.h's get_group_id() and friends pick them up. +#define GPUCA_KRNL_SECTOR_ARG constant int32_t& _iSector_internal [[buffer(2)]] +#define GPUCA_KRNL_GRID_ARGS \ + , uint _metalTgIg [[threadgroup_position_in_grid]] \ + , uint _metalTiTg [[thread_position_in_threadgroup]] \ + , uint _metalTPerTg [[threads_per_threadgroup]] \ + , uint _metalTgPerG [[threadgroups_per_grid]] + +#include "GPUReconstructionKernelList.h" + +// clang-format on +#pragma clang diagnostic pop diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h new file mode 100644 index 0000000000000..66cf5ab7bf121 --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.h @@ -0,0 +1,67 @@ +// Copyright 2019-2025 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +#ifndef GPURECONSTRUCTIONMETAL_H +#define GPURECONSTRUCTIONMETAL_H + +#include "GPUReconstructionDeviceBase.h" + +extern "C" o2::gpu::GPUReconstruction* GPUReconstruction_Create_METAL(const o2::gpu::GPUSettingsDeviceBackend& cfg); + +namespace o2::gpu +{ +struct GPUReconstructionMetalInternals; + +class GPUReconstructionMetal : public GPUReconstructionProcessing::KernelInterface +{ + public: + GPUReconstructionMetal(const GPUSettingsDeviceBackend& cfg); + ~GPUReconstructionMetal() override; + + template + void runKernelBackend(const krnlSetupTime& _xyz, const Args&... args); + + protected: + int32_t InitDevice_Runtime() override; + int32_t ExitDevice_Runtime() override; + + virtual int32_t GPUChkErrInternal(const int64_t error, const char* file, int32_t line) const override; + + void SynchronizeGPU() override; + int32_t GPUDebug(const char* state = "UNKNOWN", int32_t stream = -1, bool force = false) override; + void SynchronizeStream(int32_t stream) override; + void SynchronizeEvents(deviceEvent* evList, int32_t nEvents = 1) override; + void StreamWaitForEvents(int32_t stream, deviceEvent* evList, int32_t nEvents = 1) override; + bool IsEventDone(deviceEvent* evList, int32_t nEvents = 1) override; + + size_t WriteToConstantMemory(size_t offset, const void* src, size_t size, int32_t stream = -1, deviceEvent* ev = nullptr) override; + size_t GPUMemCpy(void* dst, const void* src, size_t size, int32_t stream, int32_t toGPU, deviceEvent* ev = nullptr, deviceEvent* evList = nullptr, int32_t nEvents = 1) override; + void ReleaseEvent(deviceEvent ev) override; + void RecordMarker(deviceEvent* ev, int32_t stream) override; + + template + int32_t AddKernel(); + + GPUReconstructionMetalInternals* mInternals; + float mOclVersion; + + template + S& getKernelObject(); + + int32_t GetMetalPrograms(); + + private: + int32_t AddKernels(); +}; + +} // namespace o2::gpu + +#endif // GPURECONSTRUCTIONMETAL_H diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm new file mode 100644 index 0000000000000..8448a9d86184b --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm @@ -0,0 +1,414 @@ +// Copyright 2019-2025 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +#include "GPUReconstructionMetal.h" +#include "GPUConstantMem.h" +#include "GPUDefParametersLoad.inc" +#include "GPUReconstructionMetalIncludesHost.h" + +#include +#include + +#include +#include +#include // _mh_execute_header + +#define GPUErrorReturn(...) \ + { \ + GPUError(__VA_ARGS__); \ + return (1); \ + } + +#include "utils/qGetLdBinarySymbols.h" +QGET_LD_BINARY_SYMBOLS(GPUReconstructionMetalCode_src); + +GPUReconstruction* GPUReconstruction_Create_METAL(const GPUSettingsDeviceBackend& cfg) { return new GPUReconstructionMetal(cfg); } + +GPUReconstructionMetal::GPUReconstructionMetal(const GPUSettingsDeviceBackend& cfg) : GPUReconstructionProcessing::KernelInterface(cfg, sizeof(GPUReconstructionDeviceBase)) +{ + if (mMaster == nullptr) { + mInternals = new GPUReconstructionMetalInternals; + *mParDevice = o2::gpu::internal::GPUDefParametersLoad(); + } + mDeviceBackendSettings->deviceType = DeviceType::METAL; +} + +GPUReconstructionMetal::~GPUReconstructionMetal() +{ + Exit(); // Make sure we destroy everything (in particular the ITS tracker) before we exit + if (mMaster == nullptr) { + delete mInternals; + } +} + +int32_t GPUReconstructionMetal::InitDevice_Runtime() +{ + // Propagate processing settings to PoCL runtime. + // Won't affect other OpenCL runtimes. + if (int nThreads = mProcessingSettings->nHostThreads; nThreads > 0) { + auto nThreadsStr = std::to_string(nThreads); + setenv("PMETAL_CPU_MAX_CU_COUNT", nThreadsStr.c_str(), 1); + } + + if (mMaster == nullptr) { + mInternals->device = MTLCreateSystemDefaultDevice(); + + int64_t deviceGlobalMem, deviceLocalMem; + MTLSize deviceMaxWorkGroup = mInternals->device.maxThreadsPerThreadgroup; + + std::string device_name = [mInternals->device.name UTF8String]; + // On Apple Silicon, treat recommended working set as an upper bound + deviceGlobalMem = mInternals->device.recommendedMaxWorkingSetSize; + deviceLocalMem = mInternals->device.maxThreadgroupMemoryLength; + if (GetProcessingSettings().debugLevel >= 2) { + GPUInfo("Using Metal device %s with properties:", device_name.c_str()); + GPUInfo("\tUnified Memory Architecture = %ld ", mInternals->device.hasUnifiedMemory); + GPUInfo("\tRecommended Max Working Set = %ld bytes", deviceGlobalMem); + GPUInfo("\tMax thread group memory = %ld bytes", deviceLocalMem); + GPUInfo("\tmaxWorkGroup = (%ld, %ld, %ld)", deviceMaxWorkGroup.width, deviceMaxWorkGroup.height, deviceMaxWorkGroup.depth); + GPUInfo(" "); + } + + mDeviceName = device_name.c_str(); + // Basically a random number for now. + mMaxBackendThreads = 1000; + + if (GetMetalPrograms()) { + return 1; + } + + if (GetProcessingSettings().debugLevel >= 2) { + GPUInfo("Metal program and kernels loaded successfully"); + } + + // We only support internal GPUs, so the ownership of the memory is always shared + mInternals->mem_gpu = [mInternals->device newBufferWithLength:mDeviceMemorySize options:MTLResourceStorageModeShared]; + if (mInternals->mem_gpu == nil) { + GPUErrorReturn("Metal Memory Allocation Error"); + } + + // We only support internal GPUs, so the ownership of the memory is always shared. + // FIXME: until I understand how to enable the gGPUConstantMemBufferSize constexpr + int32_t tmpGPUContantMemBufferSize = 100000; // gGPUConstantMemBufferSize + mInternals->mem_constant = [mInternals->device newBufferWithLength:tmpGPUContantMemBufferSize options:MTLResourceStorageModeShared]; + if (mInternals->mem_constant) { + GPUErrorReturn("Metal Constant Memory Allocation Error"); + } + + for (int32_t i = 0; i < mNStreams; i++) { + mInternals->commandQueues[i] = [mInternals->device newCommandQueue]; + if (mInternals->commandQueues[i] == nil) { + GPUErrorReturn("Error creating Metal command queue"); + } + mInternals->commandBuffers[i] = [mInternals->commandQueues[i] commandBuffer]; + if (mInternals->commandBuffers[i] == nil) { + GPUErrorReturn("Error creating Metal command buffer"); + } + } + + mInternals->mem_host = [mInternals->device newBufferWithLength:mHostMemorySize options:MTLResourceStorageModeShared]; + if (mInternals->mem_host == nil) { + GPUErrorReturn("Error allocating pinned host memory"); + } + + mHostMemoryBase = mInternals->mem_host.contents; + mHostMemorySize = mInternals->mem_host.allocatedSize; + mDeviceMemoryBase = mInternals->mem_gpu.contents; + mDeviceMemorySize = mInternals->mem_gpu.allocatedSize; + mDeviceConstantMem = (GPUConstantMem*)mInternals->mem_constant.contents; + + if (GetProcessingSettings().debugLevel >= 1) { + GPUInfo("Memory ptrs: GPU (%ld bytes): %p - Host (%ld bytes): %p", (int64_t)mDeviceMemorySize, mDeviceMemoryBase, (int64_t)mHostMemorySize, mHostMemoryBase); + memset(mHostMemoryBase, 0xDD, mHostMemorySize); + } + + GPUInfo("Metal Initialisation successfull"); + } else { + auto* master = dynamic_cast(mMaster); + mWarpSize = master->mWarpSize; + mMaxBackendThreads = master->mMaxBackendThreads; + mDeviceName = master->mDeviceName; + mDeviceConstantMem = master->mDeviceConstantMem; + mInternals = master->mInternals; + } + + for (uint32_t i = 0; i < mEvents.size(); i++) { + auto* events = (id*)mEvents[i].data(); + new (events) id[ mEvents[i].size() ]; + } + + return (0); +} + +int32_t GPUReconstructionMetal::ExitDevice_Runtime() +{ + // Uninitialize OPENCL + SynchronizeGPU(); + + if (mMaster == nullptr) { + if (mDeviceMemoryBase) { + [mInternals->mem_gpu release]; + [mInternals->mem_constant release]; + for (uint32_t i = 0; i < mInternals->functions.size(); i++) { + [mInternals->functions[i] release]; + } + mInternals->functions.clear(); + } + if (mHostMemoryBase) { + for (int32_t i = 0; i < mNStreams; i++) { + [mInternals->commandQueues[i] release]; + [mInternals->commandBuffers[i] release]; + } + [mInternals->mem_host release]; + } + + [mInternals->library release]; + [mInternals->device release]; + GPUInfo("Metal disposed correctly"); + } + mDeviceMemoryBase = nullptr; + mHostMemoryBase = nullptr; + + return (0); +} + +size_t GPUReconstructionMetal::GPUMemCpy(void* dst, const void* src, size_t sizeBytes, int32_t stream, int32_t toGPU, deviceEvent* ev, deviceEvent* evList, int32_t nEvents) +{ + if (evList == nullptr) { + nEvents = 0; + } + if (GetProcessingSettings().debugLevel >= 3) { + stream = -1; + } + + if (stream == -1) { + SynchronizeGPU(); + } + + auto realStream = stream == -1 ? 0 : stream; + id cb = mInternals->commandBuffers[realStream]; + id blit = [cb blitCommandEncoder]; + id sourceBuffer = nil; + id destBuffer = nil; + ptrdiff_t sourceOffset = 0; + ptrdiff_t destOffset = 0; + + // Sigh. + if (src > mHostMemoryBase && src < ((char*)mHostMemoryBase + mHostMemorySize)) { + sourceBuffer = mInternals->mem_host; + sourceOffset = (char*)src - (char*)mHostMemoryBase; + } else if (src > mDeviceMemoryBase && src < ((char*)mDeviceMemoryBase + mDeviceMemorySize)) { + sourceBuffer = mInternals->mem_gpu; + sourceOffset = (char*)src - (char*)mDeviceMemoryBase; + } else { + GPUErrorReturn("Unknown buffer at %x", src); + } + + if (dst > mHostMemoryBase && dst < ((char*)mHostMemoryBase + mHostMemorySize)) { + destBuffer = mInternals->mem_host; + destOffset = (char*)src - (char*)mHostMemoryBase; + } else if (dst > mDeviceMemoryBase && dst < ((char*)mDeviceMemoryBase + mDeviceMemorySize)) { + destBuffer = mInternals->mem_gpu; + destOffset = (char*)dst - (char*)mDeviceMemoryBase; + } else { + GPUErrorReturn("Unknown buffer at %x", src); + } + + [blit copyFromBuffer:sourceBuffer + sourceOffset:sourceOffset + toBuffer:destBuffer + destinationOffset:destOffset + size:sizeBytes]; + + [blit endEncoding]; + [cb commit]; + + if (GetProcessingSettings().serializeGPU & 2) { + GPUDebug(("GPUMemCpy " + std::to_string(toGPU)).c_str(), stream, true); + } + return sizeBytes; +} + +size_t GPUReconstructionMetal::WriteToConstantMemory(size_t offset, const void* src, size_t size, int32_t stream, deviceEvent* ev) +{ + if (stream == -1) { + SynchronizeGPU(); + } + + auto realStream = stream == -1 ? 0 : stream; + id cb = mInternals->commandBuffers[realStream]; + id blit = [cb blitCommandEncoder]; + id sourceBuffer = nil; + ptrdiff_t sourceOffset = 0; + if (src > mHostMemoryBase && src < ((char*)mHostMemoryBase + mHostMemorySize)) { + sourceBuffer = mInternals->mem_host; + sourceOffset = (char*)src - (char*)mHostMemoryBase; + } else if (src > mDeviceMemoryBase && src < ((char*)mDeviceMemoryBase + mDeviceMemorySize)) { + sourceBuffer = mInternals->mem_gpu; + sourceOffset = (char*)src - (char*)mDeviceMemoryBase; + } else { + GPUErrorReturn("Unknown buffer at %x", src); + } + [blit copyFromBuffer:sourceBuffer + sourceOffset:sourceOffset + toBuffer:mInternals->mem_constant + destinationOffset:offset + size:size]; + + [blit endEncoding]; + [cb commit]; + + if (GetProcessingSettings().serializeGPU & 2) { + GPUDebug("WriteToConstantMemory", stream, true); + } + return size; +} + +void GPUReconstructionMetal::ReleaseEvent(deviceEvent ev) +{ + // FIXME: is this supposed to reset the event for it to be repurposed + // or to decrease the ref count? + auto mtlEvent = (__bridge id)(ev.get()); + [mtlEvent setSignaledValue:0]; +} + +void GPUReconstructionMetal::RecordMarker(deviceEvent* ev, int32_t stream) +{ + id cb = mInternals->commandBuffers[stream]; + // Does not change the retain count, so it's important we manage + // the lifetime of the events outside here. + auto mtlEvent = (__bridge id)(ev->get()); + [cb encodeSignalEvent:mtlEvent value:1]; + [cb commit]; +} + +void GPUReconstructionMetal::SynchronizeGPU() +{ + for (int32_t i = 0; i < mNStreams; i++) { + [mInternals->commandBuffers[i] waitUntilCompleted]; + } +} + +void GPUReconstructionMetal::SynchronizeStream(int32_t stream) +{ + [mInternals->commandBuffers[stream] waitUntilCompleted]; +} + +void GPUReconstructionMetal::SynchronizeEvents(deviceEvent* evList, int32_t nEvents) +{ + // I wait for everything to complete for now... + for (int32_t si = 0; si < mNStreams; si++) { + id cb = mInternals->commandBuffers[si]; + [cb waitUntilCompleted]; + } +} + +void GPUReconstructionMetal::StreamWaitForEvents(int32_t stream, deviceEvent* evList, int32_t nEvents) +{ + // Encode commands to wait for all the events + id cb = mInternals->commandBuffers[stream]; + for (int32_t ei = 0; ei < nEvents; ei++) { + auto mtlEvent = (__bridge id)(evList[ei].get()); + [cb encodeWaitForEvent:mtlEvent value:1]; + } + [cb commit]; + [cb waitUntilCompleted]; +} + +bool GPUReconstructionMetal::IsEventDone(deviceEvent* evList, int32_t nEvents) +{ + for (int32_t i = 0; i < nEvents; i++) { + auto mtlEvent = (__bridge id)(evList[i].get()); + if (mtlEvent.signaledValue == 0) { + return false; + } + } + return true; +} + +int32_t GPUReconstructionMetal::GPUDebug(const char* state, int32_t stream, bool force) +{ + // Wait for Metal-Kernel to finish and check for Metal errors afterwards, in case of debugmode + if (!force && GetProcessingSettings().debugLevel <= 0) { + return (0); + } + for (int32_t si = 0; si < mNStreams; si++) { + [mInternals->commandBuffers[si] waitUntilCompleted]; + } + if (GetProcessingSettings().debugLevel >= 3) { + GPUInfo("GPU Sync Done"); + } + return (0); +} + +int32_t GPUReconstructionMetal::GPUChkErrInternal(const int64_t error, const char* file, int32_t line) const +{ + // Not sure how metal returns errors. + if (error != 0) { + GPUError("Metal Error: %ld / %s (%s:%d)", error, "Unknown", file, line); + } + return error != 0; +} + +// Return pointer+size for (__DATA|__DATA_CONST, "__gpu_resource") from the image +// that matches `image_name_substr` (e.g. "libO2GPUReconstruction.dylib"). +static const uint8_t* find_gpu_resource_in_image(const char* image_name_substr, + unsigned long* out_size) +{ + uint32_t count = _dyld_image_count(); + for (uint32_t i = 0; i < count; ++i) { + const char* name = _dyld_get_image_name(i); + if (!name || !strstr(name, image_name_substr)) { + continue; + } + + const struct mach_header* mh = _dyld_get_image_header(i); + + const auto* mh64 = (const struct mach_header_64*)mh; + const auto* p = (const uint8_t*) + getsectiondata(mh64, "__DATA", "__gpu_resource", out_size); + if (!p) { + p = (const uint8_t*) + getsectiondata(mh64, "__DATA_CONST", "__gpu_resource", out_size); + } + if (p) { + return p; + } + } + return nullptr; +} + +int32_t GPUReconstructionMetal::GetMetalPrograms() +{ + GPUInfo("Loading Metal library (Platform version %s)", [mInternals->device.architecture.name cStringUsingEncoding:NSUTF8StringEncoding]); + + unsigned long sz = 0; + const uint8_t* p = find_gpu_resource_in_image("libO2GPUTrackingMETAL.dylib", &sz); + if (p == nullptr || sz == 0) { + GPUError("Metal library not found in the __gpu_resource section"); + return 1; + } + + // the section is part of the mapped image, so it outlives the dispatch_data_t + // and does not have to be copied + dispatch_data_t blob = dispatch_data_create(p, sz, nullptr, ^{}); + + NSError* error = nil; + mInternals->library = [mInternals->device newLibraryWithData:blob error:&error]; + + if (mInternals->library == nil) { + NSLog(@"%@", error); + GPUError("Error loading the Metal library"); + return 1; + } + + return AddKernels(); +} diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h new file mode 100644 index 0000000000000..e8c62623cef5a --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h @@ -0,0 +1,61 @@ +// Copyright 2019-2025 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +#ifndef GPURECONSTRUCTIONOMETALINCLUDESHOST_H +#define GPURECONSTRUCTIONOMETALINCLUDESHOST_H + +#define GPUCA_GPUTYPE_METAL + +#import +#ifndef __METAL_VERSION__ +// __METAL_VERSION__ is only defined in device code. +#define __METAL_HOST__ +#endif + +#import + +#include +#include +#include +#include "GPULogging.h" + +#include "GPUReconstructionMetal.h" +#include "GPUReconstructionIncludes.h" +#include "GPUCommonHelpers.h" + +using namespace o2::gpu; + +#include +#include +#include +#include + +namespace o2::gpu +{ + +struct GPUReconstructionMetalInternals { + id device; + + std::array, constants::GPU_MAX_STREAMS> commandQueues; // ~ cl_command_queue[] + std::array, constants::GPU_MAX_STREAMS> commandBuffers; + + std::vector> functions; // ~ cl_kernel (symbols) + std::vector> pipelines; // compiled kernels + + id mem_gpu; // ~ cl_mem (device/global) + id mem_constant; // ~ cl_mem (constant-like) + id mem_host; // ~ cl_mem (host-visible) + + id library; // ~ cl_program +}; +} // namespace o2::gpu + +#endif // GPURECONSTRUCTIONOMETALINCLUDESHOST_H diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm new file mode 100644 index 0000000000000..54adb3076b60d --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm @@ -0,0 +1,82 @@ +// Copyright 2019-2025 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +#include +#include "GPUReconstructionKernelIncludes.h" +#include "GPUReconstructionMetalIncludesHost.h" + +#include "GPUReconstructionMetalKernelsSpecialize.inc" +#include "GPUReconstructionProcessingKernels.inc" + +template void GPUReconstructionProcessing::KernelInterface::runKernelVirtual(const int num, const void* args); + +template +inline void GPUReconstructionMetal::runKernelBackend(const krnlSetupTime& _xyz, const Args&... args) +{ + id function = mInternals->functions[GetKernelNum()]; + auto& kExec = _xyz.x; + auto& runRange = _xyz.y; + // auto& events = _xyz.z; + // auto& t = _xyz.t; + + NSError* error = nil; + auto pso = [mInternals->device newComputePipelineStateWithFunction:function error:&error]; + id computeEncoder = [mInternals->commandBuffers[kExec.stream] computeCommandEncoder]; + + // Map buffers and states + [computeEncoder setComputePipelineState:pso]; + [computeEncoder setBuffer:mInternals->mem_gpu offset:0 atIndex:0]; + [computeEncoder setBuffer:mInternals->mem_constant offset:0 atIndex:1]; + [computeEncoder setBuffer:mInternals->mem_host offset:0 atIndex:2]; + + MTLSize gridSize = MTLSizeMake(runRange.index, 1, 1); + + NSUInteger threadGroupSize = pso.maxTotalThreadsPerThreadgroup; + if (threadGroupSize > runRange.index) { + threadGroupSize = runRange.index; + } + + MTLSize threadgroupSize = MTLSizeMake(threadGroupSize, 1, 1); + [computeEncoder dispatchThreads:gridSize + threadsPerThreadgroup:threadgroupSize]; +} + +template +int32_t GPUReconstructionMetal::AddKernel() +{ + NSString* kname = [[NSString alloc] initWithFormat:@"krnl_%s", GetKernelName()]; + + id krnl = [mInternals->library newFunctionWithName:kname]; + if (krnl == nil) { + GPUError("Error creating Metal Kernel: %s", [kname cStringUsingEncoding:NSUTF8StringEncoding]); + return 1; + } + + mInternals->functions.emplace_back(krnl); + return 0; +} + +template +S& GPUReconstructionMetal::getKernelObject() +{ + return mInternals->functions[GetKernelNum()]; +} + +int32_t GPUReconstructionMetal::AddKernels() +{ +#define GPUCA_KRNL(x_class, ...) \ + if (AddKernel()) { \ + return 1; \ + } +#include "GPUReconstructionKernelList.h" +#undef GPUCA_KRNL + return 0; +} diff --git a/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc new file mode 100644 index 0000000000000..1ee192cff822f --- /dev/null +++ b/GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernelsSpecialize.inc @@ -0,0 +1,27 @@ +// Copyright 2019-2025 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +/// \file GPUReconstructionMetalKernelsSpecialize.inc + +template <> +inline void GPUReconstructionMetal::runKernelBackend(const krnlSetupTime& _xyz, void* const& ptr, uint64_t const& size) +{ + const uint64_t offset = static_cast(ptr) - static_cast(mDeviceMemoryBase); + const uint64_t length = (size + 15ull) & ~15ull; + + id cb = mInternals->commandBuffers[_xyz.x.stream]; + id blit = [cb blitCommandEncoder]; + + [blit fillBuffer:mInternals->mem_gpu range:NSMakeRange(offset, length) value:0]; + [blit endEncoding]; + + [cb commit]; +} diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index ca58d91212084..cc44b2003a01d 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -463,41 +463,39 @@ if (onnxruntime_FOUND) endif() # Add CMake recipes for GPU Tracking librararies -if(CUDA_ENABLED OR OPENCL_ENABLED OR HIP_ENABLED) - if(CMAKE_SYSTEM_NAME MATCHES Darwin) - message(WARNING "GPU Tracking disabled on MacOS") - else() - make_directory(${CMAKE_CURRENT_BINARY_DIR}/genGPUArch) - set(GPU_CONST_PARAM_FILES) - foreach(GPU_ARCH ${GPU_CONST_PARAM_ARCHITECTUES}) - set(PARAMFILE ${CMAKE_CURRENT_BINARY_DIR}/genGPUArch/gpu_const_param_${GPU_ARCH}.par) - add_custom_command( - OUTPUT ${PARAMFILE} - COMMAND bash -c - "echo -e '#define GPUCA_GPUTYPE_${GPU_ARCH}\\n#define PARAMETER_FILE \"GPUDefParametersDefaults.h\"\\ngInterpreter->AddIncludePath(\"${CMAKE_CURRENT_SOURCE_DIR}/Definitions\");\\ngInterpreter->AddIncludePath(\"${ON_THE_FLY_DIR}\");\\n.x ${CMAKE_CURRENT_SOURCE_DIR}/Standalone/tools/dumpGPUDefParam.C(\"${PARAMFILE}\")\\n.q\\n'" - | root -l -b > /dev/null - VERBATIM - WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/genGPUArch - MAIN_DEPENDENCY Standalone/tools/dumpGPUDefParam.C - DEPENDS ${GPU_DEFAULT_PARAMS_HEADER} - ${GPU_DEFAULT_PARAMS_HEADER_DEVICE} - ${ON_THE_FLY_DIR}/GPUDefParametersLoadPrepare.h - ${ON_THE_FLY_DIR}/GPUDefParametersLoad.inc - COMMENT "Generating GPU parameter set for architecture ${GPU_ARCH}") - LIST(APPEND GPU_CONST_PARAM_FILES ${PARAMFILE}) - endforeach() - add_custom_target(${MODULE}_GPU_CONST_PARAM_ARCHS ALL DEPENDS ${GPU_CONST_PARAM_FILES}) - install(FILES ${GPU_CONST_PARAM_FILES} DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/GPU/arch_param) - - if(CUDA_ENABLED) - add_subdirectory(Base/cuda) - endif() - if(OPENCL_ENABLED) - add_subdirectory(Base/opencl) - endif() - if(HIP_ENABLED) - add_subdirectory(Base/hip) - endif() +if(METAL_ENABLED) + add_subdirectory(Base/metal) +elseif(CUDA_ENABLED OR OPENCL_ENABLED OR HIP_ENABLED) + make_directory(${CMAKE_CURRENT_BINARY_DIR}/genGPUArch) + set(GPU_CONST_PARAM_FILES) + foreach(GPU_ARCH ${GPU_CONST_PARAM_ARCHITECTUES}) + set(PARAMFILE ${CMAKE_CURRENT_BINARY_DIR}/genGPUArch/gpu_const_param_${GPU_ARCH}.par) + add_custom_command( + OUTPUT ${PARAMFILE} + COMMAND bash -c + "echo -e '#define GPUCA_GPUTYPE_${GPU_ARCH}\\n#define PARAMETER_FILE \"GPUDefParametersDefaults.h\"\\ngInterpreter->AddIncludePath(\"${CMAKE_CURRENT_SOURCE_DIR}/Definitions\");\\ngInterpreter->AddIncludePath(\"${ON_THE_FLY_DIR}\");\\n.x ${CMAKE_CURRENT_SOURCE_DIR}/Standalone/tools/dumpGPUDefParam.C(\"${PARAMFILE}\")\\n.q\\n'" + | root -l -b > /dev/null + VERBATIM + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/genGPUArch + MAIN_DEPENDENCY Standalone/tools/dumpGPUDefParam.C + DEPENDS ${GPU_DEFAULT_PARAMS_HEADER} + ${GPU_DEFAULT_PARAMS_HEADER_DEVICE} + ${ON_THE_FLY_DIR}/GPUDefParametersLoadPrepare.h + ${ON_THE_FLY_DIR}/GPUDefParametersLoad.inc + COMMENT "Generating GPU parameter set for architecture ${GPU_ARCH}") + LIST(APPEND GPU_CONST_PARAM_FILES ${PARAMFILE}) + endforeach() + add_custom_target(${MODULE}_GPU_CONST_PARAM_ARCHS ALL DEPENDS ${GPU_CONST_PARAM_FILES}) + install(FILES ${GPU_CONST_PARAM_FILES} DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/GPU/arch_param) + + if(CUDA_ENABLED) + add_subdirectory(Base/cuda) + endif() + if(OPENCL_ENABLED) + add_subdirectory(Base/opencl) + endif() + if(HIP_ENABLED) + add_subdirectory(Base/hip) endif() endif() diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx index b499ea10e679b..7e74d209cbd68 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx +++ b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx @@ -33,7 +33,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); #else // GPUCA_DETERMINISTIC_MODE if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZPadTime) { - CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZTimePad) { - CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortPad) { - CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortTime) { - CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); + CAAlgo::sortInBlock(nThreads, iThread, sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); } #endif // GPUCA_DETERMINISTIC_MODE GPUbarrier(); } - for (uint32_t j = get_local_id(0); j < count; j += get_local_size(0)) { + for (uint32_t j = iThread; j < count; j += nThreads) { int32_t outidx = idOffsetOut + totalCount + j; const ClusterNative& GPUrestrict() orgCl = clusters -> clusters[iSector][iRow][sortBuffer[j]]; diff --git a/GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx index 0d2140c32e4a9..f245e3b53d2bd 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx +++ b/GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx @@ -31,7 +31,7 @@ GPUdii() void GPUTPCDecompressionKernels::ThreadnClusters[sector][row]; k++) { @@ -125,7 +125,7 @@ GPUdii() void GPUTPCDecompressionUtilKernels::ThreadclusterOffset[sector][row]; diff --git a/GPU/GPUTracking/Definitions/GPUDef.h b/GPU/GPUTracking/Definitions/GPUDef.h index 692e0c5ebe231..ece7cc57f7096 100644 --- a/GPU/GPUTracking/Definitions/GPUDef.h +++ b/GPU/GPUTracking/Definitions/GPUDef.h @@ -21,17 +21,27 @@ #include "GPUDefParametersWrapper.h" #include "GPUCommonRtypes.h" -// Macros for masking ptrs in OpenCL kernel calls as uint64_t (The API only allows us to pass buffer objects) +// Macros for kernel arguments. OpenCL can only pass buffer objects, so pointers +// are masked as uint64_t and cast back inside the kernel. MSL needs an explicit +// buffer index on every parameter, but can bind a pointer directly. The index is +// emitted per argument by o2_gpu_add_kernel; 0, 1 and 2 are taken by gpu_mem, +// the constant memory and the sector index. #ifdef __OPENCL__ - #define GPUPtr1(a, b) uint64_t b - #ifdef __OPENCL__ - #define GPUPtr2(a, b) ((__generic a) (a) b) - #else - #define GPUPtr2(a, b) ((__global a) (a) b) - #endif + #define GPUPtr1(idx, a, b) uint64_t b + #define GPUPtr2(a, b) ((__generic a) (a) b) + #define GPUArg1(idx, a, b) a b +#elif defined(__METAL__) + // As for OpenCL, pointers travel as a 64-bit address: a pointer to a derived + // class is not a valid kernel argument type in MSL either. + #define GPUPtr1(idx, a, b) constant uint64_t& b [[buffer(idx)]] + // through device and then to generic: the kernel's own buffers are device + // memory, but the Thread() entry points take the pointer unannotated + #define GPUPtr2(a, b) ((a)((device a)(b))) + #define GPUArg1(idx, a, b) constant a& b [[buffer(idx)]] #else - #define GPUPtr1(a, b) a b + #define GPUPtr1(idx, a, b) a b #define GPUPtr2(a, b) b + #define GPUArg1(idx, a, b) a b #endif #define GPUCA_EVDUMP_FILE "event" @@ -39,19 +49,19 @@ #ifdef GPUCA_GPUCODE #define GPUCA_MAKE_SHARED_REF(vartype, varname, varglobal, varshared) const GPUsharedref() vartype& __restrict__ varname = varshared; #define GPUCA_SHARED_STORAGE(storage) storage - #define GPUCA_SHARED_CACHE(target, src, size) \ + #define GPUCA_SHARED_CACHE(nThreads, iThread, target, src, size) \ static_assert((size) % sizeof(int32_t) == 0, "Invalid shared cache size"); \ - for (uint32_t i_shared_cache = get_local_id(0); i_shared_cache < (size) / sizeof(int32_t); i_shared_cache += get_local_size(0)) { \ + for (uint32_t i_shared_cache = (iThread); i_shared_cache < (size) / sizeof(int32_t); i_shared_cache += (nThreads)) { \ reinterpret_cast(target)[i_shared_cache] = reinterpret_cast(src)[i_shared_cache]; \ } - #define GPUCA_SHARED_CACHE_REF(target, src, size, reftype, ref) \ - GPUCA_SHARED_CACHE(target, src, size) \ + #define GPUCA_SHARED_CACHE_REF(nThreads, iThread, target, src, size, reftype, ref) \ + GPUCA_SHARED_CACHE(nThreads, iThread, target, src, size) \ GPUsharedref() const reftype* __restrict__ ref = (target) #else #define GPUCA_MAKE_SHARED_REF(vartype, varname, varglobal, varshared) const GPUglobalref() vartype & __restrict__ varname = varglobal; #define GPUCA_SHARED_STORAGE(storage) - #define GPUCA_SHARED_CACHE(target, src, size) - #define GPUCA_SHARED_CACHE_REF(target, src, size, reftype, ref) GPUglobalref() const reftype* __restrict__ ref = src + #define GPUCA_SHARED_CACHE(nThreads, iThread, target, src, size) + #define GPUCA_SHARED_CACHE_REF(nThreads, iThread, target, src, size, reftype, ref) GPUglobalref() const reftype* __restrict__ ref = src #endif #endif //GPUTPCDEF_H diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index 4e3e0cd32d752..81949f0b8ed2a 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -221,7 +221,7 @@ AddSubConfig(GPUSettingsRecDynamic, dyn) AddHelp("help", 'h') EndConfig() -#ifndef __OPENCL__ +#if !defined(__OPENCL__) && !defined(__METAL__) // these use std::string / std::vector, which device code does not have // Parameters that might affect the RTC code (if these change, the cache cannot be used) BeginSubConfig(GPUSettingsProcessingRTC, rtc, configStandalone.proc, "RTC", 0, "Processing settings", proc_rtc) AddOption(cacheOutput, bool, false, "", 0, "Cache RTC compilation results") @@ -428,7 +428,7 @@ AddSubConfig(GPUSettingsProcessingNNclusterizer, nn) AddSubConfig(GPUSettingsProcessingScaling, scaling) AddHelp("help", 'h') EndConfig() -#endif // __OPENCL__ +#endif // !__OPENCL__ && !__METAL__ #ifndef GPUCA_GPUCODE_DEVICE // Light settings concerning the event display (can be changed without rebuilding vertices) diff --git a/GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv b/GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv index 12c87afceb86f..f846641cefc9e 100644 --- a/GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv +++ b/GPU/GPUTracking/Definitions/Parameters/GPUParameters.csv @@ -1,117 +1,117 @@ -Architecture,default,default_cpu,MI100,VEGA,TAHITI,TESLA,FERMI,PASCAL,KEPLER,AMPERE,TURING,HOPPER,ADA,OPENCL,RDNA,MI210,BLACKWELL,MI300 -,,,,,,,,,,,,,,,,,, -CORE:,,,,,,,,,,,,,,,,,, -WARP_SIZE,0,,64,64,32,32,32,32,32,32,32,32,32,32,32,64,32,64 -THREAD_COUNT_DEFAULT,256,,256,256,,,,,,512,512,,512,256,512,512,512, -,,,,,,,,,,,,,,,,,, -LB:,,,,,,,,,,,,,,,,,, -GPUTPCCreateTrackingData,256,,"[256, 7]","[192, 2]",,,,,,"[224, 7]",256,"[128, 14]",416,,"[64, 21]",,384,"[320, 2]" -GPUTPCTrackletConstructor,256,,"[768, 8]","[512, 10]","[256, 2]","[256, 1]","[256, 2]","[1024, 2]","[512, 4]",1024,"[256, 2]",1024,"[1024, 1]",,"[768, 2]",,768,512 -GPUTPCTrackletSelector,256,,"[384, 5]","[192, 10]","[256, 3]","[256, 1]","[256, 3]","[512, 4]","[256, 3]","[288, 3]","[192, 3]","[544, 1]","[32, 2]",,"[384, 3]",,992,"[256, 6]" -GPUTPCNeighboursFinder,256,,"[192, 8]","[960, 8]",256,256,256,512,256,864,"[640, 1]","[512, 2]","[736, 1]",,"[480, 3]",,992,"[704, 1]" -GPUTPCNeighboursCleaner,256,,"[128, 5]","[384, 9]",256,256,256,256,256,544,512,"[192, 9]","[512, 1]",,"[384, 5]",,672,"[640, 1]" -GPUTPCExtrapolationTracking,256,,"[256, 7]","[256, 2]",,,,,,"[352, 4]","[192, 2]","[896, 1]","[352, 1]",,"[1024, 1]",,896,1024 -GPUTRDTrackerKernels_gpuVersion,512,,,,,,,,,512,,512,512,,512,,,512 -GPUTPCCreateOccupancyMap_fill,256,,,,,,,,,256,,256,256,,256,,,256 -GPUTPCCreateOccupancyMap_fold,256,,,,,,,,,256,,256,256,,256,,,256 -GPUTRDTrackerKernels_o2Version,512,,,,,,,,,512,,512,512,,512,,,512 -GPUTPCCompressionKernels_step0attached,256,,"[128, 1]","[64, 2]",,,,,,"[160, 2]",128,"[448, 1]",352,,"[1024, 1]",,"[96, 3]","[128, 4]" -GPUTPCCompressionKernels_step1unattached,256,,"[512, 2]","[512, 2]",,,,,,"[288, 4]","[512, 2]","[256, 4]","[512, 2]",,"[512, 3]",,"[512, 2]","[512, 3]" -GPUTPCDecompressionKernels_step0attached,256,,"[128, 2]","[128, 2]",,,,,,"[32, 1]","[32, 1]","[32, 1]","[32, 1]",,"[128, 1]",,"[32, 1]","[128, 1]" -GPUTPCDecompressionKernels_step1unattached,256,,"[64, 2]","[64, 2]",,,,,,"[32, 1]","[32, 1]","[32, 1]","[32, 1]",,"[64, 1]",,"[32, 1]","[64, 1]" -GPUTPCDecompressionUtilKernels_sortPerSectorRow,256,,,,,,,,,256,,256,256,,256,,,256 -GPUTPCDecompressionUtilKernels_countFilteredClusters,256,,,,,,,,,256,,256,256,,256,,,256 -GPUTPCDecompressionUtilKernels_storeFilteredClusters,256,,,,,,,,,256,,256,256,,256,,,256 -GPUTPCCFDecodeZS,"[128, 4]",,"[64, 4]","[64, 1]",,,,,,"[32, 10]","[64, 8]","[32, 10]","[32, 10]",,"[64, 1]",,"[64, 10]","[64, 1]" -GPUTPCCFDecodeZSLink,"""GPUCA_WARP_SIZE""",,"""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""",,,,,,"""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""",,64,,"""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""" -GPUTPCCFDecodeZSDenseLink,"""GPUCA_WARP_SIZE""",,"[""GPUCA_WARP_SIZE"", 4]","[""GPUCA_WARP_SIZE"", 14]",,,,,,"[""GPUCA_WARP_SIZE"", 14]","""GPUCA_WARP_SIZE""","[""GPUCA_WARP_SIZE"", 22]","[""GPUCA_WARP_SIZE"", 22]",,"[64, 17]",,"[""GPUCA_WARP_SIZE"", 8]","[""GPUCA_WARP_SIZE"", 5]" -GPUTPCCFGather,"[1024, 1]",,"[1024, 5]","[1024, 1]",,,,,,"[160, 11]","[1024, 1]",736,896,,"[928, 1]",,"[1024, 1]","[320, 2]" -COMPRESSION_GATHER,1024,,1024,1024,,,,,,1024,1024,,1024,,,,, -GPUTPCGMMergerTrackFit,256,,"[192, 2]","[64, 7]",,,,,,"[32, 16]","[32, 8]","[32, 14]","[160, 2]",,"[32, 24]",,"[64, 8]","[64, 6]" -GPUTPCGMMergerFollowLoopers,256,,"[256, 5]","[256, 4]",,,,,,"[256, 4]","[128, 4]","[1024, 1]",640,,"[128, 16]",,"[224, 3]","[256, 7]" -GPUTPCGMMergerSectorRefit,256,,"[64, 4]","[256, 2]",,,,,,"[32, 8]","[64, 5]","[32, 7]","[32, 7]",,"[32, 20]",,"[32, 10]","[64, 4]" -GPUTPCGMMergerUnpackResetIds,256,,256,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerUnpackGlobal,256,,256,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerResolve_step0,256,,512,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerResolve_step1,256,,512,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerResolve_step2,256,,512,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerResolve_step3,256,,512,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerResolve_step4,256,,512,256,,,,,,"[256, 4]","[256, 4]","[256, 4]","[256, 4]",,256,,"[256, 4]",256 -GPUTPCGMMergerClearLinks,256,,256,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerMergeWithinPrepare,256,,256,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerMergeSectorsPrepare,256,,256,256,,,,,,"[256, 2]","[256, 2]","[256, 2]","[256, 2]",,256,,"[256, 2]",256 -GPUTPCGMMergerMergeBorders_step0,256,,512,256,,,,,,192,192,192,192,,256,,192,256 -GPUTPCGMMergerMergeBorders_step2,256,,512,256,,,,,,"[64, 2]",256,"[64, 2]","[64, 2]",,256,,"[64, 2]",256 -GPUTPCGMMergerMergeCE,256,,512,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerLinkExtrapolatedTracks,256,,256,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerCollect,256,,"[768, 1]","[1024, 1]",,,,,,"[864, 1]","[128, 2]","[896, 1]",128,,1024,,"[288, 1]","[384, 4]" -GPUTPCGMMergerSortTracksPrepare,256,,256,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerPrepareForFit_step0,256,,256,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerPrepareForFit_step1,256,,256,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerPrepareForFit_step2,256,,256,256,,,,,,256,256,256,256,,256,,256,256 -GPUTPCGMMergerFinalize_step0,256,,,256,,,,,,256,,256,256,,256,,256,256 -GPUTPCGMMergerFinalize_step1,256,,,256,,,,,,256,,256,256,,256,,256,256 -GPUTPCGMMergerFinalize_step2,256,,,256,,,,,,256,,256,256,,256,,256,256 -GPUTPCGMMergerMergeLoopers_step0,256,,,,,,,,,256,,256,256,,256,,256,256 -GPUTPCGMMergerMergeLoopers_step1,256,,,,,,,,,256,,256,256,,256,,256,256 -GPUTPCGMMergerMergeLoopers_step2,256,,,,,,,,,256,,256,256,,256,,256,256 -GPUTPCGMO2Output_prepare,256,,,,,,,,,256,,256,256,,256,,256,256 -GPUTPCGMO2Output_output,256,,,,,,,,,256,,256,256,,256,,256,256 -GPUTPCStartHitsFinder,256,,"[1024, 2]","[1024, 7]",256,256,256,256,256,"[224, 1]",512,"[416, 4]",928,,"[320, 5]",,608,"[448, 3]" -GPUTPCStartHitsSorter,256,,"[1024, 5]","[512, 7]",256,256,256,256,256,"[320, 2]","[512, 1]","[864, 1]","[96, 2]",,"[192, 5]",,608,"[448, 1]" -GPUTPCCFCheckPadBaseline,576,,"[576, 2]","[576, 2]",,,,,,"[576, 3]",,"[576, 1]","[576, 1]",,"[576, 2]",,"[576, 2]",576 -GPUTPCCFHIPTailConnector,256,,256,256,,,,,,"[224, 2]",,"[320, 5]","[704, 1]",,"[128, 7]",,,"[448, 4]" -GPUTPCCFHIPClusterizer,256,,256,256,,,,,,"[288, 5]",,"[480, 3]","[448, 3]",,352,,,"[512, 3]" -GPUTPCCFChargeMapFiller_fillIndexMap,512,,512,512,,,,,,448,,448,448,,512,,448,512 -GPUTPCCFChargeMapFiller_fillFromDigits,512,,512,512,,,,,,448,,448,448,,512,,448,512 -GPUTPCCFChargeMapFiller_findFragmentStart,512,,512,512,,,,,,448,,448,448,,512,,448,512 -GPUTPCCFPeakFinder,512,,"[512, 9]","[512, 4]",,,,,,416,,992,"[672, 1]",,"[384, 2]",,"[128, 5]","[192, 10]" -GPUTPCCFNoiseSuppression,512,,512,512,,,,,,608,,896,480,,160,,,448 -GPUTPCCFDeconvolution,512,,"[512, 5]","[512, 5]",,,,,,"[480, 4]",,224,512,,480,,384,"[448, 3]" -GPUTPCCFClusterizer,512,,"[448, 3]","[512, 2]",,,,,,"[608, 3]",,736,"[192, 3]",,576,,"[160, 5]","[832, 2]" -GPUTPCNNClusterizerKernels,512,,,,,,,,,,,,,,,,, -GPUTrackingRefitKernel_mode0asGPU,256,,,,,,,,,256,,256,256,,256,,256,256 -GPUTrackingRefitKernel_mode1asTrackParCov,256,,,,,,,,,256,,256,256,,256,,256,256 -GPUMemClean16,"[""GPUCA_THREAD_COUNT_DEFAULT"", 1]",,,,,,,,,"[512, 1]",,"[512, 1]","[512, 1]",,"[256, 1]",,,"[256, 1]" -GPUitoa,"[""GPUCA_THREAD_COUNT_DEFAULT"", 1]",,,,,,,,,"[512, 1]",,"[512, 1]","[512, 1]",,"[256, 1]",,,"[256, 1]" -GPUTPCCFNoiseSuppression_noiseSuppression,"""GPUCA_LB_GPUTPCCFNoiseSuppression""",,,,,,,,,,,,,,,,448, -GPUTPCCFNoiseSuppression_updatePeaks,"""GPUCA_LB_GPUTPCCFNoiseSuppression""",,,,,,,,,,,,,,,,448, -GPUTPCNNClusterizerKernels_runCfClusterizer,"""GPUCA_LB_GPUTPCCFClusterizer""",,,,,,,,,,,,,,,,, -GPUTPCNNClusterizerKernels_fillInputNNCPU,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,, -GPUTPCNNClusterizerKernels_fillInputNNGPU,1024,,,,,,,,,,,,,,,,, -GPUTPCNNClusterizerKernels_determineClass1Labels,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,, -GPUTPCNNClusterizerKernels_determineClass2Labels,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,, -GPUTPCNNClusterizerKernels_publishClass1Regression,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,, -GPUTPCNNClusterizerKernels_publishClass2Regression,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,, -GPUTPCNNClusterizerKernels_publishDeconvolutionFlags,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,, -GPUTPCCFStreamCompaction_scanStart,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,, -GPUTPCCFStreamCompaction_scanUp,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,, -GPUTPCCFStreamCompaction_scanTop,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,, -GPUTPCCFStreamCompaction_scanDown,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,, -GPUTPCCFStreamCompaction_compactDigits,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,, -GPUTPCCompressionGatherKernels_unbuffered,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,, -GPUTPCCompressionGatherKernels_buffered32,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,, -GPUTPCCompressionGatherKernels_buffered64,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,, -GPUTPCCompressionGatherKernels_buffered128,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,, -GPUTPCCompressionGatherKernels_multiBlock,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,, -GPUTPCGMMergerFinalize_0,256,,256,,,,,,,256,256,,256,,,,256, -GPUTPCGMMergerFinalize_1,256,,256,,,,,,,256,256,,256,,,,256, -GPUTPCGMMergerFinalize_2,256,,256,,,,,,,256,256,,256,,,,256, -GPUTPCConvertKernel,,,,,,,,,,,,,,,256,,,256 -,,,,,,,,,,,,,,,,,, -PAR:,,,,,,,,,,,,,,,,,, -AMD_EUS_PER_CU,0,0,4,4,,,,,,,,,,,4,,0,4 -SORT_STARTHITS,1,0,,,,,,,,1,,1,1,,1,,1,1 -NEIGHBOURS_FINDER_MAX_NNEIGHUP,6,0,10,4,,,,,,4,4,4,4,,5,,2,5 -NEIGHBOURS_FINDER_UNROLL_GLOBAL,4,0,4,2,,,,,,2,,8,8,,4,,2,2 -NEIGHBOURS_FINDER_UNROLL_SHARED,1,0,0,0,,,,,,1,,1,0,,1,,1,1 -TRACKLET_SELECTOR_HITS_REG_SIZE,12,0,9,27,,,,,,20,20,20,20,,20,,2,20 -ALTERNATE_BORDER_SORT,1,0,1,1,,,,,,1,1,1,1,,1,,1,1 -SORT_BEFORE_FIT,1,0,1,1,,,,,,1,1,1,1,,1,,1,1 -NO_ATOMIC_PRECHECK,0,0,1,1,,,,,,1,1,1,1,,1,,1,1 -DEDX_STORAGE_TYPE,"""half""","""float""",,,,,,,,,,,,,,,, -MERGER_INTERPOLATION_ERROR_TYPE,"""half""","""float""",,,,,,,,,,,,,,,, -COMP_GATHER_KERNEL,4,0,4,4,,,,,,4,4,4,4,,4,,4,4 -COMP_GATHER_MODE,3,0,3,3,,,,,,3,3,3,3,,3,,3,3 -CF_SCAN_WORKGROUP_SIZE,512,0,,,,,,,,224,,992,448,,1024,,,448 -MERGER_SPLIT_LOOP_INTERPOLATION,,,,,,,,,,,,,,,1,,,1 +Architecture,default,default_cpu,MI100,VEGA,TAHITI,TESLA,FERMI,PASCAL,KEPLER,AMPERE,TURING,HOPPER,ADA,OPENCL,METAL,RDNA,MI210,BLACKWELL,MI300 +,,,,,,,,,,,,,,,,,,, +CORE:,,,,,,,,,,,,,,,,,,, +WARP_SIZE,0,,64,64,32,32,32,32,32,32,32,32,32,32,32,32,64,32,64 +THREAD_COUNT_DEFAULT,256,,256,256,,,,,,512,512,,512,256,256,512,512,512, +,,,,,,,,,,,,,,,,,,, +LB:,,,,,,,,,,,,,,,,,,, +GPUTPCCreateTrackingData,256,,"[256, 7]","[192, 2]",,,,,,"[224, 7]",256,"[128, 14]",416,,,"[64, 21]",,384,"[320, 2]" +GPUTPCTrackletConstructor,256,,"[768, 8]","[512, 10]","[256, 2]","[256, 1]","[256, 2]","[1024, 2]","[512, 4]",1024,"[256, 2]",1024,"[1024, 1]",,,"[768, 2]",,768,512 +GPUTPCTrackletSelector,256,,"[384, 5]","[192, 10]","[256, 3]","[256, 1]","[256, 3]","[512, 4]","[256, 3]","[288, 3]","[192, 3]","[544, 1]","[32, 2]",,,"[384, 3]",,992,"[256, 6]" +GPUTPCNeighboursFinder,256,,"[192, 8]","[960, 8]",256,256,256,512,256,864,"[640, 1]","[512, 2]","[736, 1]",,,"[480, 3]",,992,"[704, 1]" +GPUTPCNeighboursCleaner,256,,"[128, 5]","[384, 9]",256,256,256,256,256,544,512,"[192, 9]","[512, 1]",,,"[384, 5]",,672,"[640, 1]" +GPUTPCExtrapolationTracking,256,,"[256, 7]","[256, 2]",,,,,,"[352, 4]","[192, 2]","[896, 1]","[352, 1]",,,"[1024, 1]",,896,1024 +GPUTRDTrackerKernels_gpuVersion,512,,,,,,,,,512,,512,512,,,512,,,512 +GPUTPCCreateOccupancyMap_fill,256,,,,,,,,,256,,256,256,,,256,,,256 +GPUTPCCreateOccupancyMap_fold,256,,,,,,,,,256,,256,256,,,256,,,256 +GPUTRDTrackerKernels_o2Version,512,,,,,,,,,512,,512,512,,,512,,,512 +GPUTPCCompressionKernels_step0attached,256,,"[128, 1]","[64, 2]",,,,,,"[160, 2]",128,"[448, 1]",352,,,"[1024, 1]",,"[96, 3]","[128, 4]" +GPUTPCCompressionKernels_step1unattached,256,,"[512, 2]","[512, 2]",,,,,,"[288, 4]","[512, 2]","[256, 4]","[512, 2]",,,"[512, 3]",,"[512, 2]","[512, 3]" +GPUTPCDecompressionKernels_step0attached,256,,"[128, 2]","[128, 2]",,,,,,"[32, 1]","[32, 1]","[32, 1]","[32, 1]",,,"[128, 1]",,"[32, 1]","[128, 1]" +GPUTPCDecompressionKernels_step1unattached,256,,"[64, 2]","[64, 2]",,,,,,"[32, 1]","[32, 1]","[32, 1]","[32, 1]",,,"[64, 1]",,"[32, 1]","[64, 1]" +GPUTPCDecompressionUtilKernels_sortPerSectorRow,256,,,,,,,,,256,,256,256,,,256,,,256 +GPUTPCDecompressionUtilKernels_countFilteredClusters,256,,,,,,,,,256,,256,256,,,256,,,256 +GPUTPCDecompressionUtilKernels_storeFilteredClusters,256,,,,,,,,,256,,256,256,,,256,,,256 +GPUTPCCFDecodeZS,"[128, 4]",,"[64, 4]","[64, 1]",,,,,,"[32, 10]","[64, 8]","[32, 10]","[32, 10]",,,"[64, 1]",,"[64, 10]","[64, 1]" +GPUTPCCFDecodeZSLink,"""GPUCA_WARP_SIZE""",,"""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""",,,,,,"""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""",,,64,,"""GPUCA_WARP_SIZE""","""GPUCA_WARP_SIZE""" +GPUTPCCFDecodeZSDenseLink,"""GPUCA_WARP_SIZE""",,"[""GPUCA_WARP_SIZE"", 4]","[""GPUCA_WARP_SIZE"", 14]",,,,,,"[""GPUCA_WARP_SIZE"", 14]","""GPUCA_WARP_SIZE""","[""GPUCA_WARP_SIZE"", 22]","[""GPUCA_WARP_SIZE"", 22]",,,"[64, 17]",,"[""GPUCA_WARP_SIZE"", 8]","[""GPUCA_WARP_SIZE"", 5]" +GPUTPCCFGather,"[1024, 1]",,"[1024, 5]","[1024, 1]",,,,,,"[160, 11]","[1024, 1]",736,896,,,"[928, 1]",,"[1024, 1]","[320, 2]" +COMPRESSION_GATHER,1024,,1024,1024,,,,,,1024,1024,,1024,,,,,, +GPUTPCGMMergerTrackFit,256,,"[192, 2]","[64, 7]",,,,,,"[32, 16]","[32, 8]","[32, 14]","[160, 2]",,,"[32, 24]",,"[64, 8]","[64, 6]" +GPUTPCGMMergerFollowLoopers,256,,"[256, 5]","[256, 4]",,,,,,"[256, 4]","[128, 4]","[1024, 1]",640,,,"[128, 16]",,"[224, 3]","[256, 7]" +GPUTPCGMMergerSectorRefit,256,,"[64, 4]","[256, 2]",,,,,,"[32, 8]","[64, 5]","[32, 7]","[32, 7]",,,"[32, 20]",,"[32, 10]","[64, 4]" +GPUTPCGMMergerUnpackResetIds,256,,256,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerUnpackGlobal,256,,256,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerResolve_step0,256,,512,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerResolve_step1,256,,512,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerResolve_step2,256,,512,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerResolve_step3,256,,512,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerResolve_step4,256,,512,256,,,,,,"[256, 4]","[256, 4]","[256, 4]","[256, 4]",,,256,,"[256, 4]",256 +GPUTPCGMMergerClearLinks,256,,256,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerMergeWithinPrepare,256,,256,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerMergeSectorsPrepare,256,,256,256,,,,,,"[256, 2]","[256, 2]","[256, 2]","[256, 2]",,,256,,"[256, 2]",256 +GPUTPCGMMergerMergeBorders_step0,256,,512,256,,,,,,192,192,192,192,,,256,,192,256 +GPUTPCGMMergerMergeBorders_step2,256,,512,256,,,,,,"[64, 2]",256,"[64, 2]","[64, 2]",,,256,,"[64, 2]",256 +GPUTPCGMMergerMergeCE,256,,512,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerLinkExtrapolatedTracks,256,,256,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerCollect,256,,"[768, 1]","[1024, 1]",,,,,,"[864, 1]","[128, 2]","[896, 1]",128,,,1024,,"[288, 1]","[384, 4]" +GPUTPCGMMergerSortTracksPrepare,256,,256,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerPrepareForFit_step0,256,,256,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerPrepareForFit_step1,256,,256,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerPrepareForFit_step2,256,,256,256,,,,,,256,256,256,256,,,256,,256,256 +GPUTPCGMMergerFinalize_step0,256,,,256,,,,,,256,,256,256,,,256,,256,256 +GPUTPCGMMergerFinalize_step1,256,,,256,,,,,,256,,256,256,,,256,,256,256 +GPUTPCGMMergerFinalize_step2,256,,,256,,,,,,256,,256,256,,,256,,256,256 +GPUTPCGMMergerMergeLoopers_step0,256,,,,,,,,,256,,256,256,,,256,,256,256 +GPUTPCGMMergerMergeLoopers_step1,256,,,,,,,,,256,,256,256,,,256,,256,256 +GPUTPCGMMergerMergeLoopers_step2,256,,,,,,,,,256,,256,256,,,256,,256,256 +GPUTPCGMO2Output_prepare,256,,,,,,,,,256,,256,256,,,256,,256,256 +GPUTPCGMO2Output_output,256,,,,,,,,,256,,256,256,,,256,,256,256 +GPUTPCStartHitsFinder,256,,"[1024, 2]","[1024, 7]",256,256,256,256,256,"[224, 1]",512,"[416, 4]",928,,,"[320, 5]",,608,"[448, 3]" +GPUTPCStartHitsSorter,256,,"[1024, 5]","[512, 7]",256,256,256,256,256,"[320, 2]","[512, 1]","[864, 1]","[96, 2]",,,"[192, 5]",,608,"[448, 1]" +GPUTPCCFCheckPadBaseline,576,,"[576, 2]","[576, 2]",,,,,,"[576, 3]",,"[576, 1]","[576, 1]",,,"[576, 2]",,"[576, 2]",576 +GPUTPCCFHIPTailConnector,256,,256,256,,,,,,"[224, 2]",,"[320, 5]","[704, 1]",,,"[128, 7]",,,"[448, 4]" +GPUTPCCFHIPClusterizer,256,,256,256,,,,,,"[288, 5]",,"[480, 3]","[448, 3]",,,352,,,"[512, 3]" +GPUTPCCFChargeMapFiller_fillIndexMap,512,,512,512,,,,,,448,,448,448,,,512,,448,512 +GPUTPCCFChargeMapFiller_fillFromDigits,512,,512,512,,,,,,448,,448,448,,,512,,448,512 +GPUTPCCFChargeMapFiller_findFragmentStart,512,,512,512,,,,,,448,,448,448,,,512,,448,512 +GPUTPCCFPeakFinder,512,,"[512, 9]","[512, 4]",,,,,,416,,992,"[672, 1]",,,"[384, 2]",,"[128, 5]","[192, 10]" +GPUTPCCFNoiseSuppression,512,,512,512,,,,,,608,,896,480,,,160,,,448 +GPUTPCCFDeconvolution,512,,"[512, 5]","[512, 5]",,,,,,"[480, 4]",,224,512,,,480,,384,"[448, 3]" +GPUTPCCFClusterizer,512,,"[448, 3]","[512, 2]",,,,,,"[608, 3]",,736,"[192, 3]",,,576,,"[160, 5]","[832, 2]" +GPUTPCNNClusterizerKernels,512,,,,,,,,,,,,,,,,,, +GPUTrackingRefitKernel_mode0asGPU,256,,,,,,,,,256,,256,256,,,256,,256,256 +GPUTrackingRefitKernel_mode1asTrackParCov,256,,,,,,,,,256,,256,256,,,256,,256,256 +GPUMemClean16,"[""GPUCA_THREAD_COUNT_DEFAULT"", 1]",,,,,,,,,"[512, 1]",,"[512, 1]","[512, 1]",,,"[256, 1]",,,"[256, 1]" +GPUitoa,"[""GPUCA_THREAD_COUNT_DEFAULT"", 1]",,,,,,,,,"[512, 1]",,"[512, 1]","[512, 1]",,,"[256, 1]",,,"[256, 1]" +GPUTPCCFNoiseSuppression_noiseSuppression,"""GPUCA_LB_GPUTPCCFNoiseSuppression""",,,,,,,,,,,,,,,,,448, +GPUTPCCFNoiseSuppression_updatePeaks,"""GPUCA_LB_GPUTPCCFNoiseSuppression""",,,,,,,,,,,,,,,,,448, +GPUTPCNNClusterizerKernels_runCfClusterizer,"""GPUCA_LB_GPUTPCCFClusterizer""",,,,,,,,,,,,,,,,,, +GPUTPCNNClusterizerKernels_fillInputNNCPU,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,,, +GPUTPCNNClusterizerKernels_fillInputNNGPU,1024,,,,,,,,,,,,,,,,,, +GPUTPCNNClusterizerKernels_determineClass1Labels,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,,, +GPUTPCNNClusterizerKernels_determineClass2Labels,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,,, +GPUTPCNNClusterizerKernels_publishClass1Regression,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,,, +GPUTPCNNClusterizerKernels_publishClass2Regression,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,,, +GPUTPCNNClusterizerKernels_publishDeconvolutionFlags,"""GPUCA_LB_GPUTPCNNClusterizerKernels""",,,,,,,,,,,,,,,,,, +GPUTPCCFStreamCompaction_scanStart,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,,, +GPUTPCCFStreamCompaction_scanUp,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,,, +GPUTPCCFStreamCompaction_scanTop,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,,, +GPUTPCCFStreamCompaction_scanDown,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,,, +GPUTPCCFStreamCompaction_compactDigits,"""GPUCA_PAR_CF_SCAN_WORKGROUP_SIZE""",,,,,,,,,,,,,,,,,, +GPUTPCCompressionGatherKernels_unbuffered,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,,, +GPUTPCCompressionGatherKernels_buffered32,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,,, +GPUTPCCompressionGatherKernels_buffered64,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,,, +GPUTPCCompressionGatherKernels_buffered128,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,,, +GPUTPCCompressionGatherKernels_multiBlock,"""GPUCA_LB_COMPRESSION_GATHER""",,,,,,,,,,,,,,,,,, +GPUTPCGMMergerFinalize_0,256,,256,,,,,,,256,256,,256,,,,,256, +GPUTPCGMMergerFinalize_1,256,,256,,,,,,,256,256,,256,,,,,256, +GPUTPCGMMergerFinalize_2,256,,256,,,,,,,256,256,,256,,,,,256, +GPUTPCConvertKernel,,,,,,,,,,,,,,,,256,,,256 +,,,,,,,,,,,,,,,,,,, +PAR:,,,,,,,,,,,,,,,,,,, +AMD_EUS_PER_CU,0,0,4,4,,,,,,,,,,,,4,,0,4 +SORT_STARTHITS,1,0,,,,,,,,1,,1,1,,,1,,1,1 +NEIGHBOURS_FINDER_MAX_NNEIGHUP,6,0,10,4,,,,,,4,4,4,4,,,5,,2,5 +NEIGHBOURS_FINDER_UNROLL_GLOBAL,4,0,4,2,,,,,,2,,8,8,,,4,,2,2 +NEIGHBOURS_FINDER_UNROLL_SHARED,1,0,0,0,,,,,,1,,1,0,,,1,,1,1 +TRACKLET_SELECTOR_HITS_REG_SIZE,12,0,9,27,,,,,,20,20,20,20,,,20,,2,20 +ALTERNATE_BORDER_SORT,1,0,1,1,,,,,,1,1,1,1,,,1,,1,1 +SORT_BEFORE_FIT,1,0,1,1,,,,,,1,1,1,1,,,1,,1,1 +NO_ATOMIC_PRECHECK,0,0,1,1,,,,,,1,1,1,1,,,1,,1,1 +DEDX_STORAGE_TYPE,"""half""","""float""",,,,,,,,,,,,,,,,, +MERGER_INTERPOLATION_ERROR_TYPE,"""half""","""float""",,,,,,,,,,,,,,,,, +COMP_GATHER_KERNEL,4,0,4,4,,,,,,4,4,4,4,,,4,,4,4 +COMP_GATHER_MODE,3,0,3,3,,,,,,3,3,3,3,,,3,,3,3 +CF_SCAN_WORKGROUP_SIZE,512,0,,,,,,,,224,,992,448,,,1024,,,448 +MERGER_SPLIT_LOOP_INTERPOLATION,,,,,,,,,,,,,,,,1,,,1 diff --git a/GPU/GPUTracking/Global/GPUChainTracking.h b/GPU/GPUTracking/Global/GPUChainTracking.h index 759aaf818028e..868400401d7a1 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.h +++ b/GPU/GPUTracking/Global/GPUChainTracking.h @@ -300,11 +300,11 @@ class GPUChainTracking : public GPUChain int32_t RunTPCTrackingSectors_internal(); int32_t RunTPCClusterizer_prepare(bool restorePointers, const GPUTPCExtraADC& extraADCs); #ifndef GPUCA_RUN2 - std::pair RunTPCClusterizer_transferZS(int32_t iSector, const CfFragment& fragment, int32_t lane, const GPUTPCExtraADC& extraADCs); + std::pair RunTPCClusterizer_transferZS(int32_t iSector, const CfFragment& frag, int32_t lane, const GPUTPCExtraADC& extraADCs); void RunTPCClusterizer_compactPeaks(GPUTPCClusterFinder& clusterer, GPUTPCClusterFinder& clustererShadow, int32_t stage, bool doGPU, int32_t lane); - std::pair TPCClusterizerDecodeZSCount(uint32_t iSector, const CfFragment& fragment); - std::pair TPCClusterizerDecodeZSCountUpdate(uint32_t iSector, const CfFragment& fragment); - void TPCClusterizerEnsureZSOffsets(uint32_t iSector, const CfFragment& fragment); + std::pair TPCClusterizerDecodeZSCount(uint32_t iSector, const CfFragment& frag); + std::pair TPCClusterizerDecodeZSCountUpdate(uint32_t iSector, const CfFragment& frag); + void TPCClusterizerEnsureZSOffsets(uint32_t iSector, const CfFragment& frag); void TPCClusterizerTransferExtraADC(GPUTPCClusterFinder& clusterer, GPUTPCClusterFinder& clustererShadow, int lane, const GPUTPCExtraADC& extraADCs); void TPCClusterizerCheckExtraADCZeros(GPUTPCClusterFinder& clusterer, GPUTPCClusterFinder& clustererShadow, int lane, const GPUTPCExtraADC& extraADCs); #endif diff --git a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx index 8c6534d74b31d..1557b2fab0934 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx @@ -70,7 +70,7 @@ using namespace o2::tpc::constants; using namespace o2::dataformats; #ifndef GPUCA_RUN2 -std::pair GPUChainTracking::TPCClusterizerDecodeZSCountUpdate(uint32_t iSector, const CfFragment& fragment) +std::pair GPUChainTracking::TPCClusterizerDecodeZSCountUpdate(uint32_t iSector, const CfFragment& frag) { bool doGPU = mRec->GetRecoStepsGPU() & gpudatatypes::RecoStep::TPCClusterFinding; GPUTPCClusterFinder& clusterer = processors()->tpcClusterer[iSector]; @@ -78,7 +78,7 @@ std::pair GPUChainTracking::TPCClusterizerDecodeZSCountUpdat uint32_t digits = 0; uint32_t pages = 0; for (uint16_t j = 0; j < GPUTrackingInOutZS::NENDPOINTS; j++) { - clusterer.mMinMaxCN[j] = mCFContext->fragmentData[fragment.index].minMaxCN[iSector][j]; + clusterer.mMinMaxCN[j] = mCFContext->fragmentData[frag.index].minMaxCN[iSector][j]; if (doGPU) { uint16_t posInEndpoint = 0; uint16_t pagesEndpoint = 0; @@ -86,7 +86,7 @@ std::pair GPUChainTracking::TPCClusterizerDecodeZSCountUpdat const uint32_t pageFirst = (k == clusterer.mMinMaxCN[j].zsPtrFirst) ? clusterer.mMinMaxCN[j].zsPageFirst : 0; const uint32_t pageLast = (k + 1 == clusterer.mMinMaxCN[j].zsPtrLast) ? clusterer.mMinMaxCN[j].zsPageLast : mIOPtrs.tpcZS->sector[iSector].nZSPtr[j][k]; for (uint32_t l = pageFirst; l < pageLast; l++) { - uint16_t pageDigits = mCFContext->fragmentData[fragment.index].pageDigits[iSector][j][posInEndpoint++]; + uint16_t pageDigits = mCFContext->fragmentData[frag.index].pageDigits[iSector][j][posInEndpoint++]; if (pageDigits) { *(o++) = GPUTPCClusterFinder::ZSOffset{digits, j, pagesEndpoint}; digits += pageDigits; @@ -94,35 +94,35 @@ std::pair GPUChainTracking::TPCClusterizerDecodeZSCountUpdat pagesEndpoint++; } } - if (pagesEndpoint != mCFContext->fragmentData[fragment.index].pageDigits[iSector][j].size()) { + if (pagesEndpoint != mCFContext->fragmentData[frag.index].pageDigits[iSector][j].size()) { if (GetProcessingSettings().ignoreNonFatalGPUErrors) { - GPUError("TPC raw page count mismatch in TPCClusterizerDecodeZSCountUpdate: expected %d / buffered %lu", pagesEndpoint, mCFContext->fragmentData[fragment.index].pageDigits[iSector][j].size()); + GPUError("TPC raw page count mismatch in TPCClusterizerDecodeZSCountUpdate: expected %d / buffered %lu", pagesEndpoint, mCFContext->fragmentData[frag.index].pageDigits[iSector][j].size()); return {0, 0}; } else { - GPUFatal("TPC raw page count mismatch in TPCClusterizerDecodeZSCountUpdate: expected %d / buffered %lu", pagesEndpoint, mCFContext->fragmentData[fragment.index].pageDigits[iSector][j].size()); + GPUFatal("TPC raw page count mismatch in TPCClusterizerDecodeZSCountUpdate: expected %d / buffered %lu", pagesEndpoint, mCFContext->fragmentData[frag.index].pageDigits[iSector][j].size()); } } } else { clusterer.mPzsOffsets[j] = GPUTPCClusterFinder::ZSOffset{digits, j, 0}; - digits += mCFContext->fragmentData[fragment.index].nDigits[iSector][j]; - pages += mCFContext->fragmentData[fragment.index].nPages[iSector][j]; + digits += mCFContext->fragmentData[frag.index].nDigits[iSector][j]; + pages += mCFContext->fragmentData[frag.index].nPages[iSector][j]; } } if (doGPU) { pages = o - processors()->tpcClusterer[iSector].mPzsOffsets; } if (GetProcessingSettings().clusterizerZSSanityCheck && mCFContext->zsVersion >= ZSVersion::ZSVersionDenseLinkBased) { - TPCClusterizerEnsureZSOffsets(iSector, fragment); + TPCClusterizerEnsureZSOffsets(iSector, frag); } return {digits, pages}; } -void GPUChainTracking::TPCClusterizerEnsureZSOffsets(uint32_t iSector, const CfFragment& fragment) +void GPUChainTracking::TPCClusterizerEnsureZSOffsets(uint32_t iSector, const CfFragment& frag) { GPUTPCClusterFinder& clusterer = processors()->tpcClusterer[iSector]; uint32_t nAdcs = 0; for (uint16_t endpoint = 0; endpoint < GPUTrackingInOutZS::NENDPOINTS; endpoint++) { - const auto& data = mCFContext->fragmentData[fragment.index]; + const auto& data = mCFContext->fragmentData[frag.index]; uint32_t pagesEndpoint = 0; const uint32_t nAdcsExpected = data.nDigits[iSector][endpoint]; const uint32_t nPagesExpected = data.nPages[iSector][endpoint]; @@ -144,15 +144,15 @@ void GPUChainTracking::TPCClusterizerEnsureZSOffsets(uint32_t iSector, const CfF } if (pagesEndpoint != nPagesExpected) { - GPUFatal("Sector %d, Endpoint %d, Fragment %d: TPC raw page count mismatch: expected %d / buffered %u", iSector, endpoint, fragment.index, pagesEndpoint, nPagesExpected); + GPUFatal("Sector %d, Endpoint %d, Fragment %d: TPC raw page count mismatch: expected %d / buffered %u", iSector, endpoint, frag.index, pagesEndpoint, nPagesExpected); } if (nAdcDecoded != nAdcsExpected) { - GPUFatal("Sector %d, Endpoint %d, Fragment %d: TPC ADC count mismatch: expected %u, buffered %u", iSector, endpoint, fragment.index, nAdcsExpected, nAdcDecoded); + GPUFatal("Sector %d, Endpoint %d, Fragment %d: TPC ADC count mismatch: expected %u, buffered %u", iSector, endpoint, frag.index, nAdcsExpected, nAdcDecoded); } if (nAdcs != clusterer.mPzsOffsets[endpoint].offset) { - GPUFatal("Sector %d, Endpoint %d, Fragment %d: TPC ADC offset mismatch: expected %u, buffered %u", iSector, endpoint, fragment.index, nAdcs, clusterer.mPzsOffsets[endpoint].offset); + GPUFatal("Sector %d, Endpoint %d, Fragment %d: TPC ADC offset mismatch: expected %u, buffered %u", iSector, endpoint, frag.index, nAdcs, clusterer.mPzsOffsets[endpoint].offset); } nAdcs += nAdcsExpected; @@ -162,10 +162,10 @@ void GPUChainTracking::TPCClusterizerEnsureZSOffsets(uint32_t iSector, const CfF void GPUChainTracking::TPCClusterizerTransferExtraADC(GPUTPCClusterFinder& clusterer, GPUTPCClusterFinder& clustererShadow, int lane, const GPUTPCExtraADC& extraADCs) { const int32_t iSector = clusterer.mISector; - const auto& fragment = clusterer.mPmemory->fragment; + const auto& frag = clusterer.mPmemory->frag; const auto& digits = extraADCs.digitsBySector[iSector]; - if (fragment.index != 0) { + if (frag.index != 0) { return; } @@ -188,11 +188,11 @@ void GPUChainTracking::TPCClusterizerTransferExtraADC(GPUTPCClusterFinder& clust SynchronizeStream(lane); for (const auto& d : digits) { - if (!fragment.contains(d.getTimeStamp())) { + if (!frag.contains(d.getTimeStamp())) { continue; } - CfChargePos pos{(tpccf::Row)d.getRow(), (tpccf::Pad)d.getPad(), (tpccf::TPCFragmentTime)(d.getTimeStamp() - fragment.start)}; + CfChargePos pos{(tpccf::Row)d.getRow(), (tpccf::Pad)d.getPad(), (tpccf::TPCFragmentTime)(d.getTimeStamp() - frag.start)}; chargeMapHost[pos] = PackedCharge(d.getChargeFloat()); extraPositions.push_back(pos); @@ -208,10 +208,10 @@ void GPUChainTracking::TPCClusterizerTransferExtraADC(GPUTPCClusterFinder& clust void GPUChainTracking::TPCClusterizerCheckExtraADCZeros(GPUTPCClusterFinder& clusterer, GPUTPCClusterFinder& clustererShadow, int lane, const GPUTPCExtraADC& extraADCs) { const int32_t iSector = clusterer.mISector; - const auto& fragment = clusterer.mPmemory->fragment; + const auto& frag = clusterer.mPmemory->frag; const auto& digits = extraADCs.digitsBySector[iSector]; - if (fragment.index != 0) { + if (frag.index != 0) { return; } @@ -233,11 +233,11 @@ void GPUChainTracking::TPCClusterizerCheckExtraADCZeros(GPUTPCClusterFinder& clu size_t nNonZeroADCs = 0; for (const auto& d : digits) { - if (!fragment.contains(d.getTimeStamp())) { + if (!frag.contains(d.getTimeStamp())) { continue; } - CfChargePos pos{(tpccf::Row)d.getRow(), (tpccf::Pad)d.getPad(), (tpccf::TPCFragmentTime)(d.getTimeStamp() - fragment.start)}; + CfChargePos pos{(tpccf::Row)d.getRow(), (tpccf::Pad)d.getPad(), (tpccf::TPCFragmentTime)(d.getTimeStamp() - frag.start)}; auto adc = chargeMapHost[pos].unpack(); @@ -326,7 +326,7 @@ GPUTPCExtraADC GenerateSaturatedSignals(size_t seed = 42) } // namespace -std::pair GPUChainTracking::TPCClusterizerDecodeZSCount(uint32_t iSector, const CfFragment& fragment) +std::pair GPUChainTracking::TPCClusterizerDecodeZSCount(uint32_t iSector, const CfFragment& frag) { mRec->getGeneralStepTimer(GeneralStep::Prepare).Start(); uint32_t nDigits = 0; @@ -349,7 +349,7 @@ std::pair GPUChainTracking::TPCClusterizerDecodeZSCount(uint std::vector> fragments; fragments.reserve(mCFContext->nFragments); - fragments.emplace_back(std::pair{fragment, {0, 0, 0, 0, 0, -1}}); + fragments.emplace_back(std::pair{frag, {0, 0, 0, 0, 0, -1}}); for (uint32_t i = 1; i < mCFContext->nFragments; i++) { fragments.emplace_back(std::pair{fragments.back().first.next(), {0, 0, 0, 0, 0, -1}}); } @@ -601,14 +601,14 @@ void GPUChainTracking::RunTPCClusterizer_compactPeaks(GPUTPCClusterFinder& clust } } -std::pair GPUChainTracking::RunTPCClusterizer_transferZS(int32_t iSector, const CfFragment& fragment, int32_t lane, const GPUTPCExtraADC& extraADCs) +std::pair GPUChainTracking::RunTPCClusterizer_transferZS(int32_t iSector, const CfFragment& frag, int32_t lane, const GPUTPCExtraADC& extraADCs) { bool doGPU = GetRecoStepsGPU() & RecoStep::TPCClusterFinding; if (mCFContext->abandonTimeframe) { return {0, 0}; } - auto retVal = TPCClusterizerDecodeZSCountUpdate(iSector, fragment); - if (fragment.index == 0) { + auto retVal = TPCClusterizerDecodeZSCountUpdate(iSector, frag); + if (frag.index == 0) { retVal.first += extraADCs.digitsBySector[iSector].size(); } if (doGPU) { @@ -994,12 +994,12 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) std::vector laneHasData(GetProcessingSettings().nTPCClustererLanes, false); static_assert(NSECTORS <= constants::GPU_MAX_STREAMS, "Stream events must be able to hold all sectors"); const int32_t maxLane = std::min(GetProcessingSettings().nTPCClustererLanes, NSECTORS - iSectorBase); - for (CfFragment fragment = mCFContext->fragmentFirst; !fragment.isEnd(); fragment = fragment.next()) { + for (CfFragment frag = mCFContext->fragmentFirst; !frag.isEnd(); frag = frag.next()) { if (GetProcessingSettings().debugLevel >= 3) { - GPUInfo("Processing time bins [%d, %d) for sectors %d to %d", fragment.start, fragment.last(), iSectorBase, iSectorBase + GetProcessingSettings().nTPCClustererLanes - 1); + GPUInfo("Processing time bins [%d, %d) for sectors %d to %d", frag.start, frag.last(), iSectorBase, iSectorBase + GetProcessingSettings().nTPCClustererLanes - 1); } mRec->runParallelOuterLoop(doGPU, maxLane, [&](uint32_t lane) { - if (doGPU && fragment.index != 0) { + if (doGPU && frag.index != 0) { SynchronizeStream(lane); // Don't overwrite charge map from previous iteration until cluster computation is finished } @@ -1007,7 +1007,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) GPUTPCClusterFinder& clusterer = processors()->tpcClusterer[iSector]; GPUTPCClusterFinder& clustererShadow = doGPU ? processorsShadow()->tpcClusterer[iSector] : clusterer; clusterer.mPmemory->counters.nPeaks = clusterer.mPmemory->counters.nClusters = 0; - clusterer.mPmemory->fragment = fragment; + clusterer.mPmemory->frag = frag; if (mIOPtrs.tpcPackedDigits) { bool setDigitsOnGPU = doGPU && not mIOPtrs.tpcZS; @@ -1037,7 +1037,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) using PeakMapType = decltype(*clustererShadow.mPpeakMap); runKernel({GetGridAutoStep(lane, RecoStep::TPCClusterFinding)}, clustererShadow.mPchargeMap, TPCMapMemoryLayout::items(GetProcessingSettings().overrideClusterizerFragmentLen) * sizeof(ChargeMapType)); runKernel({GetGridAutoStep(lane, RecoStep::TPCClusterFinding)}, clustererShadow.mPpeakMap, TPCMapMemoryLayout::items(GetProcessingSettings().overrideClusterizerFragmentLen) * sizeof(PeakMapType)); - if (fragment.index == 0) { + if (frag.index == 0) { runKernel({GetGridAutoStep(lane, RecoStep::TPCClusterFinding)}, clustererShadow.mPpadIsNoisy, TPC_CLUSTERER_STRIDED_PAD_COUNT * sizeof(*clustererShadow.mPpadIsNoisy)); } DoDebugAndDump(RecoStep::TPCClusterFinding, GPUChainTrackingDebugFlags::TPCClustererZeroedCharges, clusterer, &GPUTPCClusterFinder::DumpChargeMap, *mDebugFile, "Zeroed Charges"); @@ -1060,7 +1060,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) } if (propagateMCLabels) { - if (fragment.index == 0) { + if (frag.index == 0) { // Must be only called on the first fragment as some buffers are used across the whole timeframe clusterer.AllocMCBuffers(); } @@ -1112,7 +1112,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) SynchronizeStream(lane); } if (mIOPtrs.tpcZS) { - CfFragment f = fragment.next(); + CfFragment f = frag.next(); int32_t nextSector = iSector; if (f.isEnd()) { nextSector += GetProcessingSettings().nTPCClustererLanes; @@ -1138,7 +1138,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) } bool checkForNoisyPads = (rec()->GetParam().rec.tpc.maxTimeBinAboveThresholdIn1000Bin > 0) || (rec()->GetParam().rec.tpc.maxConsecTimeBinAboveThreshold > 0); - checkForNoisyPads &= (rec()->GetParam().rec.tpc.noisyPadsQuickCheck ? fragment.index == 0 : true); + checkForNoisyPads &= (rec()->GetParam().rec.tpc.noisyPadsQuickCheck ? frag.index == 0 : true); checkForNoisyPads &= !GetProcessingSettings().disableTPCNoisyPadFilter; // TODO Move hipTailFilter flag to ProcessingSettings? // TODO Add some warning when re enabling pad filter with this flag, so it's not just silently enabled when disabling was requested @@ -1152,7 +1152,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) const int32_t nBlocks = GPUTPCGeometry::NROWS; runKernel({GetGridBlk(nBlocks, lane), {iSector}}); - getKernelTimer(RecoStep::TPCClusterFinding, iSector, TPC_REAL_PADS_IN_SECTOR * fragment.lengthWithoutOverlap() * sizeof(PackedCharge), false); + getKernelTimer(RecoStep::TPCClusterFinding, iSector, TPC_REAL_PADS_IN_SECTOR * frag.lengthWithoutOverlap() * sizeof(PackedCharge), false); } DoDebugAndDump(RecoStep::TPCClusterFinding, GPUChainTrackingDebugFlags::TPCClustererDigits, clusterer, &GPUTPCClusterFinder::DumpDigits, *mDebugFile); @@ -1197,7 +1197,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) GPUTPCClusterFinder& clusterer = processors()->tpcClusterer[iSector]; GPUTPCClusterFinder& clustererShadow = doGPU ? processorsShadow()->tpcClusterer[iSector] : clusterer; - const bool resetClusterCounters = fragment.index == 0; + const bool resetClusterCounters = frag.index == 0; // The reset must also run for an empty first fragment since later fragments can contain data. if (clusterer.mPmemory->counters.nPositions == 0 && !resetClusterCounters) { return; @@ -1403,7 +1403,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) } if (GetProcessingSettings().debugLevel >= 3) { - GPUInfo("Sector %02d Fragment %02d Lane %d: Found clusters: digits %u peaks %u clusters %u", iSector, fragment.index, lane, (int32_t)clusterer.mPmemory->counters.nPositions, (int32_t)clusterer.mPmemory->counters.nPeaks, (int32_t)clusterer.mPmemory->counters.nClusters); + GPUInfo("Sector %02d Fragment %02d Lane %d: Found clusters: digits %u peaks %u clusters %u", iSector, frag.index, lane, (int32_t)clusterer.mPmemory->counters.nPositions, (int32_t)clusterer.mPmemory->counters.nPeaks, (int32_t)clusterer.mPmemory->counters.nClusters); } TransferMemoryResourcesToHost(RecoStep::TPCClusterFinding, &clusterer, lane); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index e2ddcc524c683..4598e61993201 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -1201,6 +1201,7 @@ GPUd() void GPUTPCGMMerger::ResolveMergeSectors(GPUResolveSharedMemory& smem, in // PrintMergeGraph(track1, std::cout); // PrintMergeGraph(track2, std::cout); + bool nextTrack = false; while (track2->PrevSegmentNeighbour() >= 0) { track2 = &mSectorTrackInfos[track2->PrevSegmentNeighbour()]; } @@ -1211,26 +1212,41 @@ GPUd() void GPUTPCGMMerger::ResolveMergeSectors(GPUResolveSharedMemory& smem, in while (track1->PrevSegmentNeighbour() >= 0) { track1 = &mSectorTrackInfos[track1->PrevSegmentNeighbour()]; if (track1 == track2) { - goto NextTrack; + nextTrack = true; + break; } } + if (nextTrack) { + continue; + } GPUCommonAlgorithm::swap(track1, track1Base); for (int32_t k = 0; k < 2; k++) { GPUTPCGMSectorTrack* tmp = track1Base; while (tmp->Neighbour(k) >= 0) { tmp = &mSectorTrackInfos[tmp->Neighbour(k)]; if (tmp == track2) { - goto NextTrack; + nextTrack = true; + break; } } + if (nextTrack) { + break; + } + } + if (nextTrack) { + continue; } while (track1->NextSegmentNeighbour() >= 0) { track1 = &mSectorTrackInfos[track1->NextSegmentNeighbour()]; if (track1 == track2) { - goto NextTrack; + nextTrack = true; + break; } } + if (nextTrack) { + continue; + } } else { while (track1->PrevSegmentNeighbour() >= 0) { track1 = &mSectorTrackInfos[track1->PrevSegmentNeighbour()]; @@ -1244,9 +1260,16 @@ GPUd() void GPUTPCGMMerger::ResolveMergeSectors(GPUResolveSharedMemory& smem, in while (tmp->Neighbour(k) >= 0) { tmp = &mSectorTrackInfos[tmp->Neighbour(k)]; if (tmp == track2) { - goto NextTrack; + nextTrack = true; + break; } } + if (nextTrack) { + break; + } + } + if (nextTrack) { + continue; } float z1min, z1max, z2min, z2max; @@ -1318,7 +1341,6 @@ GPUd() void GPUTPCGMMerger::ResolveMergeSectors(GPUResolveSharedMemory& smem, in } // GPUInfo("Result"); // PrintMergeGraph(track1, std::cout); - NextTrack:; } } } @@ -1934,7 +1956,7 @@ GPUd() void GPUTPCGMMerger::Finalize2(int32_t nBlocks, int32_t nThreads, int32_t GPUd() void GPUTPCGMMerger::MergeLoopersInit(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) { const float lowPtThresh = Param().rec.tpc.rejectQPtB5 * 1.1f; // Might need to merge tracks above the threshold with parts below the rejection threshold - for (uint32_t i = get_global_id(0); i < mMemory->nMergedTracks; i += get_global_size(0)) { + for (uint32_t i = (iBlock * nThreads + iThread); i < mMemory->nMergedTracks; i += (nBlocks * nThreads)) { const auto& trk = mMergedTracks[i]; const auto& p = trk.GetParam(); const float qptabs = CAMath::Abs(p.GetQPt()); @@ -2003,7 +2025,7 @@ GPUd() void GPUTPCGMMerger::MergeLoopersMain(int32_t nBlocks, int32_t nThreads, } #endif - for (uint32_t i = get_global_id(0); i < mMemory->nLooperMatchCandidates; i += get_global_size(0)) { + for (uint32_t i = (iBlock * nThreads + iThread); i < mMemory->nLooperMatchCandidates; i += (nBlocks * nThreads)) { for (uint32_t j = i + 1; j < mMemory->nLooperMatchCandidates; j++) { // int32_t bs = 0; assert(CAMath::Abs(candidates[i].refz) <= CAMath::Abs(candidates[j].refz)); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx index 2a111b8ce89af..ea0620d3ea4c1 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx @@ -22,7 +22,7 @@ template <> GPUdii() void GPUTPCGMMergerTrackFit::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int32_t mode) { const int32_t iEnd = mode == -1 ? merger.Memory()->nRetryRefit : merger.NMergedTracks(); - GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), int32_t, ii, iEnd, { + GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), nBlocks, nThreads, iBlock, iThread, int32_t, ii, iEnd, { const int32_t i = mode == -1 ? merger.RetryRefitIds()[ii] : mode ? merger.TrackOrderProcess()[ii] : ii; GPUTPCGMTrackParam::RefitTrack(merger.MergedTracks()[i], i, &merger, mode == -1); }); @@ -31,7 +31,7 @@ GPUdii() void GPUTPCGMMergerTrackFit::Thread<0>(int32_t nBlocks, int32_t nThread template <> GPUdii() void GPUTPCGMMergerFollowLoopers::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger) { - GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), uint32_t, i, merger.Memory()->nLoopData, { + GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), nBlocks, nThreads, iBlock, iThread, uint32_t, i, merger.Memory()->nLoopData, { GPUTPCGMTrackParam::PropagateLooper(&merger, i); }); } diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h index 5d00451516aa8..d703f1b890443 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h @@ -81,7 +81,7 @@ class GPUTPCGMMergerResolve : public GPUTPCGMMergerGeneral }; template - GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, Args... args); + GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& clusterer, Args... args); }; class GPUTPCGMMergerClearLinks : public GPUTPCGMMergerGeneral diff --git a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx index d7c8eb9c44aab..39690ae078599 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx @@ -58,7 +58,7 @@ GPUdii() void GPUTPCGMO2Output::Thread(int32_t nBlock GPUTPCGMMerger::tmpSort* GPUrestrict() trackSort = merger.TrackSortO2(); uint2* GPUrestrict() tmpData = merger.ClusRefTmp(); - for (uint32_t i = get_global_id(0); i < nTracks; i += get_global_size(0)) { + for (uint32_t i = (iBlock * nThreads + iThread); i < nTracks; i += (nBlocks * nThreads)) { if (!tracks[i].OK()) { continue; } @@ -120,7 +120,7 @@ GPUdii() void GPUTPCGMO2Output::Thread(int32_t nBlocks uint2* GPUrestrict() tmpData = merger.ClusRefTmp(); float const SNPThresh = 0.999990f; - for (int32_t iTmp = get_global_id(0); iTmp < nTracks; iTmp += get_global_size(0)) { + for (int32_t iTmp = (iBlock * nThreads + iThread); iTmp < nTracks; iTmp += (nBlocks * nThreads)) { TrackTPC oTrack; const int32_t i = trackSort[iTmp].x; const auto& track = tracks[i]; @@ -288,7 +288,7 @@ GPUdii() void GPUTPCGMO2Output::Thread(int32_t nBlocks, in auto labelAssigner = GPUTPCTrkLbl(clusters->clustersMCTruth, 0.1f); uint32_t* clusRefs = merger.OutputClusRefsTPCO2(); - for (uint32_t i = get_global_id(0); i < merger.NOutputTracksTPCO2(); i += get_global_size(0)) { + for (uint32_t i = (iBlock * nThreads + iThread); i < merger.NOutputTracksTPCO2(); i += (nBlocks * nThreads)) { labelAssigner.reset(); const auto& trk = merger.OutputTracksTPCO2()[i]; for (int32_t j = 0; j < trk.getNClusters(); j++) { diff --git a/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx b/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx index f99544f239bb7..bc67075b4f820 100644 --- a/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx +++ b/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx @@ -22,7 +22,7 @@ template GPUdii() void GPUTrackingRefitKernel::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors) { auto& refit = processors.trackingRefit; - for (uint32_t i = get_global_id(0); i < processors.ioPtrs.nMergedTracks; i += get_global_size(0)) { + for (uint32_t i = (iBlock * nThreads + iThread); i < processors.ioPtrs.nMergedTracks; i += (nBlocks * nThreads)) { if (refit.mPTracks[i].OK()) { GPUTPCGMMergedTrack trk = refit.mPTracks[i]; int32_t retval; diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx index 784b60baec3d6..6280d6e7afaf2 100644 --- a/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx +++ b/GPU/GPUTracking/SectorTracker/GPUTPCExtrapolationTracking.cxx @@ -160,7 +160,7 @@ GPUd() void GPUTPCExtrapolationTracking::PerformExtrapolationTracking(int32_t nB template <> GPUdii() void GPUTPCExtrapolationTracking::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() tracker) { - GPUCA_SHARED_CACHE(&smem.mRows[0], tracker.TrackingDataRows(), GPUTPCGeometry::NROWS * sizeof(GPUTPCRow)); + GPUCA_SHARED_CACHE(nThreads, iThread, &smem.mRows[0], tracker.TrackingDataRows(), GPUTPCGeometry::NROWS * sizeof(GPUTPCRow)); GPUbarrier(); if (tracker.NHitsTotal() == 0) { @@ -202,7 +202,7 @@ GPUd() void GPUTPCExtrapolationTracking::ExtrapolationTrackingSectorLeftRight(ui template <> GPUdii() void GPUTPCExtrapolationTrackingCopyNumbers::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() tracker, int32_t n) { - for (int32_t i = get_global_id(0); i < n; i += get_global_size(0)) { + for (int32_t i = (iBlock * nThreads + iThread); i < n; i += (nBlocks * nThreads)) { GPUconstantref() GPUTPCTracker& GPUrestrict() trk = (&tracker)[i]; trk.CommonMemory()->nLocalTracks = trk.CommonMemory()->nTracks; trk.CommonMemory()->nLocalTrackHits = trk.CommonMemory()->nTrackHits; diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx index 33a3264a87ab3..ba31c30840631 100644 --- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx +++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx @@ -479,14 +479,14 @@ GPUdic(2, 1) void GPUTPCTrackletConstructor::DoTracklet(GPUconstantref() GPUTPCT template <> GPUdii() void GPUTPCTrackletConstructor::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& sMem, processorType& GPUrestrict() tracker) { - if (get_local_id(0) == 0) { + if (iThread == 0) { sMem.mNStartHits = *tracker.NStartHits(); } - GPUCA_SHARED_CACHE(&sMem.mRows[0], tracker.TrackingDataRows(), GPUTPCGeometry::NROWS * sizeof(GPUTPCRow)); + GPUCA_SHARED_CACHE(nThreads, iThread, &sMem.mRows[0], tracker.TrackingDataRows(), GPUTPCGeometry::NROWS * sizeof(GPUTPCRow)); GPUbarrier(); GPUTPCThreadMemory rMem; - for (rMem.mISH = get_global_id(0); rMem.mISH < sMem.mNStartHits; rMem.mISH += get_global_size(0)) { + for (rMem.mISH = (iBlock * nThreads + iThread); rMem.mISH < sMem.mNStartHits; rMem.mISH += (nBlocks * nThreads)) { rMem.mGo = 1; DoTracklet(tracker, sMem, rMem); } diff --git a/GPU/GPUTracking/TPCClusterFinder/CfChargePos.h b/GPU/GPUTracking/TPCClusterFinder/CfChargePos.h index 3f1265e6d0634..64aa0f6fcfe6d 100644 --- a/GPU/GPUTracking/TPCClusterFinder/CfChargePos.h +++ b/GPU/GPUTracking/TPCClusterFinder/CfChargePos.h @@ -34,6 +34,14 @@ struct CfChargePos { : gpad(tpcGlobalPadIdx(row, pad)), timePadded(t + GPUCF_PADDING_TIME) { } +#ifdef __METAL__ + // INVALID_CHARGE_POS below lives in the constant address space, which a + // generic `this` does not reach in MSL. + constexpr GPUhdi() CfChargePos(tpccf::Row row, tpccf::Pad pad, tpccf::TPCFragmentTime t) constant + : gpad(tpcGlobalPadIdx(row, pad)), timePadded(t + GPUCF_PADDING_TIME) + { + } +#endif GPUdi() CfChargePos(const tpccf::GlobalPad& p, const tpccf::TPCFragmentTime& t) : gpad(p), timePadded(t) {} diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.cxx index 752c85634f928..95f78bd9890ee 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.cxx @@ -24,22 +24,22 @@ template <> GPUdii() void GPUTPCCFChargeMapFiller::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer) { CfArray2D indexMap(clusterer.mPindexMap); - fillIndexMapImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer.mPmemory->fragment, clusterer.mPdigits, indexMap, clusterer.mPmemory->counters.nDigitsInFragment); + fillIndexMapImpl(nBlocks, nThreads, iBlock, iThread, clusterer.mPmemory->frag, clusterer.mPdigits, indexMap, clusterer.mPmemory->counters.nDigitsInFragment); } GPUd() void GPUTPCCFChargeMapFiller::fillIndexMapImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, - const CfFragment& fragment, + const CfFragment& frag, const tpc::Digit* digits, CfArray2D& indexMap, size_t maxDigit) { - size_t idx = get_global_id(0); + size_t idx = (iBlock * nThreads + iThread); if (idx >= maxDigit) { return; } - CPU_ONLY(idx += fragment.digitsStart); + CPU_ONLY(idx += frag.digitsStart); CPU_ONLY(tpc::Digit digit = digits[idx]); - CPU_ONLY(CfChargePos pos(digit.getRow(), digit.getPad(), fragment.toLocal(digit.getTimeStamp()))); + CPU_ONLY(CfChargePos pos(digit.getRow(), digit.getPad(), frag.toLocal(digit.getTimeStamp()))); CPU_ONLY(indexMap.safeWrite(pos, idx)); } @@ -47,21 +47,21 @@ template <> GPUdii() void GPUTPCCFChargeMapFiller::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer) { CfArray2D chargeMap(reinterpret_cast(clusterer.mPchargeMap)); - fillFromDigitsImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->fragment, clusterer.mPmemory->counters.nPositions, clusterer.mPdigits, clusterer.mPpositions, chargeMap); + fillFromDigitsImpl(nBlocks, nThreads, iBlock, iThread, clusterer, clusterer.mPmemory->frag, clusterer.mPmemory->counters.nPositions, clusterer.mPdigits, clusterer.mPpositions, chargeMap); } -GPUd() void GPUTPCCFChargeMapFiller::fillFromDigitsImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, processorType& clusterer, const CfFragment& fragment, size_t digitNum, +GPUd() void GPUTPCCFChargeMapFiller::fillFromDigitsImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, processorType& clusterer, const CfFragment& frag, size_t digitNum, const tpc::Digit* digits, CfChargePos* positions, CfArray2D& chargeMap) { - size_t idx = get_global_id(0); + size_t idx = (iBlock * nThreads + iThread); if (idx >= digitNum) { return; } - tpc::Digit digit = digits[fragment.digitsStart + idx]; + tpc::Digit digit = digits[frag.digitsStart + idx]; - CfChargePos pos(digit.getRow(), digit.getPad(), fragment.toLocal(digit.getTimeStamp())); + CfChargePos pos(digit.getRow(), digit.getPad(), frag.toLocal(digit.getTimeStamp())); positions[idx] = pos; float q = digit.getChargeFloat(); q *= clusterer.GetConstantMem()->calibObjects.tpcPadGain->getGainCorrection(clusterer.mISector, digit.getRow(), digit.getPad()); @@ -77,10 +77,10 @@ GPUdii() void GPUTPCCFChargeMapFiller::Threadcounters.nDigits; const tpc::Digit* digits = clusterer.mPdigits; - size_t st = findTransition(clusterer.mPmemory->fragment.first(), digits, nDigits, 0); - size_t end = findTransition(clusterer.mPmemory->fragment.last(), digits, nDigits, st); + size_t st = findTransition(clusterer.mPmemory->frag.first(), digits, nDigits, 0); + size_t end = findTransition(clusterer.mPmemory->frag.last(), digits, nDigits, st); - clusterer.mPmemory->fragment.digitsStart = st; + clusterer.mPmemory->frag.digitsStart = st; size_t elems = end - st; diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx index 7470f86877490..c1a459d138dbd 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx @@ -224,7 +224,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineGPU(int32_t nBlocks, int32_t return; } - const CfFragment& fragment = clusterer.mPmemory->fragment; + const CfFragment& frag = clusterer.mPmemory->frag; const bool hipFilterOn = clusterer.Param().rec.tpc.hipTailFilter; const Charge hipTailThreshold = clusterer.Param().rec.tpc.hipTailFilterThreshold; const Charge hipTailFilterAlpha = clusterer.Param().rec.tpc.hipTailFilterAlpha; @@ -253,7 +253,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineGPU(int32_t nBlocks, int32_t // saturated signal in overlap region can create tails in the next fragment // even when cleared in current fragment as they're decoded twice const TPCFragmentTime firstTB = 0; - const TPCFragmentTime lastTB = fragment.length; + const TPCFragmentTime lastTB = frag.length; for (uint16_t t = firstTB; t < lastTB; t += NumOfCachedTBs) { @@ -384,7 +384,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t const int32_t nPads = geo.NPads(row); const int32_t nVecPads = (nPads + PadsPerCacheline - 1) / PadsPerCacheline; - const CfFragment& fragment = clusterer.mPmemory->fragment; + const CfFragment& frag = clusterer.mPmemory->frag; const bool hipFilterOn = clusterer.Param().rec.tpc.hipTailFilter; const Charge hipTailThreshold = clusterer.Param().rec.tpc.hipTailFilterThreshold; const Charge hipTailFilterAlpha = clusterer.Param().rec.tpc.hipTailFilterAlpha; @@ -409,7 +409,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t std::vector activeHIPTailEndV(nVecPads, -1); std::vector tailFilterChargeV(nVecPads, Charge8{Vc::Zero}); - for (int16_t t = 0; t < fragment.length; t += NumOfCachedTBs) { + for (int16_t t = 0; t < frag.length; t += NumOfCachedTBs) { bool hasAnyTrigger = false; @@ -432,7 +432,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t for (tpccf::TPCFragmentTime localtime = 0; localtime < NumOfCachedTBs; localtime++) { const uint16_t* packedChargeStart = reinterpret_cast(&chargeMap[basePos.delta({0, localtime})]); - const UShort8 packedCharges = t + localtime < fragment.length + const UShort8 packedCharges = t + localtime < frag.length ? UShort8{packedChargeStart, Vc::Aligned} : UShort8{Vc::Zero}; const auto isCharge = packedCharges != 0; @@ -585,7 +585,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t auto activeHIPTailEnd = activeHIPTailEndV[iVecPad]; const auto shouldCloseTail = activeHIPTailStart > -1; - activeHIPTailEnd(shouldCloseTail && activeHIPTailEnd < 0) = fragment.length; + activeHIPTailEnd(shouldCloseTail && activeHIPTailEnd < 0) = frag.length; if (hipFilterOn && shouldCloseTail.isNotEmpty()) { for (int16_t p = 0; p < PadsPerCacheline; p++) { @@ -639,8 +639,8 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t GPUd() void GPUTPCCFCheckPadBaseline::updatePadBaseline(int32_t pad, const GPUTPCClusterFinder& clusterer, int32_t totalCharges, int32_t consecCharges, Charge maxCharge) { - const CfFragment& fragment = clusterer.mPmemory->fragment; - const int32_t totalChargesBaseline = clusterer.Param().rec.tpc.maxTimeBinAboveThresholdIn1000Bin * fragment.lengthWithoutOverlap() / 1000; + const CfFragment& frag = clusterer.mPmemory->frag; + const int32_t totalChargesBaseline = clusterer.Param().rec.tpc.maxTimeBinAboveThresholdIn1000Bin * frag.lengthWithoutOverlap() / 1000; const int32_t consecChargesBaseline = clusterer.Param().rec.tpc.maxConsecTimeBinAboveThreshold; const uint16_t saturationThreshold = clusterer.Param().rec.tpc.noisyPadSaturationThreshold; const bool isNoisy = (!saturationThreshold || maxCharge < saturationThreshold) && ((totalChargesBaseline > 0 && totalCharges >= totalChargesBaseline) || (consecChargesBaseline > 0 && consecCharges >= consecChargesBaseline)); @@ -668,7 +668,7 @@ GPUd() void GPUTPCCFHIPTailConnector::Thread<0>(int32_t nBlocks, int32_t nThread #ifdef GPUCA_DETERMINISTIC_MODE // Races in tail comparisons and atomic swap can lead to slightly different clusters. // So need a sequential fallback for deterministic mode - GPUCommonAlgorithm::sortInBlock(tails + 1, tails + nTails + 1, [](auto&& t1, auto&& t2) { + GPUCommonAlgorithm::sortInBlock(nThreads, iThread, tails + 1, tails + nTails + 1, [](auto&& t1, auto&& t2) { if (t1.pad != t2.pad) { return t1.pad < t2.pad; } else if (t1.tailStart != t2.tailStart) { @@ -728,7 +728,7 @@ GPUd() void GPUTPCCFHIPClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads, nTails = CAMath::Min(nTails, (uint32_t)MaxHIPTailsPerRow - 1); const auto* tails = GetHIPTails(clusterer, row); - const auto& fragment = clusterer.mPmemory->fragment; + const auto& frag = clusterer.mPmemory->frag; auto* clusterPosInRow = clusterer.mPhipClusterPosInRow ? clusterer.mPhipClusterPosInRow + row * MaxHIPTailsPerRow @@ -776,7 +776,7 @@ GPUd() void GPUTPCCFHIPClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads, cn.qMax = qMax; cn.setSaturatedQtot(qTot); cn.setSaturatedTailLength(tailEnd - tailStart); - float clusterTime = fragment.start + timeMean - clusterer.Param().rec.tpc.clustersShiftTimebinsClusterizer; + float clusterTime = frag.start + timeMean - clusterer.Param().rec.tpc.clustersShiftTimebinsClusterizer; cn.setTimeFlags(clusterTime, 0); cn.setPad(padMean); cn.setSigmaPad(padSigma); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx index c9a8c093153a2..37db02d1b8559 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx @@ -35,5 +35,5 @@ GPUdii() void GPUTPCCFClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads, tpc::ClusterNative* clusterOut = onlyMC ? nullptr : clusterer.mPclusterByRow; - GPUTPCCFClusterizer::computeClustersImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->fragment, smem, chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow, true); + GPUTPCCFClusterizer::computeClustersImpl(nBlocks, nThreads, iBlock, iThread, clusterer, clusterer.mPmemory->frag, smem, chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow, true); } diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h index ce673c778e42d..ab40e1c3bc2c3 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.h @@ -59,7 +59,7 @@ class GPUTPCCFClusterizer : public GPUKernelTemplate static GPUd() void computeClustersImpl(int32_t, int32_t, int32_t, int32_t, processorType&, const CfFragment&, GPUSharedMemory&, const CfArray2D&, const CfChargePos*, const GPUSettingsRec&, MCLabelAccumulator*, uint32_t, uint32_t, uint32_t*, tpc::ClusterNative*, uint32_t*, int8_t); - static GPUd() void buildCluster(const GPUSettingsRec&, const CfArray2D&, CfChargePos, CfChargePos*, PackedCharge*, uint8_t*, ClusterAccumulator*, MCLabelAccumulator*); + static GPUd() void buildCluster(const GPUSettingsRec&, uint16_t, const CfArray2D&, CfChargePos, CfChargePos*, PackedCharge*, uint8_t*, ClusterAccumulator*, MCLabelAccumulator*); static GPUd() uint32_t sortIntoBuckets(processorType&, const tpc::ClusterNative&, uint32_t, uint32_t, uint32_t*, tpc::ClusterNative*); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc index ca396f8aab83e..f21f9b47480af 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc @@ -17,7 +17,7 @@ GPUdii() void GPUTPCCFClusterizer::computeClustersImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, processorType& clusterer, - const CfFragment& fragment, + const CfFragment& frag, GPUSharedMemory& smem, const CfArray2D& chargeMap, const CfChargePos* filteredPeakPositions, @@ -30,7 +30,7 @@ GPUdii() void GPUTPCCFClusterizer::computeClustersImpl(int32_t nBlocks, int32_t uint32_t* clusterPosInRow, int8_t isAccepted) { - uint32_t idx = get_global_id(0); + uint32_t idx = (iBlock * nThreads + iThread); // For certain configurations dummy work items are added, so the total // number of work items is dividable by 64. @@ -43,6 +43,7 @@ GPUdii() void GPUTPCCFClusterizer::computeClustersImpl(int32_t nBlocks, int32_t buildCluster( calib, + iThread, chargeMap, pos, smem.posBcast, @@ -54,14 +55,14 @@ GPUdii() void GPUTPCCFClusterizer::computeClustersImpl(int32_t nBlocks, int32_t if (idx >= clusternum) { return; } - if (fragment.isOverlap(pos.time())) { + if (frag.isOverlap(pos.time())) { if (clusterPosInRow) { clusterPosInRow[idx] = maxClusterPerRow; } return; } tpc::ClusterNative myCluster; - pc.finalize(pos, charge, fragment.start); + pc.finalize(pos, charge, frag.start); bool rejectCluster = !pc.toNative(pos, charge, myCluster, clusterer.Param(), chargeMap); if (!isAccepted) { rejectCluster = true; @@ -145,6 +146,7 @@ GPUdii() void GPUTPCCFClusterizer::updateClusterOuter( GPUdii() void GPUTPCCFClusterizer::buildCluster( const GPUSettingsRec& calib, + uint16_t ll, const CfArray2D& chargeMap, CfChargePos pos, CfChargePos* posBcast, @@ -153,8 +155,6 @@ GPUdii() void GPUTPCCFClusterizer::buildCluster( ClusterAccumulator* myCluster, MCLabelAccumulator* labelAcc) { - uint16_t ll = get_local_id(0); - posBcast[ll] = pos; GPUbarrier(); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx index a1c4a3dc4aadd..0f2d30362b4f0 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx @@ -37,12 +37,12 @@ using namespace o2::tpc::constants; // =========================================================================== template <> -GPUdii() void GPUTPCCFDecodeZS::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, int32_t firstHBF, int32_t tpcTimeBinCut) +GPUdii() void GPUTPCCFDecodeZS::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& clusterer, int32_t firstHBF, int32_t tpcTimeBinCut) { GPUTPCCFDecodeZS::decode(clusterer, smem, nBlocks, nThreads, iBlock, iThread, firstHBF, tpcTimeBinCut); } -GPUdii() void GPUTPCCFDecodeZS::decode(GPUTPCClusterFinder& clusterer, GPUSharedMemory& s, int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t firstHBF, int32_t tpcTimeBinCut) +GPUdii() void GPUTPCCFDecodeZS::decode(GPUTPCClusterFinder& clusterer, GPUsharedref() GPUSharedMemory& s, int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t firstHBF, int32_t tpcTimeBinCut) { const uint32_t sector = clusterer.mISector; #ifdef GPUCA_GPUCODE @@ -80,7 +80,7 @@ GPUdii() void GPUTPCCFDecodeZS::decode(GPUTPCClusterFinder& clusterer, GPUShared for (uint32_t j = minJ; j < maxJ; j++) { #endif const uint32_t* pageSrc = (const uint32_t*)(((const uint8_t*)zs.zsPtr[endpoint][i]) + j * TPCZSHDR::TPC_ZS_PAGE_SIZE); - GPUCA_SHARED_CACHE_REF(&s.ZSPage[0], pageSrc, TPCZSHDR::TPC_ZS_PAGE_SIZE, uint32_t, pageCache); + GPUCA_SHARED_CACHE_REF(nThreads, iThread, &s.ZSPage[0], pageSrc, TPCZSHDR::TPC_ZS_PAGE_SIZE, uint32_t, pageCache); GPUbarrier(); const uint8_t* page = (const uint8_t*)pageCache; const o2::header::RAWDataHeader* rdh = (const o2::header::RAWDataHeader*)page; @@ -172,13 +172,13 @@ GPUdii() void GPUTPCCFDecodeZS::decode(GPUTPCClusterFinder& clusterer, GPUShared seqLen = rowData[(nSeq + 1) * 2] - rowData[nSeq * 2]; pad = rowData[nSeq++ * 2 + 1]; } - const CfFragment& fragment = clusterer.mPmemory->fragment; + const CfFragment& frag = clusterer.mPmemory->frag; TPCTime globalTime = timeBin + l; - bool discardTimeBin = not fragment.contains(globalTime); + bool discardTimeBin = not frag.contains(globalTime); discardTimeBin |= (tpcTimeBinCut > 0 && globalTime > tpcTimeBinCut); Row row = rowOffset + m; - CfChargePos pos(row, Pad(pad), discardTimeBin ? INVALID_TIME_BIN : fragment.toLocal(globalTime)); + CfChargePos pos(row, Pad(pad), discardTimeBin ? INVALID_TIME_BIN : frag.toLocal(globalTime)); positions[nDigitsTmp++] = pos; if (!discardTimeBin) { @@ -219,7 +219,7 @@ GPUdii() void GPUTPCCFDecodeZSLink::Thread<0>(int32_t nBlocks, int32_t nThreads, GPUd() size_t GPUTPCCFDecodeZSLink::DecodePage(GPUSharedMemory& smem, DecodeCtx& ctx) { - const CfFragment& fragment = ctx.clusterer.mPmemory->fragment; + const CfFragment& frag = ctx.clusterer.mPmemory->frag; const auto* rdHdr = ConsumeHeader(ctx.page); @@ -246,7 +246,7 @@ GPUd() size_t GPUTPCCFDecodeZSLink::DecodePage(GPUSharedMemory& smem, DecodeCtx& nDecoded += nAdc; - bool discardTimeBin = not fragment.contains(timeBin); + bool discardTimeBin = not frag.contains(timeBin); discardTimeBin |= (ctx.tpcTimeBinCut > 0 && timeBin > ctx.tpcTimeBinCut); if (discardTimeBin) { @@ -331,9 +331,9 @@ GPUd() void GPUTPCCFDecodeZSLink::DecodeTB( } o2::tpc::PadPos padAndRow = GetPadAndRowFromFEC(ctx.clusterer, cru, rawFECChannel, fecInPartition); - const CfFragment& fragment = ctx.clusterer.mPmemory->fragment; + const CfFragment& frag = ctx.clusterer.mPmemory->frag; float charge = ADCToFloat(adc, DECODE_MASK, DECODE_BITS_FACTOR); - WriteCharge(ctx.clusterer, charge, padAndRow, fragment.toLocal(timeBin), ctx.pageDigitOffset + myOffset); + WriteCharge(ctx.clusterer, charge, padAndRow, frag.toLocal(timeBin), ctx.pageDigitOffset + myOffset); } // for (uint8_t i = iThread; blockOffset < nAdc; i += NThreads) } @@ -393,7 +393,7 @@ GPUd() void GPUTPCCFDecodeZSLinkBase::Decode(int32_t nBlocks, int32_t nThreads, #endif const uint32_t* pageSrc = (const uint32_t*)(((const uint8_t*)zs.zsPtr[endpoint][i]) + j * TPCZSHDR::TPC_ZS_PAGE_SIZE); // Cache zs page in shared memory. Curiously this actually degrades performance... - // GPUCA_SHARED_CACHE_REF(&smem.ZSPage[0], pageSrc, TPCZSHDR::TPC_ZS_PAGE_SIZE, uint32_t, pageCache); + // GPUCA_SHARED_CACHE_REF(nThreads, iThread, &smem.ZSPage[0], pageSrc, TPCZSHDR::TPC_ZS_PAGE_SIZE, uint32_t, pageCache); // GPUbarrier(); // const uint8_t* page = (const uint8_t*)pageCache; const uint8_t* page = (const uint8_t*)pageSrc; @@ -651,7 +651,7 @@ GPUd() int16_t GPUTPCCFDecodeZSDenseLink::DecodeTB( constexpr int32_t NTHREADS = GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCCFDecodeZSDenseLink); static_assert(NTHREADS == GPUCA_WARP_SIZE, "Decoding TB Headers in parallel assumes block size is a single warp."); - const CfFragment& fragment = ctx.clusterer.mPmemory->fragment; + const CfFragment& frag = ctx.clusterer.mPmemory->frag; // Read timebin block header uint16_t tbbHdr = ConsumeByte(ctx.page); @@ -721,7 +721,7 @@ GPUd() int16_t GPUTPCCFDecodeZSDenseLink::DecodeTB( const uint8_t* adcData = ConsumeBytes(ctx.page, (nSamplesInTB * DECODE_BITS + 7) / 8); MAYBE_PAGE_OVERFLOW(ctx.page); - bool discardTimeBin = not fragment.contains(timeBin); + bool discardTimeBin = not frag.contains(timeBin); discardTimeBin |= (ctx.tpcTimeBinCut > 0 && timeBin > ctx.tpcTimeBinCut); if (discardTimeBin) { @@ -754,7 +754,7 @@ GPUd() int16_t GPUTPCCFDecodeZSDenseLink::DecodeTB( o2::tpc::PadPos padAndRow = GetPadAndRowFromFEC(ctx.clusterer, cru, rawFECChannelLink, smem.linkIds[iLink]); float charge = ADCToFloat(adc, DECODE_MASK, DECODE_BITS_FACTOR); - WriteCharge(ctx.clusterer, charge, padAndRow, fragment.toLocal(timeBin), ctx.pageDigitOffset + sample); + WriteCharge(ctx.clusterer, charge, padAndRow, frag.toLocal(timeBin), ctx.pageDigitOffset + sample); } // for (uint16_t sample = iThread; sample < nSamplesInTB; sample += NTHREADS) diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.h index 21d4ec0a28958..eed5e0d6f30ac 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.h @@ -45,7 +45,7 @@ class GPUTPCCFDecodeZS : public GPUKernelTemplate decodeZS, }; - static GPUd() void decode(GPUTPCClusterFinder& clusterer, GPUSharedMemory& s, int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t firstHBF, int32_t tpcTimeBinCut); + static GPUd() void decode(GPUTPCClusterFinder& clusterer, GPUsharedref() GPUSharedMemory& s, int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t firstHBF, int32_t tpcTimeBinCut); typedef GPUTPCClusterFinder processorType; GPUhdi() static processorType* Processor(GPUConstantMem& processors) @@ -59,7 +59,7 @@ class GPUTPCCFDecodeZS : public GPUKernelTemplate } template - GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, Args... args); + GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& clusterer, Args... args); }; class GPUTPCCFDecodeZSLinkBase : public GPUKernelTemplate diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDeconvolution.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDeconvolution.cxx index d6b8703a9b35d..b234094648eee 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDeconvolution.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDeconvolution.cxx @@ -26,7 +26,7 @@ GPUdii() void GPUTPCCFDeconvolution::Thread<0>(int32_t nBlocks, int32_t nThreads { CfArray2D chargeMap(reinterpret_cast(clusterer.mPchargeMap)); CfArray2D isPeakMap(clusterer.mPpeakMap); - GPUTPCCFDeconvolution::deconvolutionImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, isPeakMap, chargeMap, clusterer.mPpositions, clusterer.mPmemory->counters.nPositions, overwriteCharge); + GPUTPCCFDeconvolution::deconvolutionImpl(nBlocks, nThreads, iBlock, iThread, smem, isPeakMap, chargeMap, clusterer.mPpositions, clusterer.mPmemory->counters.nPositions, overwriteCharge); } GPUdii() void GPUTPCCFDeconvolution::deconvolutionImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, @@ -36,7 +36,7 @@ GPUdii() void GPUTPCCFDeconvolution::deconvolutionImpl(int32_t nBlocks, int32_t const uint32_t digitnum, uint8_t overwriteCharge) { - SizeT idx = get_global_id(0); + SizeT idx = (iBlock * nThreads + iThread); bool iamDummy = (idx >= digitnum); idx = iamDummy ? digitnum - 1 : idx; @@ -47,7 +47,7 @@ GPUdii() void GPUTPCCFDeconvolution::deconvolutionImpl(int32_t nBlocks, int32_t int8_t peakCount = (iamPeak) ? 1 : 0; - uint16_t ll = get_local_id(0); + uint16_t ll = iThread; uint16_t partId = ll; uint16_t in3x3 = 0; diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFMCLabelFlattener.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFMCLabelFlattener.cxx index 8b4f28f517782..a9e4fdbbca064 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFMCLabelFlattener.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFMCLabelFlattener.cxx @@ -46,7 +46,7 @@ template <> GPUd() void GPUTPCCFMCLabelFlattener::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory&, processorType& clusterer) { #if !defined(GPUCA_GPUCODE) - const Row row = get_global_id(0); + const Row row = (iBlock * nThreads + iThread); const size_t clusterInRow = clusterer.mPclusterInRow[row]; auto& labels = clusterer.mPlabelsByRow[row].data; @@ -66,7 +66,7 @@ template <> GPUd() void GPUTPCCFMCLabelFlattener::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory&, processorType& clusterer, GPUTPCLinearLabels* out) { #if !defined(GPUCA_GPUCODE) - uint32_t row = get_global_id(0); + uint32_t row = (iBlock * nThreads + iThread); uint32_t headerOffset = clusterer.mPlabelsHeaderGlobalOffset; uint32_t dataOffset = clusterer.mPlabelsDataGlobalOffset; diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx index 4dfa50d9439e4..336cfb6d801a9 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.cxx @@ -26,14 +26,14 @@ GPUdii() void GPUTPCCFNoiseSuppression::Thread chargeMap(reinterpret_cast(clusterer.mPchargeMap)); CfArray2D isPeakMap(clusterer.mPpeakMap); - noiseSuppressionImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, clusterer.Param().rec, chargeMap, isPeakMap, clusterer.mPpeakPositions, clusterer.mPmemory->counters.nPeaks, clusterer.mPisPeak); + noiseSuppressionImpl(nBlocks, nThreads, iBlock, iThread, smem, clusterer.Param().rec, chargeMap, isPeakMap, clusterer.mPpeakPositions, clusterer.mPmemory->counters.nPeaks, clusterer.mPisPeak); } template <> GPUdii() void GPUTPCCFNoiseSuppression::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer) { CfArray2D isPeakMap(clusterer.mPpeakMap); - updatePeaksImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer.mPpeakPositions, clusterer.mPisPeak, clusterer.mPmemory->counters.nPeaks, isPeakMap); + updatePeaksImpl(nBlocks, nThreads, iBlock, iThread, clusterer.mPpeakPositions, clusterer.mPisPeak, clusterer.mPmemory->counters.nPeaks, isPeakMap); } GPUdii() void GPUTPCCFNoiseSuppression::noiseSuppressionImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, @@ -44,7 +44,7 @@ GPUdii() void GPUTPCCFNoiseSuppression::noiseSuppressionImpl(int32_t nBlocks, in const uint32_t peaknum, uint8_t* isPeakPredicate) { - SizeT idx = get_global_id(0); + SizeT idx = (iBlock * nThreads + iThread); CfChargePos pos = peakPositions[CAMath::Min(idx, (SizeT)(peaknum - 1))]; Charge charge = chargeMap[pos].unpack(); @@ -54,6 +54,7 @@ GPUdii() void GPUTPCCFNoiseSuppression::noiseSuppressionImpl(int32_t nBlocks, in chargeMap, peakMap, calibration, + iThread, charge, pos, smem.posBcast, @@ -80,7 +81,7 @@ GPUd() void GPUTPCCFNoiseSuppression::updatePeaksImpl(int32_t nBlocks, int32_t n const uint32_t peakNum, CfArray2D& peakMap) { - SizeT idx = get_global_id(0); + SizeT idx = (iBlock * nThreads + iThread); if (idx >= peakNum) { return; @@ -167,6 +168,7 @@ GPUd() void GPUTPCCFNoiseSuppression::findMinimaAndPeaks( const CfArray2D& chargeMap, const CfArray2D& peakMap, const GPUSettingsRec& calibration, + uint16_t ll, float q, const CfChargePos& pos, CfChargePos* posBcast, @@ -175,8 +177,6 @@ GPUd() void GPUTPCCFNoiseSuppression::findMinimaAndPeaks( uint64_t* bigger, uint64_t* peaks) { - uint16_t ll = get_local_id(0); - posBcast[ll] = pos; GPUbarrier(); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.h index bdee75dc87732..8c251f310a1a8 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFNoiseSuppression.h @@ -69,7 +69,7 @@ class GPUTPCCFNoiseSuppression : public GPUKernelTemplate static GPUdi() bool keepPeak(uint64_t, uint64_t); - static GPUd() void findMinimaAndPeaks(const CfArray2D&, const CfArray2D&, const GPUSettingsRec&, float, const CfChargePos&, CfChargePos*, PackedCharge*, uint64_t*, uint64_t*, uint64_t*); + static GPUd() void findMinimaAndPeaks(const CfArray2D&, const CfArray2D&, const GPUSettingsRec&, uint16_t, float, const CfChargePos&, CfChargePos*, PackedCharge*, uint64_t*, uint64_t*, uint64_t*); }; } // namespace o2::gpu diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx index 7c93435f8bef8..ff71af8838b83 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx @@ -27,11 +27,12 @@ GPUdii() void GPUTPCCFPeakFinder::Thread<0>(int32_t nBlocks, int32_t nThreads, i { CfArray2D chargeMap(reinterpret_cast(clusterer.mPchargeMap)); CfArray2D isPeakMap(clusterer.mPpeakMap); - findPeaksImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, chargeMap, clusterer.mPpadIsNoisy, clusterer.mPpositions, clusterer.mPmemory->counters.nPositions, clusterer.Param().rec, *clusterer.GetConstantMem()->calibObjects.tpcPadGain, clusterer.mPisPeak, isPeakMap); + findPeaksImpl(nBlocks, nThreads, iBlock, iThread, smem, chargeMap, clusterer.mPpadIsNoisy, clusterer.mPpositions, clusterer.mPmemory->counters.nPositions, clusterer.Param().rec, *clusterer.GetConstantMem()->calibObjects.tpcPadGain, clusterer.mPisPeak, isPeakMap); } GPUdii() bool GPUTPCCFPeakFinder::isPeak( GPUSharedMemory& smem, + uint16_t ll, Charge q, const CfChargePos& pos, uint16_t N, @@ -40,8 +41,6 @@ GPUdii() bool GPUTPCCFPeakFinder::isPeak( CfChargePos* posBcast, PackedCharge* buf) { - uint16_t ll = get_local_id(0); - bool belowThreshold = (uint32_t)q <= calib.tpc.cfQMaxCutoff; uint16_t lookForPeaks; @@ -100,7 +99,7 @@ GPUd() void GPUTPCCFPeakFinder::findPeaksImpl(int32_t nBlocks, int32_t nThreads, uint8_t* isPeakPredicate, CfArray2D& peakMap) { - SizeT idx = get_global_id(0); + SizeT idx = (iBlock * nThreads + iThread); // For certain configurations dummy work items are added, so the total // number of work items is dividable by 64. @@ -111,7 +110,7 @@ GPUd() void GPUTPCCFPeakFinder::findPeaksImpl(int32_t nBlocks, int32_t nThreads, bool hasLostBaseline = pos.valid() ? padHasLostBaseline[pos.gpad] : true; charge = hasLostBaseline ? 0.f : charge; - uint8_t peak = isPeak(smem, charge, pos, SCRATCH_PAD_SEARCH_N, chargeMap, calib, smem.posBcast, smem.buf); + uint8_t peak = isPeak(smem, iThread, charge, pos, SCRATCH_PAD_SEARCH_N, chargeMap, calib, smem.posBcast, smem.buf); // Exit early if dummy. See comment above. bool iamDummy = (idx >= digitnum); diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h index 0d61378d3e6f2..8cf32c1589e61 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.h @@ -53,7 +53,7 @@ class GPUTPCCFPeakFinder : public GPUKernelTemplate private: static GPUd() void findPeaksImpl(int32_t, int32_t, int32_t, int32_t, GPUSharedMemory&, const CfArray2D&, const uint8_t*, const CfChargePos*, tpccf::SizeT, const GPUSettingsRec&, const TPCPadGainCalib&, uint8_t*, CfArray2D&); - static GPUd() bool isPeak(GPUSharedMemory&, tpccf::Charge, const CfChargePos&, uint16_t, const CfArray2D&, const GPUSettingsRec&, CfChargePos*, PackedCharge*); + static GPUd() bool isPeak(GPUSharedMemory&, uint16_t, tpccf::Charge, const CfChargePos&, uint16_t, const CfArray2D&, const GPUSettingsRec&, CfChargePos*, PackedCharge*); }; } // namespace o2::gpu diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFStreamCompaction.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFStreamCompaction.cxx index 0f2fd235dc0d0..4622638674d9a 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFStreamCompaction.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFStreamCompaction.cxx @@ -30,7 +30,7 @@ GPUdii() void GPUTPCCFStreamCompaction::Thread auto* scanOffset = clusterer.GetScanBuffer(iBuf - 1); auto* scanOffsetNext = clusterer.GetScanBuffer(iBuf); - int32_t iThreadGlobal = get_global_id(0); + int32_t iThreadGlobal = (iBlock * nThreads + iThread); int32_t offsetInBlock = work_group_scan_inclusive_add((iThreadGlobal < nElems) ? scanOffset[iThreadGlobal] : 0); if (iThreadGlobal < nElems) { @@ -70,7 +70,7 @@ template <> GPUdii() void GPUTPCCFStreamCompaction::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, int32_t iBuf, int32_t nElems) { #ifdef GPUCA_GPUCODE - int32_t iThreadGlobal = get_global_id(0); + int32_t iThreadGlobal = (iBlock * nThreads + iThread); int32_t* scanOffset = clusterer.GetScanBuffer(iBuf - 1); bool inBounds = (iThreadGlobal < nElems); @@ -87,7 +87,7 @@ template <> GPUdii() void GPUTPCCFStreamCompaction::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& /*smem*/, processorType& clusterer, int32_t iBuf, uint32_t offset, int32_t nElems) { #ifdef GPUCA_GPUCODE - int32_t iThreadGlobal = get_global_id(0) + offset; + int32_t iThreadGlobal = (iBlock * nThreads + iThread) + offset; int32_t* scanOffsetPrev = clusterer.GetScanBuffer(iBuf - 1); const int32_t* scanOffset = clusterer.GetScanBuffer(iBuf); @@ -107,7 +107,7 @@ GPUdii() void GPUTPCCFStreamCompaction::Thread bufferSize) { diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h index d169440a8d972..4a790750c773b 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h @@ -63,7 +63,7 @@ class GPUTPCClusterFinder : public GPUProcessor uint32_t maxTimeBin = 0; uint32_t nPagesSubsector = 0; } counters; - CfFragment fragment; + CfFragment frag; }; struct ZSOffset { diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinderDump.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinderDump.cxx index 3b06db8efc1a3..d778dce37c4af 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinderDump.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinderDump.cxx @@ -25,7 +25,7 @@ void GPUTPCClusterFinder::DumpDigits(std::ostream& out) { const auto nPositions = mPmemory->counters.nPositions; - out << "\nClusterer - Digits - Sector " << mISector << " - Fragment " << mPmemory->fragment.index << ": " << nPositions << "\n"; + out << "\nClusterer - Digits - Sector " << mISector << " - Fragment " << mPmemory->frag.index << ": " << nPositions << "\n"; out << std::hex; for (size_t i = 0; i < mPmemory->counters.nPositions; i++) { @@ -37,7 +37,7 @@ void GPUTPCClusterFinder::DumpDigits(std::ostream& out) void GPUTPCClusterFinder::DumpChargeMap(std::ostream& out, std::string_view title) { - out << "\nClusterer - " << title << " - Sector " << mISector << " - Fragment " << mPmemory->fragment.index << "\n"; + out << "\nClusterer - " << title << " - Sector " << mISector << " - Fragment " << mPmemory->frag.index << "\n"; CfArray2D map(mPchargeMap); out << std::hex; @@ -70,7 +70,7 @@ void GPUTPCClusterFinder::DumpChargeMap(std::ostream& out, std::string_view titl void GPUTPCClusterFinder::DumpPeakMap(std::ostream& out, std::string_view title) { - out << "\nClusterer - " << title << " - Sector " << mISector << " - Fragment " << mPmemory->fragment.index << "\n"; + out << "\nClusterer - " << title << " - Sector " << mISector << " - Fragment " << mPmemory->frag.index << "\n"; CfArray2D map(mPpeakMap); @@ -106,7 +106,7 @@ void GPUTPCClusterFinder::DumpPeakMap(std::ostream& out, std::string_view title) void GPUTPCClusterFinder::DumpPeaks(std::ostream& out) { - out << "\nClusterer - Peaks - Sector " << mISector << " - Fragment " << mPmemory->fragment.index << "\n"; + out << "\nClusterer - Peaks - Sector " << mISector << " - Fragment " << mPmemory->frag.index << "\n"; for (uint32_t i = 0; i < mPmemory->counters.nPositions; i++) { out << int32_t{mPisPeak[i]}; if ((i + 1) % 100 == 0) { @@ -119,7 +119,7 @@ void GPUTPCClusterFinder::DumpPeaksCompacted(std::ostream& out) { const auto nPeaks = mPmemory->counters.nPeaks; - out << "\nClusterer - Compacted Peaks - Sector " << mISector << " - Fragment " << mPmemory->fragment.index << ": " << nPeaks << "\n"; + out << "\nClusterer - Compacted Peaks - Sector " << mISector << " - Fragment " << mPmemory->frag.index << ": " << nPeaks << "\n"; for (size_t i = 0; i < nPeaks; i++) { const auto& pos = mPpeakPositions[i]; out << pos.time() << " " << int32_t{pos.pad()} << " " << int32_t{pos.row()} << "\n"; @@ -128,10 +128,10 @@ void GPUTPCClusterFinder::DumpPeaksCompacted(std::ostream& out) void GPUTPCClusterFinder::DumpSuppressedPeaks(std::ostream& out) { - const auto& fragment = mPmemory->fragment; + const auto& frag = mPmemory->frag; const auto nPeaks = mPmemory->counters.nPeaks; - out << "\nClusterer - NoiseSuppression - Sector " << mISector << " - Fragment " << fragment.index << mISector << "\n"; + out << "\nClusterer - NoiseSuppression - Sector " << mISector << " - Fragment " << frag.index << mISector << "\n"; for (uint32_t i = 0; i < nPeaks; i++) { out << int32_t{mPisPeak[i]}; if ((i + 1) % 100 == 0) { @@ -142,10 +142,10 @@ void GPUTPCClusterFinder::DumpSuppressedPeaks(std::ostream& out) void GPUTPCClusterFinder::DumpSuppressedPeaksCompacted(std::ostream& out) { - const auto& fragment = mPmemory->fragment; + const auto& frag = mPmemory->frag; const auto nPeaks = mPmemory->counters.nClusters; - out << "\nClusterer - Noise Suppression Peaks Compacted - Sector " << mISector << " - Fragment " << fragment.index << ": " << nPeaks << "\n"; + out << "\nClusterer - Noise Suppression Peaks Compacted - Sector " << mISector << " - Fragment " << frag.index << ": " << nPeaks << "\n"; for (size_t i = 0; i < nPeaks; i++) { const auto& peak = mPfilteredPeakPositions[i]; out << peak.time() << " " << int32_t{peak.pad()} << " " << int32_t{peak.row()} << "\n"; diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx index 902c04e6e92f3..4568ffa610437 100644 --- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx +++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx @@ -53,14 +53,14 @@ static_assert(GPUTPCNNClusterizerKernels::SCRATCH_PAD_WORK_GROUP_SIZE == GPUTPCC template <> GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint32_t batchStart) { - uint32_t glo_idx = get_global_id(0); + uint32_t glo_idx = (iBlock * nThreads + iThread); auto& clusterer = processors.tpcClusterer[sector]; auto& clustererNN = processors.tpcNNClusterer[sector]; CfArray2D chargeMap(reinterpret_cast(clusterer.mPchargeMap)); CPU_ONLY(MCLabelAccumulator labelAcc(clusterer)); tpc::ClusterNative* clusterOut = clusterer.mPclusterByRow; int8_t isAccepted = (clustererNN.mNnClusterizerUseClassification ? (clustererNN.mOutputDataClass[CAMath::Min(glo_idx, (uint32_t)clusterer.mPmemory->counters.nClusters - 1)] > 0) : 1); - GPUTPCCFClusterizer::computeClustersImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->fragment, smem, chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow, isAccepted); + GPUTPCCFClusterizer::computeClustersImpl(nBlocks, nThreads, iBlock, iThread, clusterer, clusterer.mPmemory->frag, smem, chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow, isAccepted); } template <> @@ -69,7 +69,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread= clusterer.mPmemory->counters.nClusters || glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { return; } @@ -158,7 +158,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint32_t batchStart) { - uint32_t glo_idx = get_global_id(0); + uint32_t glo_idx = (iBlock * nThreads + iThread); auto& clusterer = processors.tpcClusterer[sector]; auto& clustererNN = processors.tpcNNClusterer[sector]; if (glo_idx + batchStart >= clusterer.mPmemory->counters.nClusters || glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { @@ -274,7 +274,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint32_t batchStart) { - uint32_t glo_idx = get_global_id(0); + uint32_t glo_idx = (iBlock * nThreads + iThread); auto& clusterer = processors.tpcClusterer[sector]; auto& clustererNN = processors.tpcNNClusterer[sector]; if (glo_idx + batchStart >= clusterer.mPmemory->counters.nClusters || glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { @@ -313,7 +313,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint32_t batchStart) { - uint32_t glo_idx = get_global_id(0); + uint32_t glo_idx = (iBlock * nThreads + iThread); auto& clusterer = processors.tpcClusterer[sector]; auto& clustererNN = processors.tpcNNClusterer[sector]; if (glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { @@ -338,6 +338,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadcollect(peak, central_charge)); GPUTPCCFClusterizer::buildCluster( clusterer.Param().rec, + iThread, chargeMap, peak, smem.posBcast, @@ -358,6 +359,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadcollect(peak, central_charge)); GPUTPCCFClusterizer::buildCluster( clusterer.Param().rec, + iThread, chargeMap, peak, smem.posBcast, @@ -366,7 +368,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).isOverlap(peak.time())) { + if ((clusterer.mPmemory->frag).isOverlap(peak.time())) { if (clusterer.mPclusterPosInRow) { clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; } @@ -390,7 +392,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).start + publishTimePosition, + (clusterer.mPmemory->frag).start + publishTimePosition, notSingleTime ? clustererNN.mOutputDataReg1_32[model_output_index + 3] : 0.f, clustererNN.mClusterFlags[2 * glo_idx], clustererNN.mClusterFlags[2 * glo_idx + 1]); @@ -401,7 +403,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).start + publishTimePosition, + (clusterer.mPmemory->frag).start + publishTimePosition, notSingleTime ? clustererNN.mOutputDataReg1_16[model_output_index + 3].ToFloat() : 0.f, clustererNN.mClusterFlags[2 * glo_idx], clustererNN.mClusterFlags[2 * glo_idx + 1]); @@ -509,7 +511,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint32_t batchStart) { - uint32_t glo_idx = get_global_id(0); + uint32_t glo_idx = (iBlock * nThreads + iThread); auto& clusterer = processors.tpcClusterer[sector]; auto& clustererNN = processors.tpcNNClusterer[sector]; if (glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { @@ -532,6 +534,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadcollect(peak, central_charge)); GPUTPCCFClusterizer::buildCluster( clusterer.Param().rec, + iThread, chargeMap, peak, smem.posBcast, @@ -552,6 +555,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadcollect(peak, central_charge)); GPUTPCCFClusterizer::buildCluster( clusterer.Param().rec, + iThread, chargeMap, peak, smem.posBcast, @@ -560,7 +564,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).isOverlap(peak.time())) { + if ((clusterer.mPmemory->frag).isOverlap(peak.time())) { if (clusterer.mPclusterPosInRow) { clusterer.mPclusterPosInRow[full_glo_idx] = clusterer.mNMaxClusterPerRow; } @@ -576,7 +580,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).start + publishTimePosition, + (clusterer.mPmemory->frag).start + publishTimePosition, clustererNN.mOutputDataReg2_32[model_output_index + 6], clustererNN.mClusterFlags[2 * glo_idx], clustererNN.mClusterFlags[2 * glo_idx + 1]); @@ -587,7 +591,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).start + publishTimePosition, + (clusterer.mPmemory->frag).start + publishTimePosition, clustererNN.mOutputDataReg2_16[model_output_index + 6].ToFloat(), clustererNN.mClusterFlags[2 * glo_idx], clustererNN.mClusterFlags[2 * glo_idx + 1]); @@ -630,7 +634,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).start + publishTimePosition, + (clusterer.mPmemory->frag).start + publishTimePosition, clustererNN.mOutputDataReg2_32[model_output_index + 7], clustererNN.mClusterFlags[2 * glo_idx], clustererNN.mClusterFlags[2 * glo_idx + 1]); @@ -641,7 +645,7 @@ GPUdii() void GPUTPCNNClusterizerKernels::Threadfragment).start + publishTimePosition, + (clusterer.mPmemory->frag).start + publishTimePosition, clustererNN.mOutputDataReg2_16[model_output_index + 7].ToFloat(), clustererNN.mClusterFlags[2 * glo_idx], clustererNN.mClusterFlags[2 * glo_idx + 1]); @@ -680,7 +684,7 @@ template <> GPUdii() void GPUTPCNNClusterizerKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& processors, uint8_t sector, int8_t dtype, int8_t withMC, uint batchStart) { // Implements identical publishing logic as the heuristic clusterizer and deconvolution kernel - uint32_t glo_idx = get_global_id(0); + uint32_t glo_idx = (iBlock * nThreads + iThread); auto& clusterer = processors.tpcClusterer[sector]; auto& clustererNN = processors.tpcNNClusterer[sector]; if (glo_idx + batchStart >= clusterer.mPmemory->counters.nClusters || glo_idx >= (uint32_t)clustererNN.mNnClusterizerBatchedMode) { diff --git a/GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.cxx b/GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.cxx index dea4cdbca430e..6632be71cbb75 100644 --- a/GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.cxx +++ b/GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.cxx @@ -35,8 +35,8 @@ GPUdii() void GPUTRDTrackerKernels::Thread(int32_t nBlocks, int32_t nThreads, in } } #endif - GPUCA_TBB_KERNEL_LOOP(trdTracker->GetRec(), int32_t, i, trdTracker->NTracks(), { - trdTracker->DoTrackingThread(i, get_global_id(0)); + GPUCA_TBB_KERNEL_LOOP(trdTracker->GetRec(), nBlocks, nThreads, iBlock, iThread, int32_t, i, trdTracker->NTracks(), { + trdTracker->DoTrackingThread(i, (iBlock * nThreads + iThread)); }); } diff --git a/GPU/GPUTracking/cmake/kernel_helpers.cmake b/GPU/GPUTracking/cmake/kernel_helpers.cmake index cc50d28ecef9e..459165d86cf5d 100644 --- a/GPU/GPUTracking/cmake/kernel_helpers.cmake +++ b/GPU/GPUTracking/cmake/kernel_helpers.cmake @@ -55,11 +55,13 @@ function(o2_gpu_add_kernel kernel_name kernel_files) math(EXPR n "${n} - 1") foreach(i RANGE 3 ${n} 2) math(EXPR j "${i} + 1") + # buffer indices 0, 1 and 2 are gpu_mem, the constant memory and the sector + math(EXPR TMP_ARG_IDX "3 + (${i} - 3) / 2") if(${ARGV${i}} MATCHES "\\*$") - string(APPEND OPT1 ",GPUPtr1(${ARGV${i}},${ARGV${j}})") + string(APPEND OPT1 ",GPUPtr1(${TMP_ARG_IDX},${ARGV${i}},${ARGV${j}})") string(APPEND OPT2 ",GPUPtr2(${ARGV${i}},${ARGV${j}})") else() - string(APPEND OPT1 ",${ARGV${i}} ${ARGV${j}}") + string(APPEND OPT1 ",GPUArg1(${TMP_ARG_IDX},${ARGV${i}},${ARGV${j}})") string(APPEND OPT2 ",${ARGV${j}}") endif() string(APPEND OPT3 ",${ARGV${i}}") diff --git a/GPU/Utils/GPUCommonBitSet.h b/GPU/Utils/GPUCommonBitSet.h index 302334e01e29d..eebe00b2bcdcb 100644 --- a/GPU/Utils/GPUCommonBitSet.h +++ b/GPU/Utils/GPUCommonBitSet.h @@ -37,11 +37,21 @@ class bitset public: GPUdDefault() constexpr bitset() = default; +#ifndef __METAL__ + // MSL will not use a user-declared copy constructor to build an object in the + // constant address space, where GPUconstexpr() arrays of bitset live. Leaving + // it implicit is what makes those arrays constructible. GPUdDefault() constexpr bitset(const bitset&) = default; +#endif #ifdef __OPENCL__ GPUdDefault() constexpr bitset(const __constant bitset&) = default; #endif // __OPENCL__ GPUd() constexpr bitset(uint32_t vv) : v(vv) {}; +#ifdef __METAL__ + // Objects in the constant address space are built and read through their own + // overloads: a generic `this` does not reach constant memory in MSL. + GPUd() constexpr bitset(uint32_t vv) constant : v(vv) {}; +#endif static GPUglobalconstexpr() uint32_t full_set = ((1ul << N) - 1ul); GPUd() constexpr bool all() const { return (v & full_set) == full_set; } @@ -78,6 +88,13 @@ class bitset GPUd() constexpr bool operator!=(const bitset b) const { return v != b.v; } GPUd() constexpr bool operator[](uint32_t i) const { return (v >> i) & 1u; } +#ifdef __METAL__ + GPUd() constexpr bitset operator|(const bitset b) constant { return v | b.v; } + GPUd() constexpr bitset operator&(const bitset b) constant { return v & b.v; } + GPUd() constexpr bool operator[](uint32_t i) constant { return (v >> i) & 1u; } + GPUd() constexpr bool any() constant { return v & full_set; } + GPUd() constexpr uint32_t to_ulong() constant { return v; } +#endif GPUd() constexpr uint32_t to_ulong() const { return v; } diff --git a/dependencies/FindO2GPU.cmake b/dependencies/FindO2GPU.cmake index d2f426c448e12..801bcd17b691d 100644 --- a/dependencies/FindO2GPU.cmake +++ b/dependencies/FindO2GPU.cmake @@ -10,7 +10,7 @@ # or submit itself to any jurisdiction. # NOTE!!!! - Whenever this file is changed, move it over to alidist/resources -# FindO2GPU.cmake Version 19 +# FindO2GPU.cmake Version 21 set(CUDA_COMPUTETARGET_DEFAULT_FULL 80-real;86-real;89-real;120-real;75-virtual) set(HIP_AMDGPUTARGET_DEFAULT_FULL gfx906;gfx908) @@ -20,7 +20,13 @@ set(HIP_AMDGPUTARGET_DEFAULT_MINIMAL gfx906) if(NOT DEFINED ENABLE_CUDA) set(ENABLE_CUDA "AUTO") endif() -if(NOT DEFINED ENABLE_OPENCL) +if(APPLE) + # macOS ships OpenCL 1.2 only, below the 2.x the OpenCL backend needs; Metal + # replaces it there. OFF rather than AUTO while the backend is unproven. + if(NOT DEFINED ENABLE_METAL) + set(ENABLE_METAL "OFF") + endif() +elseif(NOT DEFINED ENABLE_OPENCL) set(ENABLE_OPENCL "AUTO") endif() if(NOT DEFINED ENABLE_HIP) @@ -28,6 +34,9 @@ if(NOT DEFINED ENABLE_HIP) endif() string(TOUPPER "${ENABLE_CUDA}" ENABLE_CUDA) string(TOUPPER "${ENABLE_OPENCL}" ENABLE_OPENCL) +if(APPLE) + string(TOUPPER "${ENABLE_METAL}" ENABLE_METAL) +endif() string(TOUPPER "${ENABLE_HIP}" ENABLE_HIP) if(NOT DEFINED CMAKE_BUILD_TYPE_UPPER) string(TOUPPER "${CMAKE_BUILD_TYPE}" CMAKE_BUILD_TYPE_UPPER) @@ -112,6 +121,9 @@ function(detect_gpu_arch backend) # Detect GPU architecture, optionally filterri if(OPENCL_ENABLED OR backend STREQUAL "ALL") list(APPEND TARGET_ARCH "OPENCL") endif() + if(METAL_ENABLED OR backend STREQUAL "ALL") + list(APPEND TARGET_ARCH "METAL") + endif() set(TARGET_ARCH "${TARGET_ARCH}" PARENT_SCOPE) else() message(FATAL_ERROR "Unknown backend provided: ${backend}") @@ -149,6 +161,7 @@ if(GPUCA_DETERMINISTIC_NO_FTZ) set(GPUCA_CXX_DENORMALS_FLAGS "") set(GPUCA_CUDA_DENORMALS_FLAGS "--ftz=false") set(GPUCA_OCL_DENORMALS_FLAGS "") + set(GPUCA_METAL_DENORMALS_FLAGS "-fno-fast-math") set(GPUCA_HIP_DENORMALS_FLAGS "-fno-gpu-flush-denormals-to-zero") else() if (CMAKE_SYSTEM_NAME MATCHES Darwin OR NOT CMAKE_SYSTEM_PROCESSOR MATCHES "(x86)|(X86)|(amd64)|(AMD64)") @@ -158,11 +171,13 @@ else() endif() set(GPUCA_CUDA_DENORMALS_FLAGS "--ftz=true") set(GPUCA_OCL_DENORMALS_FLAGS "-cl-denorms-are-zero") + set(GPUCA_METAL_DENORMALS_FLAGS "-ffast-math") set(GPUCA_HIP_DENORMALS_FLAGS "-fgpu-flush-denormals-to-zero") endif() set(GPUCA_CXX_NO_FAST_MATH_FLAGS "-fno-fast-math -ffp-contract=off") set(GPUCA_CUDA_NO_FAST_MATH_FLAGS "--prec-div=true --prec-sqrt=true --fmad false -Xcompiler -fno-fast-math -Xcompiler -ffp-contract=off") set(GPUCA_OCL_NO_FAST_MATH_FLAGS -cl-fp32-correctly-rounded-divide-sqrt ) +set(GPUCA_METAL_NO_FAST_MATH_FLAGS "-fno-fast-math") if(GPUCA_DETERMINISTIC_MODE GREATER_EQUAL ${GPUCA_DETERMINISTIC_MODE_MAP_WHOLEO2}) add_definitions(-DGPUCA_DETERMINISTIC_MODE) string(APPEND CMAKE_CXX_FLAGS_${CMAKE_BUILD_TYPE_UPPER} " ${GPUCA_CXX_NO_FAST_MATH_FLAGS}") @@ -429,6 +444,41 @@ if(ENABLE_HIP) endif() endif() +# =================================== Metal ================================== +if(ENABLE_METAL) + find_library(METAL_FRAMEWORK Metal) + find_library(COREFOUNDATION_FRAMEWORK CoreFoundation) + find_library(FOUNDATION_FRAMEWORK Foundation) + find_library(QUARTZCORE_FRAMEWORK QuartzCore) + # The frameworks are there on every macOS, but the backend needs MSL 4.1, which + # is macOS 27 and its toolchain, so ask the compiler instead of assuming. A + # library built as MSL 4.1 also only loads on macOS 27 and later. + if(NOT DEFINED GPUCA_METAL_MSL41) + set(GPUCA_METAL_PROBE "${CMAKE_CURRENT_BINARY_DIR}/metal_msl41_probe.metal") + file(WRITE "${GPUCA_METAL_PROBE}" "#include \nkernel void probe(device float* o [[buffer(0)]], uint i [[thread_position_in_grid]]) { o[i] = o[i] * 2.0f; }\n") + execute_process(COMMAND xcrun -sdk macosx metal -std=metal4.1 -c "${GPUCA_METAL_PROBE}" -o "${GPUCA_METAL_PROBE}.air" + RESULT_VARIABLE GPUCA_METAL_PROBE_RESULT OUTPUT_QUIET ERROR_QUIET) + if(GPUCA_METAL_PROBE_RESULT EQUAL 0) + set(GPUCA_METAL_MSL41 ON CACHE INTERNAL "Metal toolchain compiles MSL 4.1") + else() + set(GPUCA_METAL_MSL41 OFF CACHE INTERNAL "Metal toolchain compiles MSL 4.1") + endif() + endif() + if(METAL_FRAMEWORK AND COREFOUNDATION_FRAMEWORK AND FOUNDATION_FRAMEWORK AND QUARTZCORE_FRAMEWORK AND GPUCA_METAL_MSL41) + set(METAL_ENABLED ON) + set(METAL_FRAMEWORKS ${METAL_FRAMEWORK} ${COREFOUNDATION_FRAMEWORK} + ${FOUNDATION_FRAMEWORK} ${QUARTZCORE_FRAMEWORK}) + message(STATUS "Found Metal frameworks") + elseif(NOT ENABLE_METAL STREQUAL "AUTO") + if(NOT GPUCA_METAL_MSL41) + message(FATAL_ERROR "The Metal backend needs MSL 4.1: this toolchain rejects -std=metal4.1, and the result would need macOS 27 to run") + endif() + message(FATAL_ERROR "Metal frameworks not available") + else() + set(METAL_ENABLED OFF) + endif() +endif() + # if we end up here without a FATAL, it means we have found the "O2GPU" package set(O2GPU_FOUND TRUE) if (NOT GPUCA_FINDO2GPU_CHECK_ONLY)