Skip to content

Commit 514e953

Browse files
davidrohrdavidrohr
committed
GPU: Remove some template magic that was only needed for OpenCL 1.2
1 parent e01c675 commit 514e953

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

43 files changed

+261
-526
lines changed

‎GPU/GPUTracking/Base/GPUConstantMem.h‎

Lines changed: 5 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -71,12 +71,10 @@ namespace GPUCA_NAMESPACE
7171
{
7272
namespace gpu
7373
{
74-
MEM_CLASS_PRE()
7574
struct GPUConstantMem {
76-
MEM_CONSTANT(GPUParam)
77-
param;
78-
MEM_GLOBAL(GPUTPCTracker)
79-
tpcTrackers[GPUCA_NSLICES];
75+
GPUParam param;
76+
GPUTPCTracker
77+
tpcTrackers[GPUCA_NSLICES];
8078
GPUTPCConvert tpcConverter;
8179
GPUTPCCompression tpcCompressor;
8280
GPUTPCDecompression tpcDecompressor;
@@ -150,7 +148,7 @@ namespace gpu
150148
{
151149

152150
// Must be placed here, to avoid circular header dependency
153-
GPUdi() GPUconstantref() const MEM_CONSTANT(GPUConstantMem) * GPUProcessor::GetConstantMem() const
151+
GPUdi() GPUconstantref() const GPUConstantMem* GPUProcessor::GetConstantMem() const
154152
{
155153
#if defined(GPUCA_GPUCODE_DEVICE) && defined(GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM) && !defined(GPUCA_GPUCODE_HOSTONLY)
156154
return &GPUCA_CONSMEM;
@@ -159,7 +157,7 @@ GPUdi() GPUconstantref() const MEM_CONSTANT(GPUConstantMem) * GPUProcessor::GetC
159157
#endif
160158
}
161159

162-
GPUdi() GPUconstantref() const MEM_CONSTANT(GPUParam) & GPUProcessor::Param() const
160+
GPUdi() GPUconstantref() const GPUParam& GPUProcessor::Param() const
163161
{
164162
return GetConstantMem()->param;
165163
}

‎GPU/GPUTracking/Base/GPUGeneralKernels.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@
1717
using namespace GPUCA_NAMESPACE::gpu;
1818

1919
template <>
20-
GPUdii() void GPUMemClean16::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& GPUrestrict() processors, GPUglobalref() void* ptr, uint64_t size)
20+
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)
2121
{
2222
const uint64_t stride = get_global_size(0);
2323
int4 i0;
@@ -30,7 +30,7 @@ GPUdii() void GPUMemClean16::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_
3030
}
3131

3232
template <>
33-
GPUdii() void GPUitoa::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& GPUrestrict() processors, GPUglobalref() int32_t* ptr, uint64_t size)
33+
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)
3434
{
3535
const uint64_t stride = get_global_size(0);
3636
for (uint64_t i = get_global_id(0); i < size; i += stride) {

‎GPU/GPUTracking/Base/GPUGeneralKernels.h‎

Lines changed: 6 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,6 @@ namespace GPUCA_NAMESPACE
3636
{
3737
namespace gpu
3838
{
39-
MEM_CLASS_PRE()
4039
struct GPUConstantMem;
4140

4241
class GPUKernelTemplate
@@ -50,7 +49,6 @@ class GPUKernelTemplate
5049
step4 = 4,
5150
step5 = 5 };
5251

53-
MEM_CLASS_PRE()
5452
struct GPUSharedMemory {
5553
};
5654

@@ -82,21 +80,20 @@ class GPUKernelTemplate
8280
#endif
8381
};
8482

85-
typedef GPUconstantref() MEM_CONSTANT(GPUConstantMem) processorType;
83+
typedef GPUconstantref() GPUConstantMem processorType;
8684
GPUhdi() CONSTEXPR static GPUDataTypes::RecoStep GetRecoStep() { return GPUCA_RECO_STEP::NoRecoStep; }
87-
MEM_TEMPLATE()
88-
GPUhdi() static processorType* Processor(MEM_TYPE(GPUConstantMem) & processors)
85+
GPUhdi() static processorType* Processor(GPUConstantMem& processors)
8986
{
9087
return &processors;
9188
}
9289
#ifdef GPUCA_NOCOMPAT
9390
template <int32_t iKernel, typename... Args>
94-
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& processors, Args... args)
91+
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, Args... args)
9592
{
9693
}
9794
#else
9895
template <int32_t iKernel>
99-
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& processors)
96+
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors)
10097
{
10198
}
10299
#endif
@@ -108,7 +105,7 @@ class GPUMemClean16 : public GPUKernelTemplate
108105
public:
109106
GPUhdi() CONSTEXPR static GPUDataTypes::RecoStep GetRecoStep() { return GPUCA_RECO_STEP::NoRecoStep; }
110107
template <int32_t iKernel = defaultKernel>
111-
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& processors, GPUglobalref() void* ptr, uint64_t size);
108+
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUglobalref() void* ptr, uint64_t size);
112109
};
113110

114111
// Fill with incrementing sequnce of integers
@@ -117,7 +114,7 @@ class GPUitoa : public GPUKernelTemplate
117114
public:
118115
GPUhdi() CONSTEXPR static GPUDataTypes::RecoStep GetRecoStep() { return GPUCA_RECO_STEP::NoRecoStep; }
119116
template <int32_t iKernel = defaultKernel>
120-
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& processors, GPUglobalref() int32_t* ptr, uint64_t size);
117+
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUglobalref() int32_t* ptr, uint64_t size);
121118
};
122119

123120
} // namespace gpu

‎GPU/GPUTracking/Base/GPUParam.h‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -79,7 +79,6 @@ struct GPUParam_t {
7979
} // namespace internal
8080

8181
#if !(defined(__CINT__) || defined(__ROOTCINT__)) || defined(__CLING__) // Hide from ROOT 5 CINT
82-
MEM_CLASS_PRE()
8382
struct GPUParam : public internal::GPUParam_t<GPUSettingsRec, GPUSettingsParam> {
8483

8584
#ifndef GPUCA_GPUCODE

‎GPU/GPUTracking/Base/GPUParam.inc‎

Lines changed: 20 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -24,17 +24,15 @@ namespace GPUCA_NAMESPACE
2424
namespace gpu
2525
{
2626

27-
MEM_CLASS_PRE()
28-
GPUdi() void MEM_LG(GPUParam)::Slice2Global(int32_t iSlice, float x, float y, float z, float* X, float* Y, float* Z) const
27+
GPUdi() void GPUParam::Slice2Global(int32_t iSlice, float x, float y, float z, float* X, float* Y, float* Z) const
2928
{
3029
// conversion of coordinates sector->global
3130
*X = x * SliceParam[iSlice].CosAlpha - y * SliceParam[iSlice].SinAlpha;
3231
*Y = y * SliceParam[iSlice].CosAlpha + x * SliceParam[iSlice].SinAlpha;
3332
*Z = z;
3433
}
3534

36-
MEM_CLASS_PRE()
37-
GPUdi() void MEM_LG(GPUParam)::Global2Slice(int32_t iSlice, float X, float Y, float Z, float* x, float* y, float* z) const
35+
GPUdi() void GPUParam::Global2Slice(int32_t iSlice, float X, float Y, float Z, float* x, float* y, float* z) const
3836
{
3937
// conversion of coordinates global->sector
4038
*x = X * SliceParam[iSlice].CosAlpha + Y * SliceParam[iSlice].SinAlpha;
@@ -44,8 +42,7 @@ GPUdi() void MEM_LG(GPUParam)::Global2Slice(int32_t iSlice, float X, float Y, fl
4442

4543
#ifdef GPUCA_TPC_GEOMETRY_O2
4644

47-
MEM_CLASS_PRE()
48-
GPUdi() void MEM_LG(GPUParam)::GetClusterErrorsSeeding2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float& ErrY2, float& ErrZ2) const
45+
GPUdi() void GPUParam::GetClusterErrorsSeeding2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float& ErrY2, float& ErrZ2) const
4946
{
5047
const int32_t rowType = tpcGeometry.GetROC(iRow);
5148
z = CAMath::Abs(tpcGeometry.TPCLength() - CAMath::Abs(z));
@@ -60,30 +57,27 @@ GPUdi() void MEM_LG(GPUParam)::GetClusterErrorsSeeding2(uint8_t sector, int32_t
6057
ErrZ2 = GetClusterErrorSeeding(1, rowType, z, angleZ2, unscaledMult); // Returns Err2
6158
}
6259

63-
MEM_CLASS_PRE()
64-
GPUdi() float MEM_LG(GPUParam)::GetClusterErrorSeeding(int32_t yz, int32_t type, float zDiff, float angle2, float unscaledMult) const // Note, returns Err2 despite the name not containing 2
60+
GPUdi() float GPUParam::GetClusterErrorSeeding(int32_t yz, int32_t type, float zDiff, float angle2, float unscaledMult) const // Note, returns Err2 despite the name not containing 2
6561
{
66-
MakeType(const float*) c = ParamErrors[yz][type]; // Note: c[0] = p[0]^2, c[1] = p[1]^2 * padHeight, c[2] = p[2]^2 / tpcLength / padHeight, c[3] = p[3]^2 * clusterErrorOccupancyScaler^2
62+
const float* c = ParamErrors[yz][type]; // Note: c[0] = p[0]^2, c[1] = p[1]^2 * padHeight, c[2] = p[2]^2 / tpcLength / padHeight, c[3] = p[3]^2 * clusterErrorOccupancyScaler^2
6763
float v = c[0] + c[1] * angle2 + c[2] * zDiff + c[3] * (unscaledMult * unscaledMult);
6864
v = CAMath::Abs(v);
6965
v *= yz ? rec.tpc.clusterError2CorrectionZ : rec.tpc.clusterError2CorrectionY;
7066
v += yz ? rec.tpc.clusterError2AdditionalZSeeding : rec.tpc.clusterError2AdditionalYSeeding;
7167
return v;
7268
}
7369

74-
MEM_CLASS_PRE()
75-
GPUdi() float MEM_LG(GPUParam)::GetClusterError2(int32_t yz, int32_t type, float zDiff, float angle2, float unscaledMult, float scaledInvAvgCharge, float scaledInvCharge) const
70+
GPUdi() float GPUParam::GetClusterError2(int32_t yz, int32_t type, float zDiff, float angle2, float unscaledMult, float scaledInvAvgCharge, float scaledInvCharge) const
7671
{
77-
MakeType(const float*) c = ParamErrors[yz][type]; // Note: c[0] = p[0]^2, c[1] = p[1]^2 * padHeight, c[2] = p[2]^2 / tpcLength / padHeight, c[3] = p[3]^2 * clusterErrorOccupancyScaler^2
72+
const float* c = ParamErrors[yz][type]; // Note: c[0] = p[0]^2, c[1] = p[1]^2 * padHeight, c[2] = p[2]^2 / tpcLength / padHeight, c[3] = p[3]^2 * clusterErrorOccupancyScaler^2
7873
float v = c[0] + c[1] * angle2 * scaledInvAvgCharge + c[2] * zDiff * scaledInvCharge + c[3] * (unscaledMult * unscaledMult) * (scaledInvAvgCharge * scaledInvAvgCharge);
7974
v = CAMath::Abs(v);
8075
v *= yz ? rec.tpc.clusterError2CorrectionZ : rec.tpc.clusterError2CorrectionY;
8176
v += yz ? rec.tpc.clusterError2AdditionalZ : rec.tpc.clusterError2AdditionalY;
8277
return v;
8378
}
8479

85-
MEM_CLASS_PRE()
86-
GPUdi() float MEM_LG(GPUParam)::GetSystematicClusterErrorIFC2(float x, float y, float z, bool sideC) const
80+
GPUdi() float GPUParam::GetSystematicClusterErrorIFC2(float x, float y, float z, bool sideC) const
8781
{
8882
float sysErr = 0.f;
8983
const float kMaxExpArg = 9.f; // limit r-dumped error to this exp. argument
@@ -116,8 +110,7 @@ GPUdi() float MEM_LG(GPUParam)::GetSystematicClusterErrorIFC2(float x, float y,
116110
return sysErr;
117111
}
118112

119-
MEM_CLASS_PRE()
120-
GPUdi() float MEM_LG(GPUParam)::GetSystematicClusterErrorC122(float x, float y, uint8_t sector) const
113+
GPUdi() float GPUParam::GetSystematicClusterErrorC122(float x, float y, uint8_t sector) const
121114
{
122115
const float dx = x - 83.f;
123116
if (dx > occupancyTotal * rec.tpc.sysClusErrorC12Box) {
@@ -131,17 +124,15 @@ GPUdi() float MEM_LG(GPUParam)::GetSystematicClusterErrorC122(float x, float y,
131124

132125
#else // GPUCA_TPC_GEOMETRY_O2
133126

134-
MEM_CLASS_PRE()
135-
GPUdi() float MEM_LG(GPUParam)::GetClusterErrorSeeding(int32_t yz, int32_t type, float zDiff, float angle2, float scaledMult) const
127+
GPUdi() float GPUParam::GetClusterErrorSeeding(int32_t yz, int32_t type, float zDiff, float angle2, float scaledMult) const
136128
{
137-
MakeType(const float*) c = ParamErrorsSeeding0[yz][type];
129+
const float* c = ParamErrorsSeeding0[yz][type];
138130
float v = c[0] + c[1] * zDiff + c[2] * angle2;
139131
v = CAMath::Abs(v);
140132
return v;
141133
}
142134

143-
MEM_CLASS_PRE()
144-
GPUdi() void MEM_LG(GPUParam)::GetClusterErrorsSeeding2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float& ErrY2, float& ErrZ2) const
135+
GPUdi() void GPUParam::GetClusterErrorsSeeding2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float& ErrY2, float& ErrZ2) const
145136
{
146137
int32_t rowType = tpcGeometry.GetROC(iRow);
147138
z = CAMath::Abs(tpcGeometry.TPCLength() - CAMath::Abs(z));
@@ -156,10 +147,9 @@ GPUdi() void MEM_LG(GPUParam)::GetClusterErrorsSeeding2(uint8_t sector, int32_t
156147
ErrZ2 = ErrZ2 * ErrZ2 * rec.tpc.clusterError2CorrectionZ + rec.tpc.clusterError2AdditionalZ;
157148
}
158149

159-
MEM_CLASS_PRE()
160-
GPUdi() float MEM_LG(GPUParam)::GetClusterError2(int32_t yz, int32_t type, float zDiff, float angle2, float unscaledMult, float avgInvCharge, float invCharge) const
150+
GPUdi() float GPUParam::GetClusterError2(int32_t yz, int32_t type, float zDiff, float angle2, float unscaledMult, float avgInvCharge, float invCharge) const
161151
{
162-
MakeType(const float*) c = ParamS0Par[yz][type];
152+
const float* c = ParamS0Par[yz][type];
163153
float v = c[0] + c[1] * zDiff + c[2] * angle2 + c[3] * zDiff * zDiff + c[4] * angle2 * angle2 + c[5] * zDiff * angle2;
164154
v = CAMath::Abs(v);
165155
if (v < 0.0001f) {
@@ -170,22 +160,19 @@ GPUdi() float MEM_LG(GPUParam)::GetClusterError2(int32_t yz, int32_t type, float
170160
return v;
171161
}
172162

173-
MEM_CLASS_PRE()
174-
GPUdi() float MEM_LG(GPUParam)::GetSystematicClusterErrorIFC2(float trackX, float trackY, float z, bool sideC) const
163+
GPUdi() float GPUParam::GetSystematicClusterErrorIFC2(float trackX, float trackY, float z, bool sideC) const
175164
{
176165
return 0;
177166
}
178167

179-
MEM_CLASS_PRE()
180-
GPUdi() float MEM_LG(GPUParam)::GetSystematicClusterErrorC122(float trackX, float trackY, uint8_t sector) const
168+
GPUdi() float GPUParam::GetSystematicClusterErrorC122(float trackX, float trackY, uint8_t sector) const
181169
{
182170
return 0;
183171
}
184172

185173
#endif // !GPUCA_TPC_GEOMETRY_O2
186174

187-
MEM_CLASS_PRE()
188-
GPUdi() void MEM_LG(GPUParam)::GetClusterErrors2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float avgInvCharge, float invCharge, float& ErrY2, float& ErrZ2) const
175+
GPUdi() void GPUParam::GetClusterErrors2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float avgInvCharge, float invCharge, float& ErrY2, float& ErrZ2) const
189176
{
190177
const int32_t rowType = tpcGeometry.GetROC(iRow);
191178
z = CAMath::Abs(tpcGeometry.TPCLength() - CAMath::Abs(z));
@@ -202,8 +189,7 @@ GPUdi() void MEM_LG(GPUParam)::GetClusterErrors2(uint8_t sector, int32_t iRow, f
202189
ErrZ2 = GetClusterError2(1, rowType, z, angleZ2, unscaledMult, scaledInvAvgCharge, scaledInvCharge);
203190
}
204191

205-
MEM_CLASS_PRE()
206-
GPUdi() void MEM_LG(GPUParam)::UpdateClusterError2ByState(int16_t clusterState, float& ErrY2, float& ErrZ2) const
192+
GPUdi() void GPUParam::UpdateClusterError2ByState(int16_t clusterState, float& ErrY2, float& ErrZ2) const
207193
{
208194
if (clusterState & GPUTPCGMMergedTrackHit::flagEdge) {
209195
ErrY2 += rec.tpc.extraClusterErrorEdgeY2;
@@ -223,8 +209,7 @@ GPUdi() void MEM_LG(GPUParam)::UpdateClusterError2ByState(int16_t clusterState,
223209
}
224210
}
225211

226-
MEM_CLASS_PRE()
227-
GPUdi() float MEM_LG(GPUParam)::GetUnscaledMult(float time) const
212+
GPUdi() float GPUParam::GetUnscaledMult(float time) const
228213
{
229214
if (!occupancyMap) {
230215
return 0.f;
@@ -233,8 +218,7 @@ GPUdi() float MEM_LG(GPUParam)::GetUnscaledMult(float time) const
233218
return occupancyMap[bin];
234219
}
235220

236-
MEM_CLASS_PRE()
237-
GPUdi() bool MEM_LG(GPUParam)::rejectEdgeClusterByY(float uncorrectedY, int32_t iRow, float trackSigmaY) const
221+
GPUdi() bool GPUParam::rejectEdgeClusterByY(float uncorrectedY, int32_t iRow, float trackSigmaY) const
238222
{
239223
return CAMath::Abs(uncorrectedY) > (tpcGeometry.NPads(iRow) - 1) * 0.5f * tpcGeometry.PadWidth(iRow) + rec.tpc.rejectEdgeClustersMargin + trackSigmaY * rec.tpc.rejectEdgeClustersSigmaMargin;
240224
}

‎GPU/GPUTracking/Base/GPUProcessor.h‎

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -29,9 +29,7 @@ namespace gpu
2929
{
3030
struct GPUTrackingInOutPointers;
3131
class GPUReconstruction;
32-
MEM_CLASS_PRE()
3332
struct GPUParam;
34-
MEM_CLASS_PRE()
3533
struct GPUConstantMem;
3634

3735
class GPUProcessor
@@ -52,8 +50,8 @@ class GPUProcessor
5250
GPUProcessor& operator=(const GPUProcessor&) CON_DELETE;
5351
#endif
5452

55-
GPUd() GPUconstantref() const MEM_CONSTANT(GPUConstantMem) * GetConstantMem() const; // Body in GPUConstantMem.h to avoid circular headers
56-
GPUd() GPUconstantref() const MEM_CONSTANT(GPUParam) & Param() const; // ...
53+
GPUd() GPUconstantref() const GPUConstantMem* GetConstantMem() const; // Body in GPUConstantMem.h to avoid circular headers
54+
GPUd() GPUconstantref() const GPUParam& Param() const; // ...
5755
GPUd() void raiseError(uint32_t code, uint32_t param1 = 0, uint32_t param2 = 0, uint32_t param3 = 0) const;
5856
const GPUReconstruction& GetRec() const { return *mRec; }
5957

@@ -152,7 +150,7 @@ class GPUProcessor
152150
GPUReconstruction* mRec;
153151
ProcessorType mGPUProcessorType;
154152
GPUProcessor* mLinkedProcessor;
155-
GPUconstantref() const MEM_CONSTANT(GPUConstantMem) * mConstantMem;
153+
GPUconstantref() const GPUConstantMem* mConstantMem;
156154

157155
private:
158156
bool mAllocateAndInitializeLate;

‎GPU/GPUTracking/Base/GPUReconstructionDeviceBase.cxx‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -28,7 +28,6 @@ using namespace GPUCA_NAMESPACE::gpu;
2828
#endif
2929
#include <cstring>
3030

31-
MEM_CLASS_PRE()
3231
class GPUTPCRow;
3332

3433
#define SemLockName "AliceHLTTPCGPUTrackerInitLockSem"

‎GPU/GPUTracking/Base/GPUReconstructionKernelMacros.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -59,7 +59,7 @@
5959
#else
6060
#define GPUCA_KRNLGPU_SINGLE(x_class, x_attributes, x_arguments, x_forward, ...) GPUCA_KRNLGPU_SINGLE_DEF(x_class, x_attributes, x_arguments, x_forward, __VA_ARGS__) \
6161
{ \
62-
GPUshared() typename GPUCA_M_STRIP_FIRST(x_class)::MEM_LOCAL(GPUSharedMemory) smem; \
62+
GPUshared() typename GPUCA_M_STRIP_FIRST(x_class)::GPUSharedMemory smem; \
6363
GPUCA_M_STRIP_FIRST(x_class)::template Thread<GPUCA_M_KRNL_NUM(x_class)>(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, GPUCA_M_STRIP_FIRST(x_class)::Processor(GPUCA_CONSMEM)[iSlice_internal] GPUCA_M_STRIP(x_forward)); \
6464
}
6565
#endif
@@ -76,7 +76,7 @@
7676
const int32_t nSliceBlockOffset = get_num_groups(0) * iSlice_internal / nSliceCount; \
7777
const int32_t sliceBlockId = get_group_id(0) - nSliceBlockOffset; \
7878
const int32_t sliceGridDim = get_num_groups(0) * (iSlice_internal + 1) / nSliceCount - get_num_groups(0) * (iSlice_internal) / nSliceCount; \
79-
GPUshared() typename GPUCA_M_STRIP_FIRST(x_class)::MEM_LOCAL(GPUSharedMemory) smem; \
79+
GPUshared() typename GPUCA_M_STRIP_FIRST(x_class)::GPUSharedMemory smem; \
8080
GPUCA_M_STRIP_FIRST(x_class)::template Thread<GPUCA_M_KRNL_NUM(x_class)>(sliceGridDim, get_local_size(0), sliceBlockId, get_local_id(0), smem, GPUCA_M_STRIP_FIRST(x_class)::Processor(GPUCA_CONSMEM)[firstSlice + iSlice_internal] GPUCA_M_STRIP(x_forward)); \
8181
}
8282
#endif

‎GPU/GPUTracking/Base/opencl-common/GPUReconstructionOCL.cl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -81,7 +81,7 @@ typedef signed char int8_t;
8181
#define GPUCA_KRNL(...) GPUCA_KRNL_WRAP(GPUCA_KRNL_LOAD_, __VA_ARGS__)
8282
#define GPUCA_KRNL_LOAD_single(...) GPUCA_KRNLGPU_SINGLE(__VA_ARGS__)
8383
#define GPUCA_KRNL_LOAD_multi(...) GPUCA_KRNLGPU_MULTI(__VA_ARGS__)
84-
#define GPUCA_CONSMEM_PTR GPUglobal() char *gpu_mem, GPUconstant() MEM_CONSTANT(GPUConstantMem) * pConstant,
84+
#define GPUCA_CONSMEM_PTR GPUglobal() char *gpu_mem, GPUconstant() GPUConstantMem* pConstant,
8585
#define GPUCA_CONSMEM (*pConstant)
8686
#include "GPUReconstructionKernelList.h"
8787
#undef GPUCA_KRNL

‎GPU/GPUTracking/CMakeLists.txt‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -137,7 +137,6 @@ set(HDRS_INSTALL
137137
Definitions/GPUDefGPUParameters.h
138138
Definitions/GPUDef.h
139139
Definitions/GPUDefMacros.h
140-
Definitions/GPUDefOpenCL12Templates.h
141140
Definitions/GPULogging.h
142141
Definitions/GPUSettingsList.h
143142
Global/GPUChainTrackingDefs.h

0 commit comments

Comments
 (0)