[Openmp-commits] [PATCH] D75581: [libomptarget][amdgcn] Implement get_wtime

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield created this revision.
JonChesterfield added reviewers: jdoerfert, ABataev, grokos, ronlieb.
Herald added subscribers: openmp-commits, jvesely.
Herald added a project: OpenMP.

[libomptarget][amdgcn] Implement get_wtime

as a scaled call to an intrinsic. Implementation identical to that in aomp.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D75581

Files:
  openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip


Index: openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
===================================================================
--- openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
+++ openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
@@ -40,9 +40,13 @@
 
 DEVICE double __kmpc_impl_get_wtick() { return ((double)1E-9); }
 
-EXTERN uint64_t __clock64();
 DEVICE double __kmpc_impl_get_wtime() {
-  return ((double)1.0 / 745000000.0) * __clock64();
+#if __AMDGCN__ > 800
+  uint64_t t = __builtin_amdgcn_s_memrealtime();
+#else
+  uint64_t t = __builtin_amdgcn_s_memtime();
+#endif
+  return ((double)1.0 / 745000000.0) * t;
 }
 
 // Warp vote function


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D75581.248063.patch
Type: text/x-patch
Size: 674 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20200304/7747f146/attachment-0001.bin>