diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index b54cf16613..e382ede4b3 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -420,17 +420,7 @@ HIP provides the following built-in functions for reading a high-resolution time clock_t clock() long long int clock64() ``` - -AMD devices employ a per-GPU timer that increments at a constant time interval regardless of any dynamic frequency changes. All compute units in the system share the timer. -Nvidia devices implement the timer as a per-compute-unit clock that increments on every clock cycle. - -To obtain the clock frequency, use the hipDeviceProp_t.clockInstructionRate field: - -``` -hipGetDeviceProperties(&deviceProps, deviceId); -// Compute time in ms--device_ticks is based on values reported from clock() device function -float time = device_ticks / (float)deviceProps.clockInstructionRate; -``` +Returns the value of counter that is incremented every clock cycle on device. Difference in values returned provides the cycles used. ## Atomic Functions diff --git a/src/device_util.cpp b/src/device_util.cpp index 2e9e9e4540..3234408e50 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -501,8 +501,8 @@ __device__ double trunc(double x) const int warpSize = 64; -__device__ long long int clock64() { return (long long int)hc::__clock_u64(); }; -__device__ clock_t clock() { return (clock_t)hc::__clock_u64(); }; +__device__ long long int clock64() { return (long long int)hc::__cycle_u64(); }; +__device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); }; //atomicAdd()