Skip to content

Commit 8f6726b

Browse files
davidrohrdavidrohr
committed
GPU: Change GPUCA_DETERMINISTIC_MODE define to GPUCA_DETERMINISTIC_CODE macro, that can be used also in RTC
1 parent 3e56e55 commit 8f6726b

File tree

10 files changed

+116
-150
lines changed

10 files changed

+116
-150
lines changed

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu‎

Lines changed: 2 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -43,13 +43,8 @@
4343
#define THRUST_NAMESPACE thrust::hip
4444
#endif
4545

46-
#ifdef GPUCA_DETERMINISTIC_MODE
47-
#define GPU_BLOCKS 1
48-
#define GPU_THREADS 1
49-
#else
50-
#define GPU_BLOCKS 99999
51-
#define GPU_THREADS 99999
52-
#endif
46+
#define GPU_BLOCKS GPUCA_DETERMINISTIC_CODE(1, 99999)
47+
#define GPU_THREADS GPUCA_DETERMINISTIC_CODE(1, 99999)
5348

5449
// O2 track model
5550
#include "ReconstructionDataFormats/Track.h"

‎GPU/Common/GPUCommonDef.h‎

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -68,10 +68,18 @@
6868
#define GPUCA_DEBUG_STREAMER_CHECK(...)
6969
#endif
7070

71-
#ifndef GPUCA_RTC_SPECIAL_CODE
71+
#ifndef GPUCA_RTC_SPECIAL_CODE // By default, we ignore special RTC code
7272
#define GPUCA_RTC_SPECIAL_CODE(...)
7373
#endif
7474

75+
#ifndef GPUCA_DETERMINISTIC_CODE
76+
#ifdef GPUCA_DETERMINISTIC_MODE
77+
#define GPUCA_DETERMINISTIC_CODE(det, indet) det // In deterministic mode, take deterministic code path
78+
#else
79+
#define GPUCA_DETERMINISTIC_CODE(det, indet) indet // otherwise the fast default code path
80+
#endif
81+
#endif
82+
7583
// API Definitions for GPU Compilation
7684
#include "GPUCommonDefAPI.h"
7785

‎GPU/Common/GPUCommonMath.h‎

Lines changed: 30 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -248,7 +248,7 @@ GPUdi() uint32_t GPUCommonMath::Float2UIntReint(const float& x)
248248
#endif
249249
}
250250

251-
#ifdef GPUCA_DETERMINISTIC_MODE
251+
GPUCA_DETERMINISTIC_CODE( // clang-format off
252252
GPUdi() constexpr float GPUCommonMath::Round(float x) { return GPUCA_CHOICE(roundf(x), roundf(x), round(x)); }
253253
GPUdi() constexpr int32_t GPUCommonMath::Float2IntRn(float x) { return (int32_t)Round(x); }
254254
GPUhdi() constexpr float GPUCommonMath::Sqrt(float x) { return GPUCA_CHOICE(sqrtf(x), (float)sqrt((double)x), sqrt(x)); }
@@ -264,7 +264,7 @@ GPUdi() constexpr float GPUCommonMath::Log(float x) { return GPUCA_CHOICE((float
264264
GPUdi() constexpr float GPUCommonMath::Exp(float x) { return GPUCA_CHOICE((float)exp((double)x), (float)exp((double)x), exp(x)); }
265265
GPUdi() constexpr bool GPUCommonMath::Finite(float x) { return GPUCA_CHOICE(std::isfinite(x), isfinite(x), isfinite(x)); }
266266
GPUdi() constexpr bool GPUCommonMath::IsNaN(float x) { return GPUCA_CHOICE(std::isnan(x), isnan(x), isnan(x)); }
267-
#else
267+
, // !GPUCA_DETERMINISTIC_CODE
268268
GPUdi() constexpr float GPUCommonMath::Round(float x) { return GPUCA_CHOICE(roundf(x), rintf(x), rint(x)); }
269269
GPUdi() constexpr int32_t GPUCommonMath::Float2IntRn(float x) { return GPUCA_CHOICE((int32_t)Round(x), __float2int_rn(x), (int32_t)Round(x)); }
270270
GPUhdi() constexpr float GPUCommonMath::Sqrt(float x) { return GPUCA_CHOICE(sqrtf(x), sqrtf(x), sqrt(x)); }
@@ -280,20 +280,22 @@ GPUdi() constexpr float GPUCommonMath::Log(float x) { return GPUCA_CHOICE(logf(x
280280
GPUdi() constexpr float GPUCommonMath::Exp(float x) { return GPUCA_CHOICE(expf(x), expf(x), exp(x)); }
281281
GPUdi() constexpr bool GPUCommonMath::Finite(float x) { return true; }
282282
GPUdi() constexpr bool GPUCommonMath::IsNaN(float x) { return false; }
283-
#endif
283+
) // clang-format on
284284

285285
GPUhdi() void GPUCommonMath::SinCos(float x, float& s, float& c)
286286
{
287-
#if defined(GPUCA_DETERMINISTIC_MODE) && !defined(__OPENCL__)
288-
s = sin((double)x);
289-
c = cos((double)x);
290-
#elif !defined(GPUCA_GPUCODE_DEVICE) && defined(__APPLE__)
291-
__sincosf(x, &s, &c);
287+
GPUCA_DETERMINISTIC_CODE( // clang-format off
288+
s = sin((double)x);
289+
c = cos((double)x);
290+
, // !GPUCA_DETERMINISTIC_CODE
291+
#if !defined(GPUCA_GPUCODE_DEVICE) && defined(__APPLE__)
292+
__sincosf(x, &s, &c);
292293
#elif !defined(GPUCA_GPUCODE_DEVICE) && (defined(__GNU_SOURCE__) || defined(_GNU_SOURCE) || defined(GPUCA_GPUCODE))
293-
sincosf(x, &s, &c);
294+
sincosf(x, &s, &c);
294295
#else
295-
GPUCA_CHOICE((void)((s = sinf(x)) + (c = cosf(x))), sincosf(x, &s, &c), s = sincos(x, &c));
296+
GPUCA_CHOICE((void)((s = sinf(x)) + (c = cosf(x))), sincosf(x, &s, &c), s = sincos(x, &c));
296297
#endif
298+
) // clang-format on
297299
}
298300

299301
GPUhdi() void GPUCommonMath::SinCosd(double x, double& s, double& c)
@@ -390,22 +392,26 @@ GPUdi() T GPUCommonMath::MaxWithRef(T x, T y, T z, T w, S refX, S refY, S refZ,
390392

391393
GPUdi() float GPUCommonMath::InvSqrt(float _x)
392394
{
393-
#if defined(GPUCA_DETERMINISTIC_MODE) || defined(__OPENCL__)
394-
return 1.f / Sqrt(_x);
395-
#elif defined(__CUDACC__) || defined(__HIPCC__)
396-
return __frsqrt_rn(_x);
397-
#elif defined(__FAST_MATH__)
398-
return 1.f / sqrtf(_x);
395+
GPUCA_DETERMINISTIC_CODE( // clang-format off
396+
return 1.f / Sqrt(_x);
397+
, // !GPUCA_DETERMINISTIC_CODE
398+
#if defined(__CUDACC__) || defined(__HIPCC__)
399+
return __frsqrt_rn(_x);
400+
#elif defined(__OPENCL__) && defined(__clang__)
401+
return 1.f / sqrt(_x);
402+
#elif !defined(__OPENCL__) && (defined(__FAST_MATH__) || defined(__clang__))
403+
return 1.f / sqrtf(_x);
399404
#else
400-
union {
401-
float f;
402-
int32_t i;
403-
} x = {_x};
404-
const float xhalf = 0.5f * x.f;
405-
x.i = 0x5f3759df - (x.i >> 1);
406-
x.f = x.f * (1.5f - xhalf * x.f * x.f);
407-
return x.f;
405+
union {
406+
float f;
407+
int32_t i;
408+
} x = {_x};
409+
const float xhalf = 0.5f * x.f;
410+
x.i = 0x5f3759df - (x.i >> 1);
411+
x.f = x.f * (1.5f - xhalf * x.f * x.f);
412+
return x.f;
408413
#endif
414+
) // clang-format on
409415
}
410416

411417
template <>

‎GPU/GPUTracking/Definitions/GPUDefGPUParameters.h‎

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,7 @@
2525
#error Please include GPUDef.h
2626
#endif
2727

28+
#include "GPUCommonDef.h"
2829
#include "GPUDefMacros.h"
2930

3031
// GPU Run Configuration
@@ -566,12 +567,8 @@
566567
#ifndef GPUCA_MERGER_INTERPOLATION_ERROR_TYPE
567568
#define GPUCA_MERGER_INTERPOLATION_ERROR_TYPE float
568569
#endif
569-
#ifdef GPUCA_DETERMINISTIC_MODE
570-
#undef GPUCA_MERGER_INTERPOLATION_ERROR_TYPE
571-
#define GPUCA_MERGER_INTERPOLATION_ERROR_TYPE float
572-
#undef GPUCA_DEDX_STORAGE_TYPE
573-
#define GPUCA_DEDX_STORAGE_TYPE float
574-
#endif
570+
#define GPUCA_MERGER_INTERPOLATION_ERROR_TYPE_A GPUCA_DETERMINISTIC_CODE(float, GPUCA_MERGER_INTERPOLATION_ERROR_TYPE)
571+
#define GPUCA_DEDX_STORAGE_TYPE_A GPUCA_DETERMINISTIC_CODE(float, GPUCA_DEDX_STORAGE_TYPE)
575572

576573
#ifndef GPUCA_WARP_SIZE
577574
#ifdef GPUCA_GPUCODE

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx‎

Lines changed: 60 additions & 100 deletions
Original file line numberDiff line numberDiff line change
@@ -723,17 +723,9 @@ GPUd() void GPUTPCGMMerger::MergeBorderTracks<1>(int32_t nBlocks, int32_t nThrea
723723

724724
if (iThread == 0) {
725725
if (iBlock == 0) {
726-
#ifdef GPUCA_DETERMINISTIC_MODE
727-
GPUCommonAlgorithm::sortDeviceDynamic(range1, range1 + N1, [](const GPUTPCGMBorderRange& a, const GPUTPCGMBorderRange& b) { return (a.fMin != b.fMin) ? (a.fMin < b.fMin) : (a.fId < b.fId); });
728-
#else
729-
GPUCommonAlgorithm::sortDeviceDynamic(range1, range1 + N1, [](const GPUTPCGMBorderRange& a, const GPUTPCGMBorderRange& b) { return a.fMin < b.fMin; });
730-
#endif
726+
GPUCommonAlgorithm::sortDeviceDynamic(range1, range1 + N1, [](const GPUTPCGMBorderRange& a, const GPUTPCGMBorderRange& b) { return GPUCA_DETERMINISTIC_CODE((a.fMin != b.fMin) ? (a.fMin < b.fMin) : (a.fId < b.fId), a.fMin < b.fMin); });
731727
} else if (iBlock == 1) {
732-
#ifdef GPUCA_DETERMINISTIC_MODE
733-
GPUCommonAlgorithm::sortDeviceDynamic(range2, range2 + N2, [](const GPUTPCGMBorderRange& a, const GPUTPCGMBorderRange& b) { return (a.fMax != b.fMax) ? (a.fMax < b.fMax) : (a.fId < b.fId); });
734-
#else
735-
GPUCommonAlgorithm::sortDeviceDynamic(range2, range2 + N2, [](const GPUTPCGMBorderRange& a, const GPUTPCGMBorderRange& b) { return a.fMax < b.fMax; });
736-
#endif
728+
GPUCommonAlgorithm::sortDeviceDynamic(range2, range2 + N2, [](const GPUTPCGMBorderRange& a, const GPUTPCGMBorderRange& b) { return GPUCA_DETERMINISTIC_CODE((a.fMax != b.fMax) ? (a.fMax < b.fMax) : (a.fId < b.fId), a.fMax < b.fMax); });
737729
}
738730
}
739731
#else
@@ -749,21 +741,13 @@ namespace // anonymous
749741
struct MergeBorderTracks_compMax {
750742
GPUd() bool operator()(const GPUTPCGMBorderRange& a, const GPUTPCGMBorderRange& b)
751743
{
752-
#ifdef GPUCA_DETERMINISTIC_MODE
753-
return (a.fMax != b.fMax) ? (a.fMax < b.fMax) : (a.fId < b.fId);
754-
#else
755-
return a.fMax < b.fMax;
756-
#endif
744+
return GPUCA_DETERMINISTIC_CODE((a.fMax != b.fMax) ? (a.fMax < b.fMax) : (a.fId < b.fId), a.fMax < b.fMax);
757745
}
758746
};
759747
struct MergeBorderTracks_compMin {
760748
GPUd() bool operator()(const GPUTPCGMBorderRange& a, const GPUTPCGMBorderRange& b)
761749
{
762-
#ifdef GPUCA_DETERMINISTIC_MODE
763-
return (a.fMin != b.fMin) ? (a.fMin < b.fMin) : (a.fId < b.fId);
764-
#else
765-
return a.fMin < b.fMin;
766-
#endif
750+
return GPUCA_DETERMINISTIC_CODE((a.fMin != b.fMin) ? (a.fMin < b.fMin) : (a.fId < b.fId), a.fMin < b.fMin);
767751
}
768752
};
769753
} // anonymous namespace
@@ -904,11 +888,7 @@ GPUd() void GPUTPCGMMerger::MergeBorderTracks<2>(int32_t nBlocks, int32_t nThrea
904888

905889
mTrackLinks[b1.TrackID()] = iBest2;
906890
if (mergeMode > 0) {
907-
#ifdef GPUCA_DETERMINISTIC_MODE
908-
CAMath::AtomicMax(&mTrackLinks[iBest2], b1.TrackID());
909-
#else
910-
mTrackLinks[iBest2] = b1.TrackID();
911-
#endif
891+
GPUCA_DETERMINISTIC_CODE(CAMath::AtomicMax(&mTrackLinks[iBest2], b1.TrackID()), mTrackLinks[iBest2] = b1.TrackID());
912892
}
913893
}
914894
// GPUInfo("STAT: sectors %d, %d: all %d merged %d", iSector1, iSector2, statAll, statMerged);
@@ -1467,14 +1447,7 @@ struct GPUTPCGMMerger_CompareClusterIdsLooper {
14671447
if (a1.row != b1.row) {
14681448
return ((a1.row > b1.row) ^ ((a.leg - leg) & 1) ^ outwards);
14691449
}
1470-
#ifdef GPUCA_DETERMINISTIC_MODE
1471-
if (a1.id != b1.id) {
1472-
return (a1.id > b1.id);
1473-
}
1474-
return aa > bb;
1475-
#else
1476-
return a1.id > b1.id;
1477-
#endif
1450+
return GPUCA_DETERMINISTIC_CODE((a1.id != b1.id) ? (a1.id > b1.id) : (aa > bb), a1.id > b1.id);
14781451
}
14791452
};
14801453

@@ -1488,14 +1461,7 @@ struct GPUTPCGMMerger_CompareClusterIds {
14881461
if (a.row != b.row) {
14891462
return (a.row > b.row);
14901463
}
1491-
#ifdef GPUCA_DETERMINISTIC_MODE
1492-
if (a.id != b.id) {
1493-
return (a.id > b.id);
1494-
}
1495-
return aa > bb;
1496-
#else
1497-
return (a.id > b.id);
1498-
#endif
1464+
return GPUCA_DETERMINISTIC_CODE((a.id != b.id) ? (a.id > b.id) : (aa > bb), a.id > b.id);
14991465
}
15001466
};
15011467
} // anonymous namespace
@@ -1567,20 +1533,20 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
15671533
// unpack and sort clusters
15681534
if (nParts > 1 && leg == 0) {
15691535
GPUCommonAlgorithm::sort(trackParts, trackParts + nParts, [](const GPUTPCGMSectorTrack* a, const GPUTPCGMSectorTrack* b) {
1570-
#ifdef GPUCA_DETERMINISTIC_MODE
1571-
if (a->X() != b->X()) {
1536+
GPUCA_DETERMINISTIC_CODE( // clang-format off
1537+
if (a->X() != b->X()) {
1538+
return (a->X() > b->X());
1539+
}
1540+
if (a->Y() != b->Y()) {
1541+
return (a->Y() > b->Y());
1542+
}
1543+
if (a->Z() != b->Z()) {
1544+
return (a->Z() > b->Z());
1545+
}
1546+
return a->QPt() > b->QPt();
1547+
, // !GPUCA_DETERMINISTIC_CODE
15721548
return (a->X() > b->X());
1573-
}
1574-
if (a->Y() != b->Y()) {
1575-
return (a->Y() > b->Y());
1576-
}
1577-
if (a->Z() != b->Z()) {
1578-
return (a->Z() > b->Z());
1579-
}
1580-
return a->QPt() > b->QPt();
1581-
#else
1582-
return (a->X() > b->X());
1583-
#endif
1549+
) // clang-format on
15841550
});
15851551
}
15861552

@@ -1832,20 +1798,18 @@ struct GPUTPCGMMergerSortTracks_comp {
18321798
if (a.Legs() != b.Legs()) {
18331799
return a.Legs() > b.Legs();
18341800
}
1835-
#ifdef GPUCA_DETERMINISTIC_MODE
1836-
if (a.NClusters() != b.NClusters()) {
1801+
GPUCA_DETERMINISTIC_CODE( // clang-format off
1802+
if (a.NClusters() != b.NClusters()) {
1803+
return a.NClusters() > b.NClusters();
1804+
} if (CAMath::Abs(a.GetParam().GetQPt()) != CAMath::Abs(b.GetParam().GetQPt())) {
1805+
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1806+
} if (a.GetParam().GetY() != b.GetParam().GetY()) {
1807+
return a.GetParam().GetY() > b.GetParam().GetY();
1808+
}
1809+
return aa > bb;
1810+
, // !GPUCA_DETERMINISTIC_CODE
18371811
return a.NClusters() > b.NClusters();
1838-
}
1839-
if (CAMath::Abs(a.GetParam().GetQPt()) != CAMath::Abs(b.GetParam().GetQPt())) {
1840-
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1841-
}
1842-
if (a.GetParam().GetY() != b.GetParam().GetY()) {
1843-
return a.GetParam().GetY() > b.GetParam().GetY();
1844-
}
1845-
return aa > bb;
1846-
#else
1847-
return a.NClusters() > b.NClusters();
1848-
#endif
1812+
) // clang-format on
18491813
}
18501814
};
18511815

@@ -1856,17 +1820,16 @@ struct GPUTPCGMMergerSortTracksQPt_comp {
18561820
{
18571821
const GPUTPCGMMergedTrack& GPUrestrict() a = mCmp[aa];
18581822
const GPUTPCGMMergedTrack& GPUrestrict() b = mCmp[bb];
1859-
#ifdef GPUCA_DETERMINISTIC_MODE
1860-
if (CAMath::Abs(a.GetParam().GetQPt()) != CAMath::Abs(b.GetParam().GetQPt())) {
1823+
GPUCA_DETERMINISTIC_CODE( // clang-format off
1824+
if (CAMath::Abs(a.GetParam().GetQPt()) != CAMath::Abs(b.GetParam().GetQPt())) {
1825+
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1826+
} if (a.GetParam().GetY() != b.GetParam().GetY()) {
1827+
return a.GetParam().GetY() > b.GetParam().GetY();
1828+
}
1829+
return a.GetParam().GetZ() > b.GetParam().GetZ();
1830+
, // !GPUCA_DETERMINISTIC_CODE
18611831
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1862-
}
1863-
if (a.GetParam().GetY() != b.GetParam().GetY()) {
1864-
return a.GetParam().GetY() > b.GetParam().GetY();
1865-
}
1866-
return a.GetParam().GetZ() > b.GetParam().GetZ();
1867-
#else
1868-
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1869-
#endif
1832+
) // clang-format on
18701833
}
18711834
};
18721835
} // anonymous namespace
@@ -1901,20 +1864,18 @@ GPUd() void GPUTPCGMMerger::SortTracks(int32_t nBlocks, int32_t nThreads, int32_
19011864
if (a.Legs() != b.Legs()) {
19021865
return a.Legs() > b.Legs();
19031866
}
1904-
#ifdef GPUCA_DETERMINISTIC_MODE
1905-
if (a.NClusters() != b.NClusters()) {
1867+
GPUCA_DETERMINISTIC_CODE( // clang-format off
1868+
if (a.NClusters() != b.NClusters()) {
1869+
return a.NClusters() > b.NClusters();
1870+
} if (CAMath::Abs(a.GetParam().GetQPt()) != CAMath::Abs(b.GetParam().GetQPt())) {
1871+
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1872+
} if (a.GetParam().GetY() != b.GetParam().GetY()) {
1873+
return a.GetParam().GetY() > b.GetParam().GetY();
1874+
}
1875+
return aa > bb;
1876+
, // !GPUCA_DETERMINISTIC_CODE
19061877
return a.NClusters() > b.NClusters();
1907-
}
1908-
if (CAMath::Abs(a.GetParam().GetQPt()) != CAMath::Abs(b.GetParam().GetQPt())) {
1909-
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1910-
}
1911-
if (a.GetParam().GetY() != b.GetParam().GetY()) {
1912-
return a.GetParam().GetY() > b.GetParam().GetY();
1913-
}
1914-
return aa > bb;
1915-
#else
1916-
return a.NClusters() > b.NClusters();
1917-
#endif
1878+
) // clang-format on
19181879
};
19191880

19201881
GPUCommonAlgorithm::sortDeviceDynamic(mTrackOrderProcess, mTrackOrderProcess + mMemory->nOutputTracks, comp);
@@ -1931,17 +1892,16 @@ GPUd() void GPUTPCGMMerger::SortTracksQPt(int32_t nBlocks, int32_t nThreads, int
19311892
auto comp = [cmp = mOutputTracks](const int32_t aa, const int32_t bb) {
19321893
const GPUTPCGMMergedTrack& GPUrestrict() a = cmp[aa];
19331894
const GPUTPCGMMergedTrack& GPUrestrict() b = cmp[bb];
1934-
#ifdef GPUCA_DETERMINISTIC_MODE
1935-
if (CAMath::Abs(a.GetParam().GetQPt()) != CAMath::Abs(b.GetParam().GetQPt())) {
1895+
GPUCA_DETERMINISTIC_CODE( // clang-format off
1896+
if (CAMath::Abs(a.GetParam().GetQPt()) != CAMath::Abs(b.GetParam().GetQPt())) {
1897+
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1898+
} if (a.GetParam().GetY() != b.GetParam().GetY()) {
1899+
return a.GetParam().GetY() > b.GetParam().GetY();
1900+
}
1901+
return a.GetParam().GetZ() > b.GetParam().GetZ();
1902+
, // !GPUCA_DETERMINISTIC_CODE
19361903
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1937-
}
1938-
if (a.GetParam().GetY() != b.GetParam().GetY()) {
1939-
return a.GetParam().GetY() > b.GetParam().GetY();
1940-
}
1941-
return a.GetParam().GetZ() > b.GetParam().GetZ();
1942-
#else
1943-
return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1944-
#endif
1904+
) // clang-format on
19451905
};
19461906

19471907
GPUCommonAlgorithm::sortDeviceDynamic(mTrackSort, mTrackSort + mMemory->nOutputTracks, comp);

‎GPU/GPUTracking/Merger/GPUTPCGMMergerTypes.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@ enum attachTypes { attachAttached = 0x40000000,
3232

3333
struct InterpolationErrorHit {
3434
float posY, posZ;
35-
GPUCA_MERGER_INTERPOLATION_ERROR_TYPE errorY, errorZ;
35+
GPUCA_MERGER_INTERPOLATION_ERROR_TYPE_A errorY, errorZ;
3636
};
3737

3838
struct InterpolationErrors {

0 commit comments

Comments
 (0)