From 64b9f0a1920cba6d8655e352d36ade221c1f6293 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Fri, 8 Jul 2022 21:34:21 -0400 Subject: [PATCH] SWDEV-286739 - Support hipDeviceAttributeWallClockRate Part 3: Add missing declaration of wall_clock64() to fix compiling issue in device code. Add querying hipDeviceAttributeWallClockRate. Change-Id: Ie54771c2f58eeaacdc0248bc116ef193f99eb9b9 [ROCm/clr commit: 9f9fd859db46d3edb921683f88a13b5356f7741c] --- .../clr/hipamd/include/hip/amd_detail/amd_device_functions.h | 5 +++-- projects/clr/hipamd/src/hip_device_runtime.cpp | 3 +++ 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h index 22654b8ccb..96353b0f54 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -689,6 +689,7 @@ __device__ long long int __clock64(); __device__ long long int __clock(); __device__ long long int clock64(); __device__ long long int clock(); +__device__ long long int wall_clock64(); // hip.amdgcn.bc - named sync __device__ void __named_sync(); @@ -713,8 +714,8 @@ __device__ inline __attribute((always_inline)) long long int __clock() { return __clock64(); } -// Clock function to return wall clock count at a constant frequency. The interface to query -// the frequency will be implemented. +// Clock function to return wall clock count at a constant frequency that can be queried +// through hipDeviceAttributeWallClockRate attribute. __device__ inline __attribute__((always_inline)) long long int wall_clock64() { diff --git a/projects/clr/hipamd/src/hip_device_runtime.cpp b/projects/clr/hipamd/src/hip_device_runtime.cpp index aada90b8f5..89e9e69ece 100644 --- a/projects/clr/hipamd/src/hip_device_runtime.cpp +++ b/projects/clr/hipamd/src/hip_device_runtime.cpp @@ -195,6 +195,9 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) case hipDeviceAttributeClockRate: *pi = prop.clockRate; break; + case hipDeviceAttributeWallClockRate: + *pi = g_devices[device]->devices()[0]->info().wallClockFrequency_; + break; case hipDeviceAttributeMemoryClockRate: *pi = prop.memoryClockRate; break;