GPU Sanitizer support for global variables

This commit extends GPU sanitizer support to global variables. Shadow variables that store the fake pointer values are added for each user-defined global. Right now, globals are registered by the host and de-allocated by the device, but this is expected to change in the future. The device-side allocation function is in the instrumentation pass, but it's commented out for now. Global variable references are rewritten as load instructions that get their respective fake pointer from the shadow globals.

Author: EthanLuisMcDonough

llvm/lib/Transforms/Instrumentation/GPUSan.cpp

@@ -79,7 +79,7 @@ template <AllocationKind AK> struct AllocationTracker {
 
   [[clang::disable_sanitizer_instrumentation]] static _AS_PTR(void, AK)
       create(_AS_PTR(void, AK) Start, uint64_t Length, int64_t AllocationId,
-             uint64_t Slot, int64_t SourceId, uint64_t PC) {
+             int64_t Slot, int64_t SourceId, uint64_t PC) {
     if constexpr (SanitizerConfig<AK>::OFFSET_BITS < 64)
       if (OMP_UNLIKELY(Length >= (1UL << (SanitizerConfig<AK>::OFFSET_BITS))))
         __sanitizer_trap_info_ptr->exceedsAllocationLength<AK>(
@@ -90,6 +90,8 @@ template <AllocationKind AK> struct AllocationTracker {
     auto &Cnt = AllocArr.Cnt;
     if constexpr (AK == AllocationKind::LOCAL)
       Slot = ++Cnt;
+    if (Slot == -1)
+      Slot = ++Cnt;
 
     uint64_t NumSlots = SanitizerConfig<AK>::SLOTS;
     if (OMP_UNLIKELY(Slot >= NumSlots))
@@ -246,10 +248,10 @@ extern "C" {
 [[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
   gnu::used, gnu::retain]] _AS_PTR(void, AllocationKind::GLOBAL)
     ompx_new_global(_AS_PTR(void, AllocationKind::GLOBAL) Start,
-                    uint64_t Length, int64_t AllocationId, uint32_t Slot,
-                    int64_t SourceId, uint64_t PC) {
+                    uint64_t Length, int64_t AllocationId, int64_t SourceId,
+                    uint64_t PC) {
   return AllocationTracker<AllocationKind::GLOBAL>::create(
-      Start, Length, AllocationId, Slot, SourceId, PC);
+      Start, Length, AllocationId, -1, SourceId, PC);
 }
 [[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
   gnu::used, gnu::retain]] void
@@ -260,13 +262,13 @@ __sanitizer_register_host(_AS_PTR(void, AllocationKind::GLOBAL) Start,
 }
 [[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
   gnu::used, gnu::retain]] void *
-ompx_new(void *Start, uint64_t Length, int64_t AllocationId, uint32_t Slot,
-         int64_t SourceId, uint64_t PC) {
+ompx_new(void *Start, uint64_t Length, int64_t AllocationId, int64_t SourceId,
+         uint64_t PC) {
   if (REAL_PTR_IS_LOCAL(Start))
     return (void *)ompx_new_local((_AS_PTR(void, AllocationKind::LOCAL))Start,
                                   Length, AllocationId, SourceId, PC);
   return (void *)ompx_new_global((_AS_PTR(void, AllocationKind::GLOBAL))Start,
-                                 Length, AllocationId, Slot, SourceId, PC);
+                                 Length, AllocationId, SourceId, PC);
 }
 [[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
   gnu::used, gnu::retain]] void

offload/DeviceRTL/src/Sanitizer.cpp

@@ -678,6 +678,8 @@ struct GPUSanTy {
   void addGPUSanNewFn(GenericKernelTy &GK) { NewFns.push_back(&GK); }
   void addGPUSanFreeFn(GenericKernelTy &GK) { FreeFns.push_back(&GK); }
   void checkAndReportError();
+  Error transferFakePtrToDevice(const char *GlobalName, void *FakeHstPtr,
+                                SmallVector<DeviceImageTy *> &Images);
 
 private:
   uint32_t SlotCnt = SanitizerConfig<AllocationKind::GLOBAL>::SLOTS - 1;
@@ -823,6 +825,13 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
     return PinnedAllocs.unlockUnmappedHostBuffer(HstPtr);
   }
 
+  /// Transfers a fake pointer to its respective shadow variable to prevent
+  /// double initializing GPUSan shadow constants. Only runs if GPUSan is
+  /// enabled
+  Error transferFakePtrToDevice(const char *GlobalName, void *FakeHstPtr) {
+    return GPUSan.transferFakePtrToDevice(GlobalName, FakeHstPtr, LoadedImages);
+  }
+
   /// Check whether the host buffer with address \p HstPtr is pinned by the
   /// underlying vendor-specific runtime (if any). Retrieve the host pointer,
   /// the device accessible pointer and the size of the original pinned buffer.

offload/plugins-nextgen/common/include/PluginInterface.h

@@ -2425,12 +2425,12 @@ void GPUSanTy::checkAndReportError() {
   case SanitizerTrapInfoTy::None:
     llvm_unreachable("Unexpected exception");
   case SanitizerTrapInfoTy::ExceedsLength:
-    fprintf(stderr, "%sERROR: OffloadSanitizer %s\n%s", Red(), "exceeds length",
-            Default());
+    fprintf(stderr, "%sERROR: OffloadSanitizer %s %d\n%s", Red(),
+            "exceeds length", STI->PtrSlot, Default());
     break;
   case SanitizerTrapInfoTy::ExceedsSlots:
-    fprintf(stderr, "%sERROR: OffloadSanitizer %s\n%s", Red(), "exceeds slots",
-            Default());
+    fprintf(stderr, "%sERROR: OffloadSanitizer %s %d\n%s", Red(),
+            "exceeds slots", STI->PtrSlot, Default());
     break;
   case SanitizerTrapInfoTy::PointerOutsideAllocation:
     fprintf(stderr, "%sERROR: OffloadSanitizer %s : %p : %i %lu (%s)\n%s",
@@ -2449,8 +2449,8 @@ void GPUSanTy::checkAndReportError() {
     DiagnoseAccess("use-after-free");
     break;
   case SanitizerTrapInfoTy::MemoryLeak:
-    fprintf(stderr, "%sERROR: OffloadSanitizer %s\n%s", Red(), "memory leak",
-            Default());
+    fprintf(stderr, "%sERROR: OffloadSanitizer memory leak at slot %d\n%s",
+            Red(), STI->PtrSlot, Default());
     break;
   case SanitizerTrapInfoTy::GarbagePointer:
     DiagnoseAccess("garbage-pointer");
@@ -2459,6 +2459,27 @@ void GPUSanTy::checkAndReportError() {
   fflush(stderr);
 }
 
+Error GPUSanTy::transferFakePtrToDevice(const char *GlobalName,
+                                        void *FakeHstPtr,
+                                        SmallVector<DeviceImageTy *> &Images) {
+  if (!FakeHstPtr)
+    return Plugin::success();
+
+  std::string ShadowName("__san.global.");
+  ShadowName.append(GlobalName);
+
+  GenericGlobalHandlerTy &GHandler = Device.Plugin.getGlobalHandler();
+  GlobalTy ShadowGlobal(ShadowName, sizeof(void *), &FakeHstPtr);
+
+  int imgCount = 0;
+  for (auto Img : Images) {
+    if (GHandler.isSymbolInImage(Device, *Img, ShadowName))
+      return GHandler.writeGlobalToDevice(Device, *Img, ShadowGlobal);
+  }
+
+  return Plugin::error("Shadow global for '%s' not found", GlobalName);
+}
+
 bool llvm::omp::target::plugin::libomptargetSupportsRPC() {
 #ifdef LIBOMPTARGET_RPC_SUPPORT
   return true;

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

@@ -450,6 +450,11 @@ static int loadImagesOntoDevice(DeviceTy &Device) {
                                     CurrHostEntry->size,
                                     Entry->FakeTgtPtrBegin))
           return OFFLOAD_FAIL;
+
+        auto &DeviceInterface = Device.RTL->getDevice(DeviceId);
+        if (DeviceInterface.transferFakePtrToDevice(CurrHostEntry->name,
+                                                    Entry->FakeTgtPtrBegin))
+          return OFFLOAD_FAIL;
       }
     }
     Device.setHasPendingImages(false);

offload/src/PluginManager.cpp

@@ -1,6 +1,6 @@
 // clang-format off
 // RUN: %libomptarget-compileopt-generic -fsanitize=offload
-// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: not %libomptarget-run-generic 2> %t.out
 // RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
 // clang-format on
 

offload/test/sanitizer/global_null.c

@@ -0,0 +1,27 @@
+// RUN: %libomptarget-compileopt-generic -g -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2> %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+// 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
+
+#include <omp.h>
+#include <stdio.h>
+
+int global_arr[3] = {1, 2, 3};
+#pragma omp declare target(global_arr)
+
+int main() {
+#pragma omp target
+  {
+    // CHECK: is located 12 bytes inside of a 12-byte region
+    for (int i = 0; i < 4; i++) {
+      global_arr[i] *= 4;
+    }
+  }
+  return 0;
+}

offload/test/sanitizer/global_variable/global_variable_array_iter_fail.c

@@ -0,0 +1,26 @@
+// RUN: %libomptarget-compileopt-generic -g -fsanitize=offload
+// RUN: %libomptarget-run-generic
+
+// 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
+
+#include <omp.h>
+#include <stdio.h>
+
+int global_arr[3] = {1, 2, 3};
+#pragma omp declare target(global_arr)
+
+int main() {
+#pragma omp target
+  {
+    for (int i = 0; i < 3; i++) {
+      global_arr[i] *= 4;
+    }
+    global_arr[1] = 22;
+  }
+  return 0;
+}

offload/test/sanitizer/global_variable/global_variable_array_pass.c

@@ -0,0 +1,25 @@
+// RUN: %libomptarget-compileopt-generic -g -fsanitize=offload
+// RUN: not %libomptarget-run-generic 2> %t.out
+// RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
+
+// 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
+
+#include <omp.h>
+#include <stdio.h>
+
+int global_arr[3] = {1, 2, 3};
+#pragma omp declare target(global_arr)
+
+int main() {
+#pragma omp target
+  {
+    // CHECK: is located 20 bytes inside of a 12-byte region
+    global_arr[5] = 27;
+  }
+  return 0;
+}

offload/test/sanitizer/global_variable/global_variable_array_static_fail.c

@@ -0,0 +1,25 @@
+// RUN: %libomptarget-compileopt-generic -g -fsanitize=offload
+// RUN: %libomptarget-run-generic
+
+// 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
+
+#include <omp.h>
+#include <stdio.h>
+
+int global1 = 0;
+int global2 = -1;
+#pragma omp declare target(global1, global2)
+
+int main() {
+#pragma omp target
+  {
+    global1 = 72;
+    global2 = global1 + 75;
+  }
+  return 0;
+}

offload/test/sanitizer/global_variable/global_variable_pass.c

@@ -1,6 +1,6 @@
 // clang-format off
 // RUN: %libomptarget-compileopt-generic -fsanitize=offload
-// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: not %libomptarget-run-generic 2> %t.out
 // RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
 // clang-format on
 

offload/test/sanitizer/heap_null.c

@@ -1,6 +1,6 @@
 // clang-format off
 // RUN: %libomptarget-compileopt-generic -fsanitize=offload
-// RUN: not %libomptarget-run-generic 2>&1 > %t.out
+// RUN: not %libomptarget-run-generic 2> %t.out
 // RUN: %fcheck-generic --check-prefixes=CHECK < %t.out
 // clang-format on
 
@@ -19,4 +19,5 @@ int main(void) {
   int *Random = &X;
 #pragma omp target
   { *Random = 99; }
+  // CHECK: 0x{{[a-f0-9]*}} is located {{[0-9]*}} bytes inside
 }