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: 9f9fd859db]

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() {

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;