Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h
+++ openmp/trunk/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);                                                            \
     }                                                                          \
   }