[Openmp-commits] [openmp] r350128 - [OPENMP][NVPTX]Outline assert into noinline function, NFC.

[Alexey Bataev via Openmp-commits]

Author: abataev
Date: Fri Dec 28 09:29:47 2018
New Revision: 350128

URL: http://llvm.org/viewvc/llvm-project?rev=350128&view=rev
Log:
[OPENMP][NVPTX]Outline assert into noinline function, NFC.

Summary:
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.

Reviewers: gtbercea, kkwli0, grokos

Subscribers: guansong, caomhin, openmp-commits

Differential Revision: https://reviews.llvm.org/D56101

Modified:
    openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h?rev=350128&r1=350127&r2=350128&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h Fri Dec 28 09:29:47 2018
@@ -138,6 +138,18 @@ static NOINLINE void log(const char *fmt
 #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 @@ static NOINLINE void log(const char *fmt
 #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 @@ static NOINLINE void log(const char *fmt
 #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);                                                            \
     }                                                                          \
   }