More tests

Author: jdoerfert

offload/include/Shared/Sanitizer.h

@@ -14,9 +14,15 @@
 #include "Types.h"
 #include "Utils.h"
 
-extern "C" int ompx_block_id(int Dim);
-extern "C" int ompx_block_dim(int Dim);
-extern "C" int ompx_thread_id(int Dim);
+extern "C" {
+int ompx_block_id(int Dim);
+int ompx_block_dim(int Dim);
+int ompx_thread_id(int Dim);
+[[clang::disable_sanitizer_instrumentation, gnu::noinline]] inline int64_t
+__san_get_location_value() {
+  return -1;
+}
+}
 
 enum class AllocationKind { LOCAL, GLOBAL, LAST = GLOBAL };
 
@@ -176,6 +182,7 @@ struct SanitizerTrapInfoTy {
   uint32_t ThreadId[3];
   uint64_t PC;
   uint64_t LocationId;
+  uint64_t CallId;
   /// }
 
   [[clang::disable_sanitizer_instrumentation]] void
@@ -185,6 +192,7 @@ struct SanitizerTrapInfoTy {
       ThreadId[Dim] = ompx_thread_id(Dim);
     }
     LocationId = SourceId;
+    CallId = __san_get_location_value();
   }
 
   template <enum AllocationKind AK>

offload/plugins-nextgen/common/src/PluginInterface.cpp

@@ -21,6 +21,7 @@
 #include "Utils/ELF.h"
 #include "omptarget.h"
 #include "llvm/Support/ErrorHandling.h"
+#include <cstdio>
 #include <string>
 
 #ifdef OMPT_SUPPORT
@@ -2235,9 +2236,7 @@ void GPUSanTy::checkAndReportError() {
       fprintf(stderr, "    no backtrace available\n");
       return;
     }
-    printf("Loc %p : %u\n", LocationsGV.getPtr(), LocationsGV.getSize());
-    printf("Nam %p : %u\n", LocationNamesGV.getPtr(),
-           LocationNamesGV.getSize());
+    fprintf(stderr, "%lu\n", STI->CallId);
     char *LocationNames = LocationNamesGV.getPtrAs<char>();
     LocationEncodingTy *Locations = LocationsGV.getPtrAs<LocationEncodingTy>();
     int32_t FrameIdx = 0;
@@ -2249,6 +2248,7 @@ void GPUSanTy::checkAndReportError() {
       LocationId = LE.ParentIdx;
       FrameIdx++;
     } while (LocationId >= 0);
+    fputc('\n', stderr);
   };
 
   auto DiagnoseAccess = [&](StringRef Name) {

offload/test/sanitizer/stack_trace_1.c

@@ -0,0 +1,43 @@
+// clang-format off
+// : %libomptarget-compileopt-generic -fsanitize=offload -O1
+// : not %libomptarget-run-generic 2> %t.out
+// : %fcheck-generic --check-prefixes=CHECK < %t.out
+// : %libomptarget-compileopt-generic -fsanitize=offload -O3
+// : not %libomptarget-run-generic 2> %t.out
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload -O3 -g
+// RUN: not %libomptarget-run-generic 2> %t.out
+// RUN: %fcheck-generic --check-prefixes=DEBUG < %t.out
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+// UNSUPPORTED: s390x-ibm-linux-gnu
+// UNSUPPORTED: s390x-ibm-linux-gnu-LTO
+
+[[clang::optnone]] int deref(int *P) { return *P; }
+
+[[gnu::always_inline]] int bar(int *P) { return deref(P); }
+
+int main(void) {
+
+#pragma omp target
+  {
+    int *NullPtr = 0;
+    // clang-format off
+    // CHECK: ERROR: OffloadSanitizer out-of-bounds access on address 0x0000000000000000 at pc [[PC:.*]]
+    // CHECK: WRITE of size 4 at 0x0000000000000000 thread <0, 0, 0> block <0, 0, 0> (acc 1, heap)
+    // CHECK:     #0 [[PC]] omp target (main:[[@LINE-6]]) in <unknown>:0
+    // 
+    // CHECK: 0x0000000000000000 is located 0 bytes inside of a 0-byte region [0x0000000000000000,0x0000000000000000)
+    //
+    // DEBUG: ERROR: OffloadSanitizer out-of-bounds access on address 0x0000000000000000 at pc [[PC:.*]]
+    // DEBUG: WRITE of size 4 at 0x0000000000000000 thread <0, 0, 0> block <0, 0, 0> (acc 1, heap)
+    // DEBUG:     #0 [[PC]] omp target (main:[[@LINE-12]]) in {{.*}}volatile_stack_null.c:[[@LINE+4]]
+    // 
+    // DEBUG: 0x0000000000000000 is located 0 bytes inside of a 0-byte region [0x0000000000000000,0x0000000000000000)
+    // clang-format on
+    bar(NullPtr);
+  }
+}

offload/test/sanitizer/stack_trace_multi_path_1.c

@@ -0,0 +1,47 @@
+// clang-format off
+// : %libomptarget-compileopt-generic -fsanitize=offload -O1
+// : not %libomptarget-run-generic 2> %t.out
+// : %fcheck-generic --check-prefixes=CHECK < %t.out
+// : %libomptarget-compileopt-generic -fsanitize=offload -O3
+// : not %libomptarget-run-generic 2> %t.out
+// RUN: %libomptarget-compileopt-generic -fsanitize=offload -O3 -g
+// RUN: not %libomptarget-run-generic 2> %t.out
+// RUN: %fcheck-generic --check-prefixes=DEBUG < %t.out
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+// UNSUPPORTED: s390x-ibm-linux-gnu
+// UNSUPPORTED: s390x-ibm-linux-gnu-LTO
+
+[[clang::optnone]] int deref(int *P) { return *P; }
+
+[[gnu::noinline]] int bar(int *P) { return deref(P); }
+[[gnu::noinline]] int baz(int *P) { return deref(P); }
+
+int main(void) {
+
+#pragma omp target
+  {
+    int *NullPtr = 0;
+    int X;
+    int *Valid = &X;
+    // clang-format off
+    // CHECK: ERROR: OffloadSanitizer out-of-bounds access on address 0x0000000000000000 at pc [[PC:.*]]
+    // CHECK: WRITE of size 4 at 0x0000000000000000 thread <0, 0, 0> block <0, 0, 0> (acc 1, heap)
+    // CHECK:     #0 [[PC]] omp target (main:[[@LINE-6]]) in <unknown>:0
+    // 
+    // CHECK: 0x0000000000000000 is located 0 bytes inside of a 0-byte region [0x0000000000000000,0x0000000000000000)
+    //
+    // DEBUG: ERROR: OffloadSanitizer out-of-bounds access on address 0x0000000000000000 at pc [[PC:.*]]
+    // DEBUG: WRITE of size 4 at 0x0000000000000000 thread <0, 0, 0> block <0, 0, 0> (acc 1, heap)
+    // DEBUG:     #0 [[PC]] omp target (main:[[@LINE-12]]) in {{.*}}volatile_stack_null.c:[[@LINE+4]]
+    // 
+    // DEBUG: 0x0000000000000000 is located 0 bytes inside of a 0-byte region [0x0000000000000000,0x0000000000000000)
+    // clang-format on
+    bar(Valid);
+    baz(NullPtr);
+  }
+}