From 3d4dfde5d7c37fd747ceefd789c11c899d3a7714 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Fri, 18 Apr 2025 19:55:43 +0200 Subject: [PATCH 1/2] GPU RTC: Don't pass a third launch bounds parameter --- GPU/GPUTracking/Base/cuda/GPUReconstructionCUDARTCCalls.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDARTCCalls.cu b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDARTCCalls.cu index 571428dc39e21..3e4d3113fb995 100644 --- a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDARTCCalls.cu +++ b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDARTCCalls.cu @@ -15,7 +15,7 @@ #define GPUCA_GPUCODE_HOSTONLY #define GPUCA_GPUCODE_NO_LAUNCH_BOUNDS -#define GPUCA_KRNL_REG(args) __launch_bounds__(GPUCA_M_STRIP(args)) +#define GPUCA_KRNL_REG(args) __launch_bounds__(GPUCA_M_MAX2_3(GPUCA_M_STRIP(args))) #include "GPUReconstructionCUDAIncludesSystem.h" #include "GPUReconstructionCUDADef.h" From d52a6fe2446bde4463421c9285aa776bc2c8fed8 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Fri, 18 Apr 2025 22:51:43 +0200 Subject: [PATCH 2/2] GPU: Fix parameter that was forgotten when moving to if constexpr with GPUCA_PAR_... --- GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h | 3 +++ GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx | 8 ++++---- 2 files changed, 7 insertions(+), 4 deletions(-) diff --git a/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h b/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h index 4ee6b23d46b51..57ad9907ca86f 100644 --- a/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h +++ b/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h @@ -574,6 +574,9 @@ #ifndef GPUCA_PAR_COMP_GATHER_MODE #define GPUCA_PAR_COMP_GATHER_MODE 0 #endif + #ifndef GPUCA_PAR_NO_ATOMIC_PRECHECK + #define GPUCA_PAR_NO_ATOMIC_PRECHECK 0 + #endif #ifndef GPUCA_PAR_DEDX_STORAGE_TYPE #define GPUCA_PAR_DEDX_STORAGE_TYPE float #endif diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx index 3b50bec45a41e..ddf01b586cd70 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx @@ -631,11 +631,11 @@ GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestric for (uint32_t ih = hitFst; ih < hitLst; ih++) { int32_t id = idOffset + ids[ih]; GPUAtomic(uint32_t)* const weight = weights + id; -#if GPUCA_NO_ATOMIC_PRECHECK == 0 - if (myWeight <= *weight) { - continue; + if constexpr (GPUCA_PAR_NO_ATOMIC_PRECHECK == 0) { + if (myWeight <= *weight) { + continue; + } } -#endif const cahit2 hh = CA_TEXTURE_FETCH(cahit2, gAliTexRefu2, hits, ih); const float y = y0 + hh.x * stepY; const float z = z0 + hh.y * stepZ;