openmp/libomptarget/DeviceRTL/src/Utils.cpp

   1 //===------- Utils.cpp - OpenMP device runtime utility functions -- 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 "Utils.h"
  13
  14 #include "Debug.h"
  15 #include "Interface.h"
  16 #include "Mapping.h"
  17
  18 #pragma omp declare target
  19
  20 using namespace _OMP;
  21
  22 namespace _OMP {
  23 /// Helper to keep code alive without introducing a performance penalty.
  24 __attribute__((used, retain, weak, optnone, cold)) void keepAlive() {
  25   __kmpc_get_hardware_thread_id_in_block();
  26   __kmpc_get_hardware_num_threads_in_block();
  27   __kmpc_get_warp_size();
  28   __kmpc_barrier_simple_spmd(nullptr, 0);
  29   __kmpc_barrier_simple_generic(nullptr, 0);
  30 }
  31 } // namespace _OMP
  32
  33 namespace impl {
  34
  35 /// AMDGCN Implementation
  36 ///
  37 ///{
  38 #pragma omp begin declare variant match(device = {arch(amdgcn)})
  39
  40 void Unpack(uint64_t Val, uint32_t *LowBits, uint32_t *HighBits) {
  41   static_assert(sizeof(unsigned long) == 8, "");
  42   *LowBits = (uint32_t)(Val & 0x00000000FFFFFFFFUL);
  43   *HighBits = (uint32_t)((Val & 0xFFFFFFFF00000000UL) >> 32);
  44 }
  45
  46 uint64_t Pack(uint32_t LowBits, uint32_t HighBits) {
  47   return (((uint64_t)HighBits) << 32) | (uint64_t)LowBits;
  48 }
  49
  50 #pragma omp end declare variant
  51
  52 /// NVPTX Implementation
  53 ///
  54 ///{
  55 #pragma omp begin declare variant match(                                       \
  56     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
  57
  58 void Unpack(uint64_t Val, uint32_t *LowBits, uint32_t *HighBits) {
  59   uint32_t LowBitsLocal, HighBitsLocal;
  60   asm("mov.b64 {%0,%1}, %2;"
  61       : "=r"(LowBitsLocal), "=r"(HighBitsLocal)
  62       : "l"(Val));
  63   *LowBits = LowBitsLocal;
  64   *HighBits = HighBitsLocal;
  65 }
  66
  67 uint64_t Pack(uint32_t LowBits, uint32_t HighBits) {
  68   uint64_t Val;
  69   asm("mov.b64 %0, {%1,%2};" : "=l"(Val) : "r"(LowBits), "r"(HighBits));
  70   return Val;
  71 }
  72
  73 #pragma omp end declare variant
  74
  75 /// AMDGCN Implementation
  76 ///
  77 ///{
  78 #pragma omp begin declare variant match(device = {arch(amdgcn)})
  79
  80 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
  81   int Width = mapping::getWarpSize();
  82   int Self = mapping::getThreadIdInWarp();
  83   int Index = SrcLane + (Self & ~(Width - 1));
  84   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
  85 }
  86
  87 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
  88                     int32_t Width) {
  89   int Self = mapping::getThreadIdInWarp();
  90   int Index = Self + LaneDelta;
  91   Index = (int)(LaneDelta + (Self & (Width - 1))) >= Width ? Self : Index;
  92   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
  93 }
  94
  95 #pragma omp end declare variant
  96 ///}
  97
  98 /// NVPTX Implementation
  99 ///
 100 ///{
 101 #pragma omp begin declare variant match(                                       \
 102     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 103
 104 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
 105   return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, 0x1f);
 106 }
 107
 108 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta, int32_t Width) {
 109   int32_t T = ((mapping::getWarpSize() - Width) << 8) | 0x1f;
 110   return __nvvm_shfl_sync_down_i32(Mask, Var, Delta, T);
 111 }
 112
 113 #pragma omp end declare variant
 114 } // namespace impl
 115
 116 uint64_t utils::pack(uint32_t LowBits, uint32_t HighBits) {
 117   return impl::Pack(LowBits, HighBits);
 118 }
 119
 120 void utils::unpack(uint64_t Val, uint32_t &LowBits, uint32_t &HighBits) {
 121   impl::Unpack(Val, &LowBits, &HighBits);
 122 }
 123
 124 int32_t utils::shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
 125   return impl::shuffle(Mask, Var, SrcLane);
 126 }
 127
 128 int32_t utils::shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta,
 129                            int32_t Width) {
 130   return impl::shuffleDown(Mask, Var, Delta, Width);
 131 }
 132
 133 extern "C" {
 134 int32_t __kmpc_shuffle_int32(int32_t Val, int16_t Delta, int16_t SrcLane) {
 135   FunctionTracingRAII();
 136   return impl::shuffleDown(lanes::All, Val, Delta, SrcLane);
 137 }
 138
 139 int64_t __kmpc_shuffle_int64(int64_t Val, int16_t Delta, int16_t Width) {
 140   FunctionTracingRAII();
 141   uint32_t lo, hi;
 142   utils::unpack(Val, lo, hi);
 143   hi = impl::shuffleDown(lanes::All, hi, Delta, Width);
 144   lo = impl::shuffleDown(lanes::All, lo, Delta, Width);
 145   return utils::pack(lo, hi);
 146 }
 147 }
 148
 149 #pragma omp end declare target