diff --git a/GPU/Common/GPUCommonAlgorithm.h b/GPU/Common/GPUCommonAlgorithm.h index d0643391246a8..8cd53ec5e0609 100644 --- a/GPU/Common/GPUCommonAlgorithm.h +++ b/GPU/Common/GPUCommonAlgorithm.h @@ -282,30 +282,27 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(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) { + GPUCommonAlgorithm::sort(begin, end, comp); + } + GPUbarrier(); #else - GPUCA_DETERMINISTIC_CODE( // clang-format off - GPUbarrier(); - if (get_local_id(0) == 0) { - GPUCommonAlgorithm::sort(begin, end, comp); - } - GPUbarrier(); - , // !GPUCA_DETERMINISTIC_CODE - 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)) { - int32_t offset = i % 2; - int32_t curPos = 2 * tIdx + offset; - int32_t nextPos = curPos + 1; - - if (nextPos < n) { - if (!comp(begin[curPos], begin[nextPos])) { - IterSwap(&begin[curPos], &begin[nextPos]); - } + 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)) { + int32_t offset = i % 2; + int32_t curPos = 2 * tIdx + offset; + int32_t nextPos = curPos + 1; + + if (nextPos < n) { + if (!comp(begin[curPos], begin[nextPos])) { + IterSwap(&begin[curPos], &begin[nextPos]); } } - GPUbarrier(); } - ) // clang-format on + GPUbarrier(); + } #endif } diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDArtc.cu b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDArtc.cu index 50a568ab345cf..805397c9b430e 100644 --- a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDArtc.cu +++ b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDArtc.cu @@ -14,9 +14,12 @@ #define GPUCA_GPUCODE_GENRTC #define GPUCA_GPUCODE_COMPILEKERNELS + +// Keep some preprocessor calls unprocessed #define GPUCA_RTC_SPECIAL_CODE(...) GPUCA_RTC_SPECIAL_CODE(__VA_ARGS__) #define GPUCA_DETERMINISTIC_CODE(...) GPUCA_DETERMINISTIC_CODE(__VA_ARGS__) -// GPUReconstructionCUDAIncludesSystem.h prependended without preprocessor running + +// GPUReconstructionCUDAIncludesSystem.h prependended by CMakewithout preprocessor running #include "GPUReconstructionCUDADef.h" #include "GPUReconstructionIncludesDeviceAll.h" diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx index 3b88c8764d0fd..bba97e9eace9b 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx +++ b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx @@ -273,19 +273,19 @@ GPUdii() void GPUTPCCompressionKernels::Thread(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])); - , // !GPUCA_DETERMINISTIC_CODE - if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZPadTime) { - CAAlgo::sortInBlock(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])); - } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortPad) { - CAAlgo::sortInBlock(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])); - } - ) // clang-format on + } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZTimePad) { + CAAlgo::sortInBlock(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])); + } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortTime) { + CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare(clusters->clusters[iSector][iRow])); + } +#endif // GPUCA_DETERMINISTIC_MODE GPUbarrier(); }