[OpenMP] Introduce assumptions into the new device runtime mapping

[jdoerfert]

Summary

We know a lot about the mapping that we can encode as assumptions for
the optimizer. Our mapping vs impl layer makes this even easier.
Assumptions are checked as assertions in debug mode and otherwise
assumed (via llvm.assume).

[tianshilei1992]

LG

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

 }
 
 uint32_t getThreadIdInWarp() {
   return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
 }
 
 uint32_t getThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
 
-uint32_t getBlockSize() {
-  // TODO: verify this logic for generic mode.
-  return getWorkgroupDim(__builtin_amdgcn_workgroup_id_x(),
-                         __builtin_amdgcn_grid_size_x(),
-                         __builtin_amdgcn_workgroup_size_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 getNumberOfProcessorElements() {
-  // TODO
-  return mapping::getBlockSize();
+  // TODO: verify this logic for generic mode.
+  return getWorkgroupDim(__builtin_amdgcn_workgroup_id_x(),
+                         __builtin_amdgcn_grid_size_x(),
+                         __builtin_amdgcn_workgroup_size_x());
 }
 
 uint32_t getWarpId() {
-  return mapping::getThreadIdInBlock() / mapping::getWarpSize();
+  return impl::getThreadIdInBlock() / impl::getWarpSize();
 }
 
 uint32_t getNumberOfWarpsInBlock() {
-  return mapping::getBlockSize() / mapping::getWarpSize();
+  return impl::getBlockSize() / impl::getWarpSize();
 }
 
 #pragma omp end declare variant
 ///}
 
 /// NVPTX Implementation
 ///
 ///{

 
 LaneMaskTy lanemaskGT() {
   __kmpc_impl_lanemask_t Res;
   asm("mov.u32 %0, %%lanemask_gt;" : "=r"(Res));
   return Res;
 }
 
 uint32_t getThreadIdInWarp() {
-  return mapping::getThreadIdInBlock() & (mapping::getWarpSize() - 1);
+  return impl::getThreadIdInBlock() & (impl::getWarpSize() - 1);
 }
 
 uint32_t getThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }
 
-uint32_t getBlockSize() {
-  return __nvvm_read_ptx_sreg_ntid_x() -
-         (!mapping::isSPMDMode() * mapping::getWarpSize());
-}
-
 uint32_t getKernelSize() { return __nvvm_read_ptx_sreg_nctaid_x(); }
 
 uint32_t getBlockId() { return __nvvm_read_ptx_sreg_ctaid_x(); }
 
 uint32_t getNumberOfBlocks() { return __nvvm_read_ptx_sreg_nctaid_x(); }
 
 uint32_t getNumberOfProcessorElements() {
   return __nvvm_read_ptx_sreg_ntid_x();
 }
 
 uint32_t getWarpId() {
-  return mapping::getThreadIdInBlock() / mapping::getWarpSize();
+  return impl::getThreadIdInBlock() / impl::getWarpSize();
 }
 
 uint32_t getNumberOfWarpsInBlock() {
-  return (mapping::getBlockSize() + mapping::getWarpSize() - 1) /
-         mapping::getWarpSize();
+  return (impl::getBlockSize() + impl::getWarpSize() - 1) / impl::getWarpSize();
 }
 
 #pragma omp end declare variant
 ///}
 
 uint32_t getWarpSize() { return getGridValue().GV_Warp_Size; }
 
 } // namespace impl

 }
 
 LaneMaskTy mapping::activemask() { return impl::activemask(); }
 
 LaneMaskTy mapping::lanemaskLT() { return impl::lanemaskLT(); }
 
 LaneMaskTy mapping::lanemaskGT() { return impl::lanemaskGT(); }
 
-uint32_t mapping::getThreadIdInWarp() { return impl::getThreadIdInWarp(); }
+uint32_t mapping::getThreadIdInWarp() {
+  uint32_t ThreadIdInWarp = impl::getThreadIdInWarp();
+  ASSERT(ThreadIdInWarp < impl::getWarpSize());
+  return ThreadIdInWarp;
+}
+
+uint32_t mapping::getThreadIdInBlock() {
+  uint32_t ThreadIdInBlock = impl::getThreadIdInBlock();
+  ASSERT(ThreadIdInBlock < impl::getNumberOfProcessorElements());
+  return ThreadIdInBlock;
+}
 
-uint32_t mapping::getThreadIdInBlock() { return impl::getThreadIdInBlock(); }
+uint32_t mapping::getWarpSize() { return impl::getWarpSize(); }
 
 uint32_t mapping::getBlockSize() {
-  uint32_t BlockSize = impl::getBlockSize();
-  ASSERT(BlockSize == 1 | BlockSize % mapping::getWarpSize() == 0);
+  uint32_t BlockSize = mapping::getNumberOfProcessorElements() -
+                       (!impl::isSPMDMode() * impl::getWarpSize());
+  ASSERT(BlockSize == 1 | (BlockSize % impl::getWarpSize() == 0));
   return BlockSize;
 }
 
-uint32_t mapping::getKernelSize() { return impl::getKernelSize(); }
-
-uint32_t mapping::getBlockId() { return impl::getBlockId(); }
-
-uint32_t mapping::getNumberOfBlocks() { return impl::getNumberOfBlocks(); }
-
-uint32_t mapping::getNumberOfProcessorElements() {
-  return impl::getNumberOfProcessorElements();
+uint32_t mapping::getKernelSize() {
+  uint32_t KernelSize = impl::getKernelSize();
+  ASSERT(KernelSize == 1 | (KernelSize % impl::getWarpSize() == 0));
+  return KernelSize;
+}
+
+uint32_t mapping::getWarpId() {
+  uint32_t WarpID = impl::getWarpId();
+  ASSERT(WarpID < impl::getNumberOfWarpsInBlock());
+  return WarpID;
+}
+
+uint32_t mapping::getBlockId() {
+  uint32_t BlockId = impl::getBlockId();
+  ASSERT(BlockId < impl::getNumberOfBlocks());
+  return BlockId;
 }
 
-uint32_t mapping::getWarpId() { return impl::getWarpId(); }
+uint32_t mapping::getNumberOfWarpsInBlock() {
+  uint32_t NumberOfWarpsInBlocks = impl::getNumberOfWarpsInBlocks();
+  ASSERT(impl::getWarpId() < NumberOfWarpsInBlocks);
+  return NumberOfWarpsInBlocks;
+}
 
-uint32_t mapping::getWarpSize() { return impl::getWarpSize(); }
+uint32_t mapping::getNumberOfBlocks() {
+  uint32_t NumberOfBlocks = impl::getNumberOfBlocks();
+  ASSERT(impl::getBlockId() < NumberOfBlocks);
+  return NumberOfBlocks;
+}
 
-uint32_t mapping::getNumberOfWarpsInBlock() {
-  return impl::getNumberOfWarpsInBlock();
+uint32_t mapping::getNumberOfProcessorElements() {
+  uint32_t NumberOfProcessorElements = impl::getNumberOfProcessorElements();
+  ASSERT(impl::getThreadIdInBlock() < NumberOfProcessorElements);
+  return NumberOfProcessorElements();
 }
 
 /// Execution mode
 ///
 ///{
 static int SHARED(IsSPMDMode);
 
 void mapping::init(bool IsSPMD) {

openmp/libomptarget/DeviceRTL/src/State.cpp

 }
 
 int omp_get_team_size(int Level) {
   return returnValIfLevelIsActive(Level, state::ParallelTeamSize, 1);
 }
 
 int omp_get_num_threads(void) { return state::ParallelTeamSize; }
 
-int omp_get_thread_limit(void) { return mapping::getKernelSize(); }
+int omp_get_thread_limit(void) { return mapping::getBlockSize(); }
 
 int omp_get_num_procs(void) { return mapping::getNumberOfProcessorElements(); }
 
 void omp_set_nested(int) {}
 
 int omp_get_nested(void) { return false; }
 
 void omp_set_max_active_levels(int Levels) {