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

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield updated this revision to Diff 226574.
JonChesterfield added a comment.

- Use functions from support.h instead of adding to target_impl


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D69476/new/

https://reviews.llvm.org/D69476

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


Index: openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
===================================================================
--- openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
+++ openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
@@ -128,12 +128,13 @@
 
 #if OMPTARGET_NVPTX_DEBUG || OMPTARGET_NVPTX_TEST || OMPTARGET_NVPTX_WARNING
 #include <stdio.h>
-#include "target_impl.h"
+#include "support.h"
 
 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)(GetWarpId()),
+         (int)(GetLaneId()), 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)(GetWarpId()),
+           (int)(GetLaneId()), parameters...);
   assert(cond);
 }
 


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