[OpenMP] Utilize the "non-uniform-workgroup" to simplify DeviceRTL

[jdoerfert]

Summary

OpenMP offloading always uses uniform workgroups, see
https://reviews.llvm.org/D135374. The runtime doesn't need to handle
non-uniform workgroups at all either.

[tianshilei1992]

LG

[estewart08]

This seems to have broken the amdgpu buildbot.
https://lab.llvm.org/buildbot/#/builders/193/builds/21152

[jdoerfert]

@estewart08 Can you try the suggested fix below?

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

82

This should be __builtin_amdgcn_grid_size_x()/__builtin_amdgcn_workgroup_size_x().

[estewart08]

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

82

That fixed the check-openmp failures on my local build.

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

 
 using namespace _OMP;
 
 namespace _OMP {
 namespace impl {
 
 // Forward declarations defined to be defined for AMDGCN and NVPTX.
 const llvm::omp::GV &getGridValue();
-uint32_t getGridDim(uint32_t n, uint16_t d);
-uint32_t getWorkgroupDim(uint32_t group_id, uint32_t grid_size,
-                         uint16_t group_size);
 uint32_t getNumHardwareThreadsInBlock();
 LaneMaskTy activemask();
 LaneMaskTy lanemaskLT();
 LaneMaskTy lanemaskGT();
 uint32_t getThreadIdInWarp();
 uint32_t getThreadIdInBlock();
 uint32_t getKernelSize();
 uint32_t getBlockId();
 uint32_t getNumberOfBlocks();
 uint32_t getWarpId();
 uint32_t getNumberOfWarpsInBlock();
 
 /// AMDGCN Implementation
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
 const llvm::omp::GV &getGridValue() {
   return llvm::omp::getAMDGPUGridValues<__AMDGCN_WAVEFRONT_SIZE>();
 }
 
-uint32_t getGridDim(uint32_t n, uint16_t d) {
-  uint32_t q = n / d;
-  return q + (n > q * d);
-}
-
-uint32_t getWorkgroupDim(uint32_t group_id, uint32_t grid_size,
-                         uint16_t group_size) {
-  uint32_t r = grid_size - group_id * group_size;
-  return (r < group_size) ? r : group_size;
-}
-
 uint32_t getNumHardwareThreadsInBlock() {
-  return getWorkgroupDim(__builtin_amdgcn_workgroup_id_x(),
-                         __builtin_amdgcn_grid_size_x(),
-                         __builtin_amdgcn_workgroup_size_x());
+  return __builtin_amdgcn_workgroup_size_x();
 }
 
 LaneMaskTy activemask() { return __builtin_amdgcn_read_exec(); }
 
 LaneMaskTy lanemaskLT() {
   uint32_t Lane = mapping::getThreadIdInWarp();
   int64_t Ballot = mapping::activemask();
   uint64_t Mask = ((uint64_t)1 << Lane) - (uint64_t)1;

 }
 
 uint32_t getThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
 
 uint32_t getKernelSize() { return __builtin_amdgcn_grid_size_x(); }
 
 uint32_t getBlockId() { return __builtin_amdgcn_workgroup_id_x(); }
 
-uint32_t getNumberOfBlocks() {
-  return getGridDim(__builtin_amdgcn_grid_size_x(),
-                    __builtin_amdgcn_workgroup_size_x());
-}
+uint32_t getNumberOfBlocks() { return __builtin_amdgcn_grid_size_x(); }
 
 uint32_t getWarpId() {
   return impl::getThreadIdInBlock() / mapping::getWarpSize();
 }
 
 uint32_t getNumberOfWarpsInBlock() {
   return mapping::getBlockSize() / mapping::getWarpSize();
 }
uint32_t mapping::getThreadIdInBlock() {
   uint32_t ThreadIdInBlock = impl::getThreadIdInBlock();
   ASSERT(ThreadIdInBlock < impl::getNumHardwareThreadsInBlock());
   return ThreadIdInBlock;
 }
 
 uint32_t mapping::getWarpSize() { return impl::getWarpSize(); }
 
 uint32_t mapping::getBlockSize(bool IsSPMD) {
-  uint32_t BlockSize = mapping::getNumberOfProcessorElements() -
-                       (!IsSPMD * impl::getWarpSize());
+  uint32_t BlockSize =
+      mapping::getNumberOfProcessorElements() - (!IsSPMD * impl::getWarpSize());
   return BlockSize;
 }
 uint32_t mapping::getBlockSize() {
   return mapping::getBlockSize(mapping::isSPMDMode());
 }
 
 uint32_t mapping::getKernelSize() { return impl::getKernelSize(); }