[Libomptarget] Correctly implement `getWTime` on AMDGPU

[jhuber6]

Summary

AMDGPU provides a fixed frequency clock since some generations back.
However, the frequency is variable by card and must be looked up at
runtime. This patch adds a new device environment line for the clock
frequency so that we can use it in the same way as NVPTX. This is the
correct implementation and the version in ASO should be replaced.

[tianshilei1992]

openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp

833

This doesn't need to be 1000000000UL?

[jhuber6]

openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp

833

A 32-bit integer fits at least two billion, so we're just under here.

[tianshilei1992]

Do we have a test?

[jhuber6]

Add test

[tianshilei1992]

openmp/libomptarget/test/offloading/wtime.c

25

Probably a more reliable way is to check the diff between duration and 0.0. as long as it is greater than epsilon we are good.

[jhuber6]

Updating test

[tianshilei1992]

openmp/libomptarget/test/offloading/wtime.c

25

The check will fail if the time is 0.01 right? You can potentially just do compile and run since you put an assertion there.

[jhuber6]

Update test, I think the host and device miught use different deltas so it failed sometimes.

[tianshilei1992]

LGTM, suppose the AMD intrinsics are correct. :-D

[jhuber6]

In D154456#4472519, @tianshilei1992 wrote:

LGTM, suppose the AMD intrinsics are correct. :-D

It's taken from the device libs, and when I tested myself I got a number that was in-line with what I got on the host, and I tested it on two cards with different frequencies.

openmp/libomptarget/DeviceRTL/include/Configuration.h

 uint32_t getDeviceNum();
 
 /// Return the user choosen debug level.
 uint32_t getDebugKind();
 
 /// Return the amount of dynamic shared memory that was allocated at launch.
 uint64_t getDynamicMemorySize();
 
+/// Returns the cycles per second of the device's fixed frequency clock.
+uint64_t getClockFrequency();
+
 /// Return if debugging is enabled for the given debug kind.
 bool isDebugMode(DebugKind Level);
 
 /// Indicates if this kernel may require thread-specific states, or if it was
 /// explicitly disabled by the user.
 bool mayUseThreadStates();
 
 /// Indicates if this kernel may require data environments for nested
 /// parallelism, or if it was explicitly disabled by the user.
 bool mayUseNestedParallelism();
 
 } // namespace config
 } // namespace ompx
 
 #endif

openmp/libomptarget/DeviceRTL/src/Configuration.cpp

 uint32_t config::getDeviceNum() {
   return __omp_rtl_device_environment.DeviceNum;
 }
 
 uint64_t config::getDynamicMemorySize() {
   return __omp_rtl_device_environment.DynamicMemSize;
 }
 
+uint64_t config::getClockFrequency() {
+  return __omp_rtl_device_environment.ClockFrequency;
+}
+
 bool config::isDebugMode(config::DebugKind Kind) {
   return config::getDebugKind() & Kind;
 }
 
 bool config::mayUseThreadStates() { return !__omp_rtl_assume_no_thread_state; }
 
 bool config::mayUseNestedParallelism() {
   return !__omp_rtl_assume_no_nested_parallelism;
 }
 
 #pragma omp end declare target

openmp/libomptarget/DeviceRTL/src/Misc.cpp

 //===--------- Misc.cpp - OpenMP device misc interfaces ----------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 //
 //===----------------------------------------------------------------------===//
 
+#include "Configuration.h"
 #include "Types.h"
 
 #include "Debug.h"
 
 #pragma omp begin declare target device_type(nohost)
 
 namespace ompx {
 namespace impl {
 
 double getWTick();
 
 double getWTime();
 
 /// AMDGCN Implementation
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
-double getWTick() { return ((double)1E-9); }
+double getWTick() {
+  // The number of ticks per second for the AMDGPU clock varies by card and can
+  // only be retrived by querying the driver. We rely on the device environment
+  // to inform us what the proper frequency is.
+  return 1.0 / config::getClockFrequency();
+}
 
 double getWTime() {
-  // The intrinsics for measuring time have undocumented frequency
-  // This will probably need to be found by measurement on a number of
-  // architectures. Until then, return 0, which is very inaccurate as a
-  // timer but resolves the undefined symbol at link time.
-  return 0;
+  uint64_t NumTicks = 0;
+  if constexpr (__has_builtin(__builtin_amdgcn_s_sendmsg_rtnl))
+    NumTicks = __builtin_amdgcn_s_sendmsg_rtnl(0x83);
+  else if constexpr (__has_builtin(__builtin_amdgcn_s_memrealtime))
+    NumTicks = __builtin_amdgcn_s_memrealtime();
+  else if constexpr (__has_builtin(__builtin_amdgcn_s_memtime))
+    NumTicks = __builtin_amdgcn_s_memtime();
+
+  return static_cast<double>(NumTicks) * getWTick();
 }
 
 #pragma omp end declare variant
 
 /// NVPTX Implementation
 ///
 ///{
 #pragma omp begin declare variant match(                                       \

openmp/libomptarget/include/DeviceEnvironment.h

 
 // deviceRTL uses <stdint> and DeviceRTL uses explicit definitions
 
 struct DeviceEnvironmentTy {
   uint32_t DebugKind;
   uint32_t NumDevices;
   uint32_t DeviceNum;
   uint32_t DynamicMemSize;
+  uint64_t ClockFrequency;
 };
 
 #endif

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

Error initImpl(GenericPluginTy &Plugin) override {
     ComputeUnitKind = GPUName;
 
     // Get the wavefront size.
     uint32_t WavefrontSize = 0;
     if (auto Err = getDeviceAttr(HSA_AGENT_INFO_WAVEFRONT_SIZE, WavefrontSize))
       return Err;
     GridValues.GV_Warp_Size = WavefrontSize;
 
+    // Get the frequency of the steady clock.
+    if (auto Err = getDeviceAttr(HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY,
+                                 ClockFrequency))
+      return Err;
+
     // Load the grid values dependending on the wavefront.
     if (WavefrontSize == 32)
       GridValues = getAMDGPUGridValues<32>();
     else if (WavefrontSize == 64)
       GridValues = getAMDGPUGridValues<64>();
     else
       return Plugin::error("Unexpected AMDGPU wavefront %d", WavefrontSize);
 
doJITPostProcessing(std::unique_ptr<MemoryBuffer> MB) const override {
 
     return std::move(
         MemoryBuffer::getFileOrSTDIN(LinkerOutputFilePath.data()).get());
   }
 
   /// See GenericDeviceTy::getComputeUnitKind().
   std::string getComputeUnitKind() const override { return ComputeUnitKind; }
 
+  /// Returns the clock frequency for the given AMDGPU device.
+  uint64_t getClockFrequency() const override { return ClockFrequency; }
+
   /// Allocate and construct an AMDGPU kernel.
   Expected<GenericKernelTy *>
   constructKernelEntry(const __tgt_offload_entry &KernelEntry,
                        DeviceImageTy &Image) override {
 
     Expected<OMPTgtExecModeFlags> ExecModeOrErr =
         getExecutionModeForKernel(KernelEntry.name, Image);
     if (!ExecModeOrErr)
private:
   AMDGPUSignalManagerTy AMDGPUSignalManager;
 
   /// The agent handler corresponding to the device.
   hsa_agent_t Agent;
 
   /// The GPU architecture.
   std::string ComputeUnitKind;
 
+  /// The frequency of the steady clock inside the device.
+  uint64_t ClockFrequency;
+
   /// Reference to the host device.
   AMDHostDeviceTy &HostDevice;
 
   /// List of device packet queues.
   std::vector<AMDGPUQueueTy> Queues;
 };
 
 Error AMDGPUDeviceImageTy::loadExecutable(const AMDGPUDeviceTy &Device) {

openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.h

struct GenericDeviceTy : public DeviceAllocatorTy {
   uint32_t getBlockLimit() const { return GridValues.GV_Max_Teams; }
   uint32_t getDefaultNumThreads() const {
     return GridValues.GV_Default_WG_Size;
   }
   uint32_t getDefaultNumBlocks() const {
     return GridValues.GV_Default_Num_Teams;
   }
   uint32_t getDynamicMemorySize() const { return OMPX_SharedMemorySize; }
+  virtual uint64_t getClockFrequency() const { return CLOCKS_PER_SEC; }
 
   /// Get target compute unit kind (e.g., sm_80, or gfx908).
   virtual std::string getComputeUnitKind() const { return "unknown"; }
 
   /// Post processing after jit backend. The ownership of \p MB will be taken.
   virtual Expected<std::unique_ptr<MemoryBuffer>>
   doJITPostProcessing(std::unique_ptr<MemoryBuffer> MB) const {
     return std::move(MB);

openmp/libomptarget/plugins-nextgen/common/PluginInterface/PluginInterface.cpp

if (!shouldSetupDeviceEnvironment())
     return Plugin::success();
 
   DeviceEnvironmentTy DeviceEnvironment;
   DeviceEnvironment.DebugKind = OMPX_DebugKind;
   DeviceEnvironment.NumDevices = Plugin.getNumDevices();
   // TODO: The device ID used here is not the real device ID used by OpenMP.
   DeviceEnvironment.DeviceNum = DeviceId;
   DeviceEnvironment.DynamicMemSize = OMPX_SharedMemorySize;
+  DeviceEnvironment.ClockFrequency = getClockFrequency();
 
   // Create the metainfo of the device environment global.
   GlobalTy DevEnvGlobal("__omp_rtl_device_environment",
                         sizeof(DeviceEnvironmentTy), &DeviceEnvironment);
 
   // Write device environment values to the device.
   GenericGlobalHandlerTy &GHandler = Plugin.getGlobalHandler();
   if (auto Err = GHandler.writeGlobalToDevice(*this, Image, DevEnvGlobal)) {

openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp

CUresult getDeviceAttrRaw(uint32_t Kind, int &Value) {
     return cuDeviceGetAttribute(&Value, (CUdevice_attribute)Kind, Device);
   }
 
   /// See GenericDeviceTy::getComputeUnitKind().
   std::string getComputeUnitKind() const override {
     return ComputeCapability.str();
   }
 
+  /// Returns the clock frequency for the given NVPTX device.
+  uint64_t getClockFrequency() const override { return 1000000000; }
+
 private:
   using CUDAStreamManagerTy = GenericDeviceResourceManagerTy<CUDAStreamRef>;
   using CUDAEventManagerTy = GenericDeviceResourceManagerTy<CUDAEventRef>;
 
   /// Stream manager for CUDA streams.
   CUDAStreamManagerTy CUDAStreamManager;
 
   /// Event manager for CUDA events.

openmp/libomptarget/test/offloading/wtime.c

-// RUN: %libomptarget-compileopt-run-and-check-generic
-
-// UNSUPPORTED: amdgcn-amd-amdhsa
+// RUN: %libomptarget-compileopt-and-run-generic
 
+#include <assert.h>
 #include <omp.h>
 #include <stdio.h>
 #include <stdlib.h>
 
 #define N (1024 * 1024 * 256)
 
 int main(int argc, char *argv[]) {
   int *data = (int *)malloc(N * sizeof(int));
-#pragma omp target map(from : data[0 : N])
+  double duration = 0.0;
+
+#pragma omp target map(from : data[0 : N]) map(from : duration)
   {
     double start = omp_get_wtime();
     for (int i = 0; i < N; ++i)
       data[i] = i;
     double end = omp_get_wtime();
-    double duration = end - start;
-    printf("duration: %lfs\n", duration);
+    duration = end - start;
   }
+  assert(duration > 0.0);
   free(data);
   return 0;
 }
-
-// CHECK: duration: {{.+[1-9]+}}