[Openmp-commits] [PATCH] D56101: [OPENMP][NVPTX]Outline assert into noinline function, NFC.

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev created this revision.
ABataev added reviewers: gtbercea, kkwli0, grokos.
Herald added a subscriber: guansong.

At high optimization level asserts lead to some unexpected results
because of auto-inserted unreachable instructions. This outlining
prevents some of such dangerous optimizations and leads to better
stability.


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D56101

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


Index: libomptarget/deviceRTLs/nvptx/src/debug.h
===================================================================
--- libomptarget/deviceRTLs/nvptx/src/debug.h
+++ libomptarget/deviceRTLs/nvptx/src/debug.h
@@ -138,6 +138,18 @@
 #endif
 #if OMPTARGET_NVPTX_TEST
 #include <assert.h>
+
+template <typename... Arguments>
+NOINLINE static void check(bool cond, const char *fmt,
+                           Arguments... parameters) {
+  if (!cond)
+    printf(fmt, (int)blockIdx.x, (int)threadIdx.x,
+           (int)(threadIdx.x / WARPSIZE), (int)(threadIdx.x & 0x1F),
+           parameters...);
+  assert(cond);
+}
+
+NOINLINE static void check(bool cond) { assert(cond); }
 #endif
 
 // set flags that are tested (inclusion properties)
@@ -207,13 +219,13 @@
 #define ASSERT0(_flag, _cond, _str)                                            \
   {                                                                            \
     if (TON(_flag)) {                                                          \
-      assert(_cond);                                                           \
+      check(_cond);                                                            \
     }                                                                          \
   }
 #define ASSERT(_flag, _cond, _str, _args...)                                   \
   {                                                                            \
     if (TON(_flag)) {                                                          \
-      assert(_cond);                                                           \
+      check(_cond);                                                            \
     }                                                                          \
   }
 
@@ -222,16 +234,15 @@
 #define TON(_flag) ((OMPTARGET_NVPTX_TEST) & (_flag))
 #define ASSERT0(_flag, _cond, _str)                                            \
   {                                                                            \
-    if (TON(_flag) && !(_cond)) {                                              \
-      log("<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n");                  \
-      assert(_cond);                                                           \
+    if (TON(_flag)) {                                                          \
+      check((_cond), "<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n");       \
     }                                                                          \
   }
 #define ASSERT(_flag, _cond, _str, _args...)                                   \
   {                                                                            \
-    if (TON(_flag) && !(_cond)) {                                              \
-      log("<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n", _args);           \
-      assert(_cond);                                                           \
+    if (TON(_flag)) {                                                          \
+      check((_cond), "<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n",        \
+            _args);                                                            \
     }                                                                          \
   }
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D56101.179550.patch
Type: text/x-patch
Size: 3239 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20181227/26bc35be/attachment.bin>