From 85c2d92b9b3217229ced6a0f22e7013f7a65e124 Mon Sep 17 00:00:00 2001 From: Ethan Stewart Date: Wed, 2 Nov 2022 11:37:42 -0500 Subject: [PATCH] [openmp][AMDGPU] - Correct getNumberOfBlocks calculation. This patch fixes the 6 amdgpu buildbot lit test failures introduced by https://reviews.llvm.org/D135444. libomptarget :: amdgcn-amd-amdhsa :: mapping/reduction_implicit_map.cpp libomptarget :: amdgcn-amd-amdhsa :: offloading/cuda_no_devices.c libomptarget :: amdgcn-amd-amdhsa :: offloading/target-teams-atomic.c libomptarget :: amdgcn-amd-amdhsa-LTO :: mapping/reduction_implicit_map.cpp libomptarget :: amdgcn-amd-amdhsa-LTO :: offloading/cuda_no_devices.c libomptarget :: amdgcn-amd-amdhsa-LTO :: offloading/target-teams-atomic.c Reviewed By: jhuber6 Differential Revision: https://reviews.llvm.org/D137261 --- openmp/libomptarget/DeviceRTL/src/Mapping.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp index 7cef92d..512577c 100644 --- a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp +++ b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp @@ -79,7 +79,9 @@ uint32_t getKernelSize() { return __builtin_amdgcn_grid_size_x(); } uint32_t getBlockId() { return __builtin_amdgcn_workgroup_id_x(); } -uint32_t getNumberOfBlocks() { return __builtin_amdgcn_grid_size_x(); } +uint32_t getNumberOfBlocks() { + return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x(); +} uint32_t getWarpId() { return impl::getThreadIdInBlock() / mapping::getWarpSize(); -- 2.7.4