offload/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 begin declare target device_type(nohost)
  19
  20 using namespace ompx;
  21
  22 namespace impl {
  23
  24 bool isSharedMemPtr(const void *Ptr) { return false; }
  25
  26 void Unpack(uint64_t Val, uint32_t *LowBits, uint32_t *HighBits) {
  27   static_assert(sizeof(unsigned long) == 8, "");
  28   *LowBits = static_cast<uint32_t>(Val & 0x00000000FFFFFFFFUL);
  29   *HighBits = static_cast<uint32_t>((Val & 0xFFFFFFFF00000000UL) >> 32);
  30 }
  31
  32 uint64_t Pack(uint32_t LowBits, uint32_t HighBits) {
  33   return (((uint64_t)HighBits) << 32) | (uint64_t)LowBits;
  34 }
  35
  36 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane);
  37 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
  38                     int32_t Width);
  39
  40 uint64_t ballotSync(uint64_t Mask, int32_t Pred);
  41
  42 /// AMDGCN Implementation
  43 ///
  44 ///{
  45 #pragma omp begin declare variant match(device = {arch(amdgcn)})
  46
  47 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
  48   int Width = mapping::getWarpSize();
  49   int Self = mapping::getThreadIdInWarp();
  50   int Index = SrcLane + (Self & ~(Width - 1));
  51   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
  52 }
  53
  54 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
  55                     int32_t Width) {
  56   int Self = mapping::getThreadIdInWarp();
  57   int Index = Self + LaneDelta;
  58   Index = (int)(LaneDelta + (Self & (Width - 1))) >= Width ? Self : Index;
  59   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
  60 }
  61
  62 uint64_t ballotSync(uint64_t Mask, int32_t Pred) {
  63   return Mask & __builtin_amdgcn_ballot_w64(Pred);
  64 }
  65
  66 bool isSharedMemPtr(const void *Ptr) {
  67   return __builtin_amdgcn_is_shared(
  68       (const __attribute__((address_space(0))) void *)Ptr);
  69 }
  70 #pragma omp end declare variant
  71 ///}
  72
  73 /// NVPTX Implementation
  74 ///
  75 ///{
  76 #pragma omp begin declare variant match(                                       \
  77         device = {arch(nvptx, nvptx64)},                                       \
  78             implementation = {extension(match_any)})
  79
  80 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
  81   return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, 0x1f);
  82 }
  83
  84 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta, int32_t Width) {
  85   int32_t T = ((mapping::getWarpSize() - Width) << 8) | 0x1f;
  86   return __nvvm_shfl_sync_down_i32(Mask, Var, Delta, T);
  87 }
  88
  89 uint64_t ballotSync(uint64_t Mask, int32_t Pred) {
  90   return __nvvm_vote_ballot_sync(static_cast<uint32_t>(Mask), Pred);
  91 }
  92
  93 bool isSharedMemPtr(const void *Ptr) { return __nvvm_isspacep_shared(Ptr); }
  94
  95 #pragma omp end declare variant
  96 ///}
  97 } // namespace impl
  98
  99 uint64_t utils::pack(uint32_t LowBits, uint32_t HighBits) {
 100   return impl::Pack(LowBits, HighBits);
 101 }
 102
 103 void utils::unpack(uint64_t Val, uint32_t &LowBits, uint32_t &HighBits) {
 104   impl::Unpack(Val, &LowBits, &HighBits);
 105 }
 106
 107 int32_t utils::shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
 108   return impl::shuffle(Mask, Var, SrcLane);
 109 }
 110
 111 int32_t utils::shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta,
 112                            int32_t Width) {
 113   return impl::shuffleDown(Mask, Var, Delta, Width);
 114 }
 115
 116 int64_t utils::shuffleDown(uint64_t Mask, int64_t Var, uint32_t Delta,
 117                            int32_t Width) {
 118   uint32_t Lo, Hi;
 119   utils::unpack(Var, Lo, Hi);
 120   Hi = impl::shuffleDown(Mask, Hi, Delta, Width);
 121   Lo = impl::shuffleDown(Mask, Lo, Delta, Width);
 122   return utils::pack(Lo, Hi);
 123 }
 124
 125 uint64_t utils::ballotSync(uint64_t Mask, int32_t Pred) {
 126   return impl::ballotSync(Mask, Pred);
 127 }
 128
 129 bool utils::isSharedMemPtr(void *Ptr) { return impl::isSharedMemPtr(Ptr); }
 130
 131 extern "C" {
 132 int32_t __kmpc_shuffle_int32(int32_t Val, int16_t Delta, int16_t SrcLane) {
 133   return impl::shuffleDown(lanes::All, Val, Delta, SrcLane);
 134 }
 135
 136 int64_t __kmpc_shuffle_int64(int64_t Val, int16_t Delta, int16_t Width) {
 137   return utils::shuffleDown(lanes::All, Val, Delta, Width);
 138 }
 139 }
 140
 141 #pragma omp end declare target