openmp/libomptarget/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   unsigned long long nsecs;
  66   asm volatile("mov.u64  %0, %%globaltimer;" : "=l"(nsecs));
  67   return (double)nsecs * getWTick();
  68 }
  69
  70 #pragma omp end declare variant
  71
  72 /// Lookup a device-side function using a host pointer /p HstPtr using the table
  73 /// provided by the device plugin. The table is an ordered pair of host and
  74 /// device pointers sorted on the value of the host pointer.
  75 void *indirectCallLookup(void *HstPtr) {
  76   if (!HstPtr)
  77     return nullptr;
  78
  79   struct IndirectCallTable {
  80     void *HstPtr;
  81     void *DevPtr;
  82   };
  83   IndirectCallTable *Table =
  84       reinterpret_cast<IndirectCallTable *>(config::getIndirectCallTablePtr());
  85   uint64_t TableSize = config::getIndirectCallTableSize();
  86
  87   // If the table is empty we assume this is device pointer.
  88   if (!Table || !TableSize)
  89     return HstPtr;
  90
  91   uint32_t Left = 0;
  92   uint32_t Right = TableSize;
  93
  94   // If the pointer is definitely not contained in the table we exit early.
  95   if (HstPtr < Table[Left].HstPtr || HstPtr > Table[Right - 1].HstPtr)
  96     return HstPtr;
  97
  98   while (Left != Right) {
  99     uint32_t Current = Left + (Right - Left) / 2;
 100     if (Table[Current].HstPtr == HstPtr)
 101       return Table[Current].DevPtr;
 102
 103     if (HstPtr < Table[Current].HstPtr)
 104       Right = Current;
 105     else
 106       Left = Current;
 107   }
 108
 109   // If we searched the whole table and found nothing this is a device pointer.
 110   return HstPtr;
 111 }
 112
 113 } // namespace impl
 114 } // namespace ompx
 115
 116 /// Interfaces
 117 ///
 118 ///{
 119
 120 extern "C" {
 121 int32_t __kmpc_cancellationpoint(IdentTy *, int32_t, int32_t) { return 0; }
 122
 123 int32_t __kmpc_cancel(IdentTy *, int32_t, int32_t) { return 0; }
 124
 125 double omp_get_wtick(void) { return ompx::impl::getWTick(); }
 126
 127 double omp_get_wtime(void) { return ompx::impl::getWTime(); }
 128
 129 void *__llvm_omp_indirect_call_lookup(void *HstPtr) {
 130   return ompx::impl::indirectCallLookup(HstPtr);
 131 }
 132 }
 133
 134 ///}
 135 #pragma omp end declare target