From 1f74a66bea87ea42da6dbf18c0ddad649cfefbf3 Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
Date: Tue, 28 Aug 2018 19:30:20 -0400
Subject: [PATCH] Add clock() and clock64()

---
 include/hip/hcc_detail/device_functions.h | 37 ++++++++++--
 tests/src/deviceLib/hipTestClock.cpp      | 69 +++++++++++++++++++++++
 2 files changed, 101 insertions(+), 5 deletions(-)
 create mode 100644 tests/src/deviceLib/hipTestClock.cpp

diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h
index 5e1c8a4e6b..ebf9de8b4a 100644
--- a/include/hip/hcc_detail/device_functions.h
+++ b/include/hip/hcc_detail/device_functions.h
@@ -611,16 +611,43 @@ __device__ static inline float __ull2float_rz(unsigned long long int x) { return
 
 #ifdef __HCC_OR_HIP_CLANG__
 
+// Clock functions
+__device__ long long int __clock64();
+__device__ long long int __clock();
+__device__ long long int clock64();
+__device__ long long int clock();
+// hip.amdgcn.bc - named sync
+__device__ void __named_sync(int a, int b);
+
 #ifdef __HIP_DEVICE_COMPILE__
 
 // Clock functions
-__device__
-inline
-long long int __clock64() { return (long long int) __builtin_amdgcn_s_memrealtime(); }
+#if __HCC__
+extern "C" uint64_t __clock_u64()  __HC__;
+#endif
 
 __device__
-inline
-long long int  __clock() { return (long long int) __builtin_amdgcn_s_memrealtime(); }
+inline  __attribute((always_inline))
+long long int __clock64() {
+// ToDo: Unify HCC and HIP implementation.
+#if __HCC__
+  return (long long int) __clock_u64();
+#else
+  return (long long int) __builtin_amdgcn_s_memrealtime();
+#endif
+}
+
+__device__
+inline __attribute((always_inline))
+long long int  __clock() { return __clock64(); }
+
+__device__
+inline  __attribute__((always_inline))
+long long int clock64() { return __clock64(); }
+
+__device__
+inline __attribute__((always_inline))
+long long int  clock() { return __clock(); }
 
 // hip.amdgcn.bc - named sync
 __device__
diff --git a/tests/src/deviceLib/hipTestClock.cpp b/tests/src/deviceLib/hipTestClock.cpp
new file mode 100644
index 0000000000..46f64e35a3
--- /dev/null
+++ b/tests/src/deviceLib/hipTestClock.cpp
@@ -0,0 +1,69 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include <hip/hip_runtime.h>
+#include <hip/hip_runtime_api.h>
+#include <iostream>
+#include "test_common.h"
+
+#define HIP_ASSERT(status) assert(status == hipSuccess)
+
+#define LEN 512
+#define SIZE 2048
+
+struct TestClock {
+     
+  static __global__ void kernel1(int* Ad) {
+      int tid = threadIdx.x + blockIdx.x * blockDim.x;
+      Ad[tid] = clock() + clock64() + __clock() + __clock64();
+  }
+
+  static __global__ void kernel2(int* Ad) {
+      int tid = threadIdx.x + blockIdx.x * blockDim.x;
+      Ad[tid] = clock() + clock64() + __clock() + __clock64() - Ad[tid];
+  }
+
+  void run() {
+    int *A, *Ad;
+    A = new int[LEN];
+    for (unsigned i = 0; i < LEN; i++) {
+        A[i] = 0;
+    }
+
+    HIP_ASSERT(hipMalloc((void**)&Ad, SIZE));
+    hipLaunchKernelGGL(kernel1, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
+    hipLaunchKernelGGL(kernel2, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
+    HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
+
+    for (unsigned i = 0; i < LEN; i++) {
+        assert(0 != A[i]);
+    }
+  }
+};
+
+int main() {
+  TestClock().run();
+  passed();
+}