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
19namespace ompx {
20namespace impl {
21
22double getWTick();
23
24double getWTime();
25
26/// AMDGCN Implementation
27///
28///{
29#pragma omp begin declare variant match(device = {arch(amdgcn)})
30
31double 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
38double 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
59double getWTick() {
60 // Timer precision is 1ns
61 return ((double)1E-9);
62}
63
64double 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.
74void *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
119extern "C" {
120int32_t __kmpc_cancellationpoint(IdentTy *, int32_t, int32_t) { return 0; }
121
122int32_t __kmpc_cancel(IdentTy *, int32_t, int32_t) { return 0; }
123
124double omp_get_wtick(void) { return ompx::impl::getWTick(); }
125
126double omp_get_wtime(void) { return ompx::impl::getWTime(); }
127
128void *__llvm_omp_indirect_call_lookup(void *HstPtr) {
129 return ompx::impl::indirectCallLookup(HstPtr);
130}
131}
132
133///}
134#pragma omp end declare target
135

source code of offload/DeviceRTL/src/Misc.cpp