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 "Configuration.h"
  13 #include "Types.h"
  14
  15 #include "Debug.h"
  16
  17 #pragma omp begin declare target device_type(nohost)
  18
  19 namespace ompx {
  20 namespace impl {
  21
  22 double getWTick();
  23
  24 double getWTime();
  25
  26 /// AMDGCN Implementation
  27 ///
  28 ///{
  29 #pragma omp begin declare variant match(device = {arch(amdgcn)})
  30
  31 double getWTick() {
  32   // The number of ticks per second for the AMDGPU clock varies by card and can
  33   // only be retrived by querying the driver. We rely on the device environment
  34   // to inform us what the proper frequency is.
  35   return 1.0 / config::getClockFrequency();
  36 }
  37
  38 double getWTime() {
  39   uint64_t NumTicks = 0;
  40   if constexpr (__has_builtin(__builtin_amdgcn_s_sendmsg_rtnl))
  41     NumTicks = __builtin_amdgcn_s_sendmsg_rtnl(0x83);
  42   else if constexpr (__has_builtin(__builtin_amdgcn_s_memrealtime))
  43     NumTicks = __builtin_amdgcn_s_memrealtime();
  44   else if constexpr (__has_builtin(__builtin_amdgcn_s_memtime))
  45     NumTicks = __builtin_amdgcn_s_memtime();
  46
  47   return static_cast<double>(NumTicks) * getWTick();
  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)},                                       \
  57             implementation = {extension(match_any)})
  58
  59 double getWTick() {
  60   // Timer precision is 1ns
  61   return ((double)1E-9);
  62 }
  63
  64 double getWTime() {
  65   uint64_t nsecs = __nvvm_read_ptx_sreg_globaltimer();
  66   return static_cast<double>(nsecs) * getWTick();
  67 }
  68
  69 #pragma omp end declare variant
  70
  71 /// Lookup a device-side function using a host pointer /p HstPtr using the table
  72 /// provided by the device plugin. The table is an ordered pair of host and
  73 /// device pointers sorted on the value of the host pointer.
  74 void *indirectCallLookup(void *HstPtr) {
  75   if (!HstPtr)
  76     return nullptr;
  77
  78   struct IndirectCallTable {
  79     void *HstPtr;
  80     void *DevPtr;
  81   };
  82   IndirectCallTable *Table =
  83       reinterpret_cast<IndirectCallTable *>(config::getIndirectCallTablePtr());
  84   uint64_t TableSize = config::getIndirectCallTableSize();
  85
  86   // If the table is empty we assume this is device pointer.
  87   if (!Table || !TableSize)
  88     return HstPtr;
  89
  90   uint32_t Left = 0;
  91   uint32_t Right = TableSize;
  92
  93   // If the pointer is definitely not contained in the table we exit early.
  94   if (HstPtr < Table[Left].HstPtr || HstPtr > Table[Right - 1].HstPtr)
  95     return HstPtr;
  96
  97   while (Left != Right) {
  98     uint32_t Current = Left + (Right - Left) / 2;
  99     if (Table[Current].HstPtr == HstPtr)
 100       return Table[Current].DevPtr;
 101
 102     if (HstPtr < Table[Current].HstPtr)
 103       Right = Current;
 104     else
 105       Left = Current;
 106   }
 107
 108   // If we searched the whole table and found nothing this is a device pointer.
 109   return HstPtr;
 110 }
 111
 112 } // namespace impl
 113 } // namespace ompx
 114
 115 /// Interfaces
 116 ///
 117 ///{
 118
 119 extern "C" {
 120 int32_t __kmpc_cancellationpoint(IdentTy *, int32_t, int32_t) { return 0; }
 121
 122 int32_t __kmpc_cancel(IdentTy *, int32_t, int32_t) { return 0; }
 123
 124 double omp_get_wtick(void) { return ompx::impl::getWTick(); }
 125
 126 double omp_get_wtime(void) { return ompx::impl::getWTime(); }
 127
 128 void *__llvm_omp_indirect_call_lookup(void *HstPtr) {
 129   return ompx::impl::indirectCallLookup(HstPtr);
 130 }
 131 }
 132
 133 ///}
 134 #pragma omp end declare target