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;