[libomptarget][amdgcn] Implement get_wtime

[JonChesterfield]

Summary

[libomptarget][amdgcn] Implement get_wtime

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

[JonChesterfield]

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

50

I don't have a source for the magic number. It's from before my time. I can offer that rough checks from running applications look ok, and that it has been shipping as part of aomp for ages without user complaints.

[grokos]

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

50

At least can you add a comment about what this magic number is (obviously the clock speed in Hz, but it would be nice to have some sort of documentation)? Or (even better) use a macro because this frequency may change in future chips?

[JonChesterfield]

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

50

The isa docs offer that the clock operates as if at a fixed frequency. It seems plausible that the various amdgcn chips would all have a timer running at the same rate, but I'd also be happier with a reference on that.

I'll ask around. It would be useful to know whether the scaling factor should be different for different chips.

[JonChesterfield]

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

50

Yes. Though I will first ask around internally to try to find out what the 745MHz was based on, and whether it's chip dependent.

[JonChesterfield]

I don't have a documentation reference for this yet. A crude benchmark suggests the frequency is 25MHz on a vega 20, but I'm hoping for a source other than experimentation.

[JonChesterfield]

I'm unable to find a good reference for what this frequency should be, so am considering unconditionally returning zero from this function. That won't 'work' in any useful sense, but it's an improvement on failing to link because __clock64 is undefined.

[jdoerfert]

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

50

My money is on: NVIDIA Tesla K40 Graphic Card - 1 GPUs - 745 MHz Core

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

if (lane == (WARPSIZE - 1))
     return 0;
   uint64_t ballot = __kmpc_impl_activemask();
   uint64_t mask = (~((uint64_t)0)) << (lane + 1);
   return mask & ballot;
 }
 
 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
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
   // 33 is ICMP_NE from llvm/include/llvm/IR/InstrTypes.h
   return __builtin_amdgcn_uicmp(1, 0, 33);
 }
 
 DEVICE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t var,