Skip to content

GPU: Fix parameter that was forgotten when moving to if constexpr with GPUCA_PAR_... #14203

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
davidrohr merged 2 commits into from
Apr 19, 2025
Merged

GPU: Fix parameter that was forgotten when moving to if constexpr with GPUCA_PAR_... #14203

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
3d4dfde
GPU RTC: Don't pass a third launch bounds parameter
davidrohr Apr 18, 2025
d52a6fe
GPU: Fix parameter that was forgotten when moving to if constexpr wit…
davidrohr Apr 18, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion GPU/GPUTracking/Base/cuda/GPUReconstructionCUDARTCCalls.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down
3 changes: 3 additions & 0 deletions GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
8 changes: 4 additions & 4 deletions GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down