Skip to content

Commit 2bafa1f

Browse files
davidrohrktf
authored andcommitted
GPU: Use launch_bounds also for HIP instead of attribute(num_vgpr)
1 parent ba279cd commit 2bafa1f

File tree

2 files changed

+2
-2
lines changed

2 files changed

+2
-2
lines changed

GPU/Common/GPUDefGPUParameters.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@
2727
#define GPUCA_THREAD_COUNT_CONSTRUCTOR 128
2828
#define GPUCA_THREAD_COUNT_SELECTOR 128
2929
#define GPUCA_THREAD_COUNT_FINDER 128
30-
#define GPUCA_NEIGHBORSFINDER_REGS REG, 64
30+
#define GPUCA_NEIGHBORSFINDER_REGS REG, (GPUCA_THREAD_COUNT_FINDER, 1)
3131
#define GPUCA_NEIGHBOURS_FINDER_MAX_NNEIGHUP 0
3232
#define GPUCA_TRACKLET_SELECTOR_HITS_REG_SIZE 12
3333
#elif defined(GPUCA_GPUTYPE_TURING)

GPU/GPUTracking/Base/hip/GPUReconstructionHIP.hip.cxx

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,7 @@ GPUg() void runKernelHIP(GPUCA_CONSMEM_PTR int iSlice, Args... args)
7272
*/
7373

7474
#undef GPUCA_KRNL_REG
75-
#define GPUCA_KRNL_REG(num) __attribute__((amdgpu_num_vgpr(num)))
75+
#define GPUCA_KRNL_REG(args) __launch_bounds__(GPUCA_M_STRIP(args))
7676
#define GPUCA_KRNL(x_class, x_attributes, x_arguments, x_forward) GPUCA_KRNL_WRAP(GPUCA_KRNL_, x_class, x_attributes, x_arguments, x_forward)
7777
#define GPUCA_KRNL_BACKEND_CLASS GPUReconstructionHIPBackend
7878
#define GPUCA_KRNL_CALL_single(x_class, x_attributes, x_arguments, x_forward) \

0 commit comments

Comments
 (0)