[Openmp-commits] [PATCH] D69476: [nfc][libomptarget] Warp size aware logging. Last part of D69423

[Jon Chesterfield via Phabricator via Openmp-commits]

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

[nfc][libomptarget] Warp size aware logging. Last part of D69423 <https://reviews.llvm.org/D69423>

Derives a mask from warp size, behind a target_impl shim. Deduplicates some
logic and facilitates making an intrinsic call instead of using theadIdx.x.

Replaces x & (sz -1) with x % sz for clarity. LLVM rewrites to & at -O1.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D69476

Files:
  openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
  openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h


Index: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
===================================================================
--- openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
@@ -147,4 +147,10 @@
 #endif // CUDA_VERSION
 }
 
+INLINE uint32_t __kmpc_impl_get_warp_id() { return threadIdx.x / WARPSIZE; }
+
+INLINE uint32_t __kmpc_impl_get_thread_id_in_warp() {
+  return threadIdx.x % WARPSIZE;
+}
+
 #endif
Index: openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
===================================================================
--- openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
+++ openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
@@ -132,8 +132,9 @@
 
 template <typename... Arguments>
 NOINLINE static void log(const char *fmt, Arguments... parameters) {
-  printf(fmt, (int)blockIdx.x, (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE),
-         (int)(threadIdx.x & 0x1F), parameters...);
+  printf(fmt, (int)blockIdx.x, (int)threadIdx.x,
+         (int)(__kmpc_impl_get_warp_id()),
+         (int)(__kmpc_impl_get_thread_id_in_warp()), parameters...);
 }
 
 #endif
@@ -145,8 +146,8 @@
                            Arguments... parameters) {
   if (!cond)
     printf(fmt, (int)blockIdx.x, (int)threadIdx.x,
-           (int)(threadIdx.x / WARPSIZE), (int)(threadIdx.x & 0x1F),
-           parameters...);
+           (int)(__kmpc_impl_get_warp_id()),
+           (int)(__kmpc_impl_get_thread_id_in_warp()), parameters...);
   assert(cond);
 }
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D69476.226572.patch
Type: text/x-patch
Size: 1526 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20191027/6fdef5c3/attachment.bin>