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" 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;