[Openmp-commits] [openmp] 85c2d92 - [openmp][AMDGPU] - Correct getNumberOfBlocks calculation.

[Ethan Stewart via Openmp-commits]

Author: Ethan Stewart
Date: 2022-11-02T11:38:56-05:00
New Revision: 85c2d92b9b3217229ced6a0f22e7013f7a65e124

URL: https://github.com/llvm/llvm-project/commit/85c2d92b9b3217229ced6a0f22e7013f7a65e124
DIFF: https://github.com/llvm/llvm-project/commit/85c2d92b9b3217229ced6a0f22e7013f7a65e124.diff

LOG: [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

Added: 
    

Modified: 
    openmp/libomptarget/DeviceRTL/src/Mapping.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
index 7cef92d304b66..512577c06f9eb 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();