Skip to content

Commit 491aeb7

Browse files
davidrohrdavidrohr
committed
GPU: shared memory variables should never be declared as restrict
1 parent 0ac25f6 commit 491aeb7

16 files changed

+59
-60
lines changed

‎GPU/GPUTracking/Base/GPUGeneralKernels.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,7 @@
1515
using namespace GPUCA_NAMESPACE::gpu;
1616

1717
template <>
18-
GPUdii() void GPUMemClean16::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & GPUrestrict() smem, processorType& GPUrestrict() processors, GPUglobalref() void* ptr, unsigned long size)
18+
GPUdii() void GPUMemClean16::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& GPUrestrict() processors, GPUglobalref() void* ptr, unsigned long size)
1919
{
2020
const unsigned long stride = get_global_size(0);
2121
int4 i0;

‎GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -176,7 +176,7 @@ GPUd() bool GPUTPCCompressionKernels::GPUTPCCompressionKernels_Compare<3>::opera
176176
}
177177

178178
template <>
179-
GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1unattached>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() processors)
179+
GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1unattached>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors)
180180
{
181181
const GPUTrackingInOutPointers& GPUrestrict() ioPtrs = processors.ioPtrs;
182182
const o2::tpc::ClusterNativeAccess* GPUrestrict() clusters = ioPtrs.clustersNative;

‎GPU/GPUTracking/ITS/GPUITSFitterKernels.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -55,7 +55,7 @@ GPUdii() bool GPUITSFitterKernel::fitTrack(GPUITSFitter& GPUrestrict() Fitter, G
5555
}
5656

5757
template <>
58-
GPUdii() void GPUITSFitterKernel::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() processors)
58+
GPUdii() void GPUITSFitterKernel::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors)
5959
{
6060
GPUITSFitter& Fitter = processors.itsFitter;
6161

‎GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx‎

Lines changed: 32 additions & 32 deletions
Large diffs are not rendered by default.

‎GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -37,7 +37,7 @@ struct tmpSort {
3737
};
3838

3939
template <>
40-
GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::prepare>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() merger)
40+
GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::prepare>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
4141
{
4242
const GPUTPCGMMergedTrack* tracks = merger.OutputTracks();
4343
const unsigned int nTracks = merger.NOutputTracks();
@@ -67,7 +67,7 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::prepare>(int nBlocks, i
6767
}
6868

6969
template <>
70-
GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::sort>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() merger)
70+
GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::sort>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
7171
{
7272
if (iThread || iBlock) {
7373
return;
@@ -96,7 +96,7 @@ void GPUCA_KRNL_BACKEND_CLASS::runKernelBackendInternal<GPUTPCGMO2Output, GPUTPC
9696
#endif // __CUDACC__ || __HIPCC__ - Specialize GPUTPCGMMergerSortTracks and GPUTPCGMMergerSortTracksQPt
9797

9898
template <>
99-
GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() merger)
99+
GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
100100
{
101101
constexpr float MinDelta = 0.1;
102102
const GPUTPCGMMergedTrack* tracks = merger.OutputTracks();

‎GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -528,8 +528,7 @@ GPUd() void GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestrict
528528
const unsigned int hitLst = CA_TEXTURE_FETCH(calink, gAliTexRefu, firsthit, mybin + ny + 1);
529529
for (unsigned int ih = hitFst; ih < hitLst; ih++) {
530530
int id = idOffset + ids[ih];
531-
GPUAtomic(unsigned int) * GPUrestrict() const weight = weights + id;
532-
;
531+
GPUAtomic(unsigned int)* const weight = weights + id;
533532
#if !defined(GPUCA_NO_ATOMIC_PRECHECK) && GPUCA_NO_ATOMIC_PRECHECK < 1
534533
if (myWeight <= *weight) {
535534
continue;

‎GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx‎

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

1919
template <int I>
20-
GPUdii() void GPUTrackingRefitKernel::Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() processors)
20+
GPUdii() void GPUTrackingRefitKernel::Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors)
2121
{
2222
auto& refit = processors.trackingRefit;
2323
for (unsigned int i = get_global_id(0); i < processors.ioPtrs.nMergedTracks; i += get_global_size(0)) {
@@ -37,5 +37,5 @@ GPUdii() void GPUTrackingRefitKernel::Thread(int nBlocks, int nThreads, int iBlo
3737
}
3838
}
3939
}
40-
template GPUd() void GPUTrackingRefitKernel::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() processors);
41-
template GPUd() void GPUTrackingRefitKernel::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& GPUrestrict() smem, processorType& GPUrestrict() processors);
40+
template GPUd() void GPUTrackingRefitKernel::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors);
41+
template GPUd() void GPUTrackingRefitKernel::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors);

‎GPU/GPUTracking/SliceTracker/GPUTPCCreateSliceData.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818
using namespace GPUCA_NAMESPACE::gpu;
1919

2020
template <>
21-
GPUdii() void GPUTPCCreateSliceData::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & GPUrestrict() s, processorType& GPUrestrict() tracker)
21+
GPUdii() void GPUTPCCreateSliceData::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & s, processorType& GPUrestrict() tracker)
2222
{
2323
tracker.Data().InitFromClusterData(nBlocks, nThreads, iBlock, iThread, tracker.GetConstantMem(), tracker.ISlice(), s.tmp);
2424
}

‎GPU/GPUTracking/SliceTracker/GPUTPCGlobalTracking.cxx‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@ using namespace GPUCA_NAMESPACE::gpu;
2323

2424
#if !defined(__OPENCL__) || defined(__OPENCLCPP__)
2525

26-
GPUd() int GPUTPCGlobalTracking::PerformGlobalTrackingRun(GPUTPCTracker& tracker, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & GPUrestrict() smem, const GPUTPCTracker& GPUrestrict() sliceSource, int iTrack, int rowIndex, float angle, int direction)
26+
GPUd() int GPUTPCGlobalTracking::PerformGlobalTrackingRun(GPUTPCTracker& tracker, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, const GPUTPCTracker& GPUrestrict() sliceSource, int iTrack, int rowIndex, float angle, int direction)
2727
{
2828
/*for (int j = 0;j < Tracks()[j].NHits();j++)
2929
{
@@ -113,7 +113,7 @@ GPUd() int GPUTPCGlobalTracking::PerformGlobalTrackingRun(GPUTPCTracker& tracker
113113
return (nHits >= GPUCA_GLOBAL_TRACKING_MIN_HITS);
114114
}
115115

116-
GPUd() void GPUTPCGlobalTracking::PerformGlobalTracking(int nBlocks, int nThreads, int iBlock, int iThread, const GPUTPCTracker& tracker, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & GPUrestrict() smem, GPUTPCTracker& GPUrestrict() sliceTarget, bool right)
116+
GPUd() void GPUTPCGlobalTracking::PerformGlobalTracking(int nBlocks, int nThreads, int iBlock, int iThread, const GPUTPCTracker& tracker, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, GPUTPCTracker& GPUrestrict() sliceTarget, bool right)
117117
{
118118
for (int i = iBlock * nThreads + iThread; i < tracker.CommonMemory()->nLocalTracks; i += nThreads * nBlocks) {
119119
{
@@ -161,7 +161,7 @@ GPUd() void GPUTPCGlobalTracking::PerformGlobalTracking(int nBlocks, int nThread
161161
}
162162

163163
template <>
164-
GPUdii() void GPUTPCGlobalTracking::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & GPUrestrict() smem, processorType& GPUrestrict() tracker)
164+
GPUdii() void GPUTPCGlobalTracking::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& GPUrestrict() tracker)
165165
{
166166
CA_SHARED_CACHE(&smem.mRows[0], tracker.SliceDataRows(), GPUCA_ROW_COUNT * sizeof(MEM_PLAIN(GPUTPCRow)));
167167
GPUbarrier();
@@ -204,7 +204,7 @@ GPUd() void GPUTPCGlobalTracking::GlobalTrackingSliceLeftRight(unsigned int iSli
204204
#endif // !__OPENCL__ || __OPENCLCPP__
205205

206206
template <>
207-
GPUdii() void GPUTPCGlobalTrackingCopyNumbers::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & GPUrestrict() smem, processorType& GPUrestrict() tracker, int n)
207+
GPUdii() void GPUTPCGlobalTrackingCopyNumbers::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& GPUrestrict() tracker, int n)
208208
{
209209
for (int i = get_global_id(0); i < n; i += get_global_size(0)) {
210210
GPUconstantref() MEM_GLOBAL(GPUTPCTracker) & GPUrestrict() trk = (&tracker)[i];

‎GPU/GPUTracking/SliceTracker/GPUTPCNeighboursCleaner.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@
1717
using namespace GPUCA_NAMESPACE::gpu;
1818

1919
template <>
20-
GPUdii() void GPUTPCNeighboursCleaner::Thread<0>(int /*nBlocks*/, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & GPUrestrict() s, processorType& GPUrestrict() tracker)
20+
GPUdii() void GPUTPCNeighboursCleaner::Thread<0>(int /*nBlocks*/, int nThreads, int iBlock, int iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & s, processorType& GPUrestrict() tracker)
2121
{
2222
// *
2323
// * kill link to the neighbour if the neighbour is not pointed to the cluster

0 commit comments

Comments
 (0)