openmp/libomptarget/DeviceRTL/src/Mapping.cpp

   1 //===------- Mapping.cpp - OpenMP device runtime mapping helpers -- C++ -*-===//
   2 //
   3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
   4 // See https://llvm.org/LICENSE.txt for license information.
   5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
   6 //
   7 //===----------------------------------------------------------------------===//
   8 //
   9 //
  10 //===----------------------------------------------------------------------===//
  11
  12 #include "Mapping.h"
  13 #include "Interface.h"
  14 #include "State.h"
  15 #include "Types.h"
  16 #include "Utils.h"
  17
  18 #pragma omp declare target
  19
  20 #include "llvm/Frontend/OpenMP/OMPGridValues.h"
  21
  22 using namespace _OMP;
  23
  24 namespace _OMP {
  25 namespace impl {
  26
  27 /// AMDGCN Implementation
  28 ///
  29 ///{
  30 #pragma omp begin declare variant match(device = {arch(amdgcn)})
  31
  32 constexpr const llvm::omp::GV &getGridValue() {
  33   return llvm::omp::getAMDGPUGridValues<__AMDGCN_WAVEFRONT_SIZE>();
  34 }
  35
  36 uint32_t getGridDim(uint32_t n, uint16_t d) {
  37   uint32_t q = n / d;
  38   return q + (n > q * d);
  39 }
  40
  41 uint32_t getWorkgroupDim(uint32_t group_id, uint32_t grid_size,
  42                          uint16_t group_size) {
  43   uint32_t r = grid_size - group_id * group_size;
  44   return (r < group_size) ? r : group_size;
  45 }
  46
  47 uint32_t getNumHardwareThreadsInBlock() {
  48   return getWorkgroupDim(__builtin_amdgcn_workgroup_id_x(),
  49                          __builtin_amdgcn_grid_size_x(),
  50                          __builtin_amdgcn_workgroup_size_x());
  51 }
  52
  53 LaneMaskTy activemask() { return __builtin_amdgcn_read_exec(); }
  54
  55 LaneMaskTy lanemaskLT() {
  56   uint32_t Lane = mapping::getThreadIdInWarp();
  57   int64_t Ballot = mapping::activemask();
  58   uint64_t Mask = ((uint64_t)1 << Lane) - (uint64_t)1;
  59   return Mask & Ballot;
  60 }
  61
  62 LaneMaskTy lanemaskGT() {
  63   uint32_t Lane = mapping::getThreadIdInWarp();
  64   if (Lane == (mapping::getWarpSize() - 1))
  65     return 0;
  66   int64_t Ballot = mapping::activemask();
  67   uint64_t Mask = (~((uint64_t)0)) << (Lane + 1);
  68   return Mask & Ballot;
  69 }
  70
  71 uint32_t getThreadIdInWarp() {
  72   return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
  73 }
  74
  75 uint32_t getThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
  76
  77 uint32_t getKernelSize() { return __builtin_amdgcn_grid_size_x(); }
  78
  79 uint32_t getBlockId() { return __builtin_amdgcn_workgroup_id_x(); }
  80
  81 uint32_t getNumberOfBlocks() {
  82   return getGridDim(__builtin_amdgcn_grid_size_x(),
  83                     __builtin_amdgcn_workgroup_size_x());
  84 }
  85
  86 uint32_t getWarpId() {
  87   return impl::getThreadIdInBlock() / mapping::getWarpSize();
  88 }
  89
  90 uint32_t getNumberOfWarpsInBlock() {
  91   return mapping::getBlockSize() / mapping::getWarpSize();
  92 }
  93
  94 #pragma omp end declare variant
  95 ///}
  96
  97 /// NVPTX Implementation
  98 ///
  99 ///{
 100 #pragma omp begin declare variant match(                                       \
 101     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 102
 103 uint32_t getNumHardwareThreadsInBlock() {
 104   return __nvvm_read_ptx_sreg_ntid_x();
 105 }
 106
 107 constexpr const llvm::omp::GV &getGridValue() {
 108   return llvm::omp::NVPTXGridValues;
 109 }
 110
 111 LaneMaskTy activemask() {
 112   unsigned int Mask;
 113   asm("activemask.b32 %0;" : "=r"(Mask));
 114   return Mask;
 115 }
 116
 117 LaneMaskTy lanemaskLT() {
 118   __kmpc_impl_lanemask_t Res;
 119   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(Res));
 120   return Res;
 121 }
 122
 123 LaneMaskTy lanemaskGT() {
 124   __kmpc_impl_lanemask_t Res;
 125   asm("mov.u32 %0, %%lanemask_gt;" : "=r"(Res));
 126   return Res;
 127 }
 128
 129 uint32_t getThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }
 130
 131 uint32_t getThreadIdInWarp() {
 132   return impl::getThreadIdInBlock() & (mapping::getWarpSize() - 1);
 133 }
 134
 135 uint32_t getKernelSize() {
 136   return __nvvm_read_ptx_sreg_nctaid_x() *
 137          mapping::getNumberOfProcessorElements();
 138 }
 139
 140 uint32_t getBlockId() { return __nvvm_read_ptx_sreg_ctaid_x(); }
 141
 142 uint32_t getNumberOfBlocks() { return __nvvm_read_ptx_sreg_nctaid_x(); }
 143
 144 uint32_t getWarpId() {
 145   return impl::getThreadIdInBlock() / mapping::getWarpSize();
 146 }
 147
 148 uint32_t getNumberOfWarpsInBlock() {
 149   return (mapping::getBlockSize() + mapping::getWarpSize() - 1) /
 150          mapping::getWarpSize();
 151 }
 152
 153 #pragma omp end declare variant
 154 ///}
 155
 156 uint32_t getWarpSize() { return getGridValue().GV_Warp_Size; }
 157
 158 } // namespace impl
 159 } // namespace _OMP
 160
 161 /// We have to be deliberate about the distinction of `mapping::` and `impl::`
 162 /// below to avoid repeating assumptions or including irrelevant ones.
 163 ///{
 164
 165 static bool isInLastWarp() {
 166   uint32_t MainTId = (mapping::getNumberOfProcessorElements() - 1) &
 167                      ~(mapping::getWarpSize() - 1);
 168   return mapping::getThreadIdInBlock() == MainTId;
 169 }
 170
 171 bool mapping::isMainThreadInGenericMode(bool IsSPMD) {
 172   if (IsSPMD || icv::Level)
 173     return false;
 174
 175   // Check if this is the last warp in the block.
 176   return isInLastWarp();
 177 }
 178
 179 bool mapping::isMainThreadInGenericMode() {
 180   return mapping::isMainThreadInGenericMode(mapping::isSPMDMode());
 181 }
 182
 183 bool mapping::isInitialThreadInLevel0(bool IsSPMD) {
 184   if (IsSPMD)
 185     return mapping::getThreadIdInBlock() == 0;
 186   return isInLastWarp();
 187 }
 188
 189 bool mapping::isLeaderInWarp() {
 190   __kmpc_impl_lanemask_t Active = mapping::activemask();
 191   __kmpc_impl_lanemask_t LaneMaskLT = mapping::lanemaskLT();
 192   return utils::popc(Active & LaneMaskLT) == 0;
 193 }
 194
 195 LaneMaskTy mapping::activemask() { return impl::activemask(); }
 196
 197 LaneMaskTy mapping::lanemaskLT() { return impl::lanemaskLT(); }
 198
 199 LaneMaskTy mapping::lanemaskGT() { return impl::lanemaskGT(); }
 200
 201 uint32_t mapping::getThreadIdInWarp() {
 202   uint32_t ThreadIdInWarp = impl::getThreadIdInWarp();
 203   ASSERT(ThreadIdInWarp < impl::getWarpSize());
 204   return ThreadIdInWarp;
 205 }
 206
 207 uint32_t mapping::getThreadIdInBlock() {
 208   uint32_t ThreadIdInBlock = impl::getThreadIdInBlock();
 209   ASSERT(ThreadIdInBlock < impl::getNumHardwareThreadsInBlock());
 210   return ThreadIdInBlock;
 211 }
 212
 213 uint32_t mapping::getWarpSize() { return impl::getWarpSize(); }
 214
 215 uint32_t mapping::getBlockSize() {
 216   uint32_t BlockSize = mapping::getNumberOfProcessorElements() -
 217                        (!mapping::isSPMDMode() * impl::getWarpSize());
 218   return BlockSize;
 219 }
 220
 221 uint32_t mapping::getKernelSize() { return impl::getKernelSize(); }
 222
 223 uint32_t mapping::getWarpId() {
 224   uint32_t WarpID = impl::getWarpId();
 225   ASSERT(WarpID < impl::getNumberOfWarpsInBlock());
 226   return WarpID;
 227 }
 228
 229 uint32_t mapping::getBlockId() {
 230   uint32_t BlockId = impl::getBlockId();
 231   ASSERT(BlockId < impl::getNumberOfBlocks());
 232   return BlockId;
 233 }
 234
 235 uint32_t mapping::getNumberOfWarpsInBlock() {
 236   uint32_t NumberOfWarpsInBlocks = impl::getNumberOfWarpsInBlock();
 237   ASSERT(impl::getWarpId() < NumberOfWarpsInBlocks);
 238   return NumberOfWarpsInBlocks;
 239 }
 240
 241 uint32_t mapping::getNumberOfBlocks() {
 242   uint32_t NumberOfBlocks = impl::getNumberOfBlocks();
 243   ASSERT(impl::getBlockId() < NumberOfBlocks);
 244   return NumberOfBlocks;
 245 }
 246
 247 uint32_t mapping::getNumberOfProcessorElements() {
 248   uint32_t NumberOfProcessorElements = impl::getNumHardwareThreadsInBlock();
 249   ASSERT(impl::getThreadIdInBlock() < NumberOfProcessorElements);
 250   return NumberOfProcessorElements;
 251 }
 252
 253 ///}
 254
 255 /// Execution mode
 256 ///
 257 ///{
 258 static int SHARED(IsSPMDMode);
 259
 260 void mapping::init(bool IsSPMD) {
 261   if (mapping::isInitialThreadInLevel0(IsSPMD))
 262     IsSPMDMode = IsSPMD;
 263 }
 264
 265 bool mapping::isSPMDMode() { return IsSPMDMode; }
 266
 267 bool mapping::isGenericMode() { return !isSPMDMode(); }
 268 ///}
 269
 270 extern "C" {
 271 __attribute__((noinline)) uint32_t __kmpc_get_hardware_thread_id_in_block() {
 272   FunctionTracingRAII();
 273   return mapping::getThreadIdInBlock();
 274 }
 275
 276 __attribute__((noinline)) uint32_t __kmpc_get_hardware_num_threads_in_block() {
 277   FunctionTracingRAII();
 278   return impl::getNumHardwareThreadsInBlock();
 279 }
 280
 281 __attribute__((noinline)) uint32_t __kmpc_get_warp_size() {
 282   FunctionTracingRAII();
 283   return impl::getWarpSize();
 284 }
 285 }
 286 #pragma omp end declare target