offload/DeviceRTL/src/Misc.cpp

   1 //===--------- Misc.cpp - OpenMP device misc interfaces ----------- 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 "Allocator.h"
  13 #include "Configuration.h"
  14 #include "DeviceTypes.h"
  15 #include "Shared/RPCOpcodes.h"
  16 #include "shared/rpc.h"
  17
  18 #include "Debug.h"
  19
  20 #pragma omp begin declare target device_type(nohost)
  21
  22 namespace ompx {
  23 namespace impl {
  24
  25 double getWTick();
  26
  27 double getWTime();
  28
  29 /// AMDGCN Implementation
  30 ///
  31 ///{
  32 #pragma omp begin declare variant match(device = {arch(amdgcn)})
  33
  34 double getWTick() {
  35   // The number of ticks per second for the AMDGPU clock varies by card and can
  36   // only be retrived by querying the driver. We rely on the device environment
  37   // to inform us what the proper frequency is.
  38   return 1.0 / config::getClockFrequency();
  39 }
  40
  41 double getWTime() {
  42   uint64_t NumTicks = 0;
  43   if constexpr (__has_builtin(__builtin_amdgcn_s_sendmsg_rtnl))
  44     NumTicks = __builtin_amdgcn_s_sendmsg_rtnl(0x83);
  45   else if constexpr (__has_builtin(__builtin_amdgcn_s_memrealtime))
  46     NumTicks = __builtin_amdgcn_s_memrealtime();
  47   else if constexpr (__has_builtin(__builtin_amdgcn_s_memtime))
  48     NumTicks = __builtin_amdgcn_s_memtime();
  49
  50   return static_cast<double>(NumTicks) * getWTick();
  51 }
  52
  53 #pragma omp end declare variant
  54
  55 /// NVPTX Implementation
  56 ///
  57 ///{
  58 #pragma omp begin declare variant match(                                       \
  59         device = {arch(nvptx, nvptx64)},                                       \
  60             implementation = {extension(match_any)})
  61
  62 double getWTick() {
  63   // Timer precision is 1ns
  64   return ((double)1E-9);
  65 }
  66
  67 double getWTime() {
  68   uint64_t nsecs = __nvvm_read_ptx_sreg_globaltimer();
  69   return static_cast<double>(nsecs) * getWTick();
  70 }
  71
  72 #pragma omp end declare variant
  73
  74 /// Lookup a device-side function using a host pointer /p HstPtr using the table
  75 /// provided by the device plugin. The table is an ordered pair of host and
  76 /// device pointers sorted on the value of the host pointer.
  77 void *indirectCallLookup(void *HstPtr) {
  78   if (!HstPtr)
  79     return nullptr;
  80
  81   struct IndirectCallTable {
  82     void *HstPtr;
  83     void *DevPtr;
  84   };
  85   IndirectCallTable *Table =
  86       reinterpret_cast<IndirectCallTable *>(config::getIndirectCallTablePtr());
  87   uint64_t TableSize = config::getIndirectCallTableSize();
  88
  89   // If the table is empty we assume this is device pointer.
  90   if (!Table || !TableSize)
  91     return HstPtr;
  92
  93   uint32_t Left = 0;
  94   uint32_t Right = TableSize;
  95
  96   // If the pointer is definitely not contained in the table we exit early.
  97   if (HstPtr < Table[Left].HstPtr || HstPtr > Table[Right - 1].HstPtr)
  98     return HstPtr;
  99
 100   while (Left != Right) {
 101     uint32_t Current = Left + (Right - Left) / 2;
 102     if (Table[Current].HstPtr == HstPtr)
 103       return Table[Current].DevPtr;
 104
 105     if (HstPtr < Table[Current].HstPtr)
 106       Right = Current;
 107     else
 108       Left = Current;
 109   }
 110
 111   // If we searched the whole table and found nothing this is a device pointer.
 112   return HstPtr;
 113 }
 114
 115 /// The openmp client instance used to communicate with the server.
 116 /// FIXME: This is marked as 'retain' so that it is not removed via
 117 /// `-mlink-builtin-bitcode`
 118 [[gnu::visibility("protected"), gnu::weak,
 119   gnu::retain]] rpc::Client Client asm("__llvm_rpc_client");
 120
 121 } // namespace impl
 122 } // namespace ompx
 123
 124 /// Interfaces
 125 ///
 126 ///{
 127
 128 extern "C" {
 129 int32_t __kmpc_cancellationpoint(IdentTy *, int32_t, int32_t) { return 0; }
 130
 131 int32_t __kmpc_cancel(IdentTy *, int32_t, int32_t) { return 0; }
 132
 133 double omp_get_wtick(void) { return ompx::impl::getWTick(); }
 134
 135 double omp_get_wtime(void) { return ompx::impl::getWTime(); }
 136
 137 void *__llvm_omp_indirect_call_lookup(void *HstPtr) {
 138   return ompx::impl::indirectCallLookup(HstPtr);
 139 }
 140
 141 void *omp_alloc(size_t size, omp_allocator_handle_t allocator) {
 142   switch (allocator) {
 143   case omp_default_mem_alloc:
 144   case omp_large_cap_mem_alloc:
 145   case omp_const_mem_alloc:
 146   case omp_high_bw_mem_alloc:
 147   case omp_low_lat_mem_alloc:
 148     return malloc(size);
 149   default:
 150     return nullptr;
 151   }
 152 }
 153
 154 void omp_free(void *ptr, omp_allocator_handle_t allocator) {
 155   switch (allocator) {
 156   case omp_default_mem_alloc:
 157   case omp_large_cap_mem_alloc:
 158   case omp_const_mem_alloc:
 159   case omp_high_bw_mem_alloc:
 160   case omp_low_lat_mem_alloc:
 161     free(ptr);
 162   case omp_null_allocator:
 163   default:
 164     return;
 165   }
 166 }
 167
 168 unsigned long long __llvm_omp_host_call(void *fn, void *data, size_t size) {
 169   rpc::Client::Port Port = ompx::impl::Client.open<OFFLOAD_HOST_CALL>();
 170   Port.send_n(data, size);
 171   Port.send([=](rpc::Buffer *buffer, uint32_t) {
 172     buffer->data[0] = reinterpret_cast<uintptr_t>(fn);
 173   });
 174   unsigned long long Ret;
 175   Port.recv([&](rpc::Buffer *Buffer, uint32_t) {
 176     Ret = static_cast<unsigned long long>(Buffer->data[0]);
 177   });
 178   Port.close();
 179   return Ret;
 180 }
 181 }
 182
 183 ///}
 184 #pragma omp end declare target