Repair the rebase

Author: jdoerfert

offload/plugins-nextgen/amdgpu/src/rtl.cpp

@@ -1601,7 +1601,6 @@ struct AMDGPUStreamManagerTy final
   using ResourceRef = AMDGPUResourceRef<AMDGPUStreamTy>;
   using ResourcePoolTy = GenericDeviceResourceManagerTy<ResourceRef>;
 
-  GenericDeviceTy &Device;
   AMDGPUStreamManagerTy(GenericDeviceTy &Device, hsa_agent_t HSAAgent)
       : GenericDeviceResourceManagerTy(Device), Device(Device),
         OMPX_QueueTracking("LIBOMPTARGET_AMDGPU_HSA_QUEUE_BUSY_TRACKING", true),
@@ -3528,6 +3527,9 @@ void AMDGPUQueueTy::callbackError(hsa_status_t Status, hsa_queue_t *Source,
                                   void *Data) {
   auto &AMDGPUDevice = *reinterpret_cast<AMDGPUDeviceTy *>(Data);
 
+  auto *Device = reinterpret_cast<AMDGPUDeviceTy *>(Data);
+  Device->reportSanitizerError();
+
   if (Status == HSA_STATUS_ERROR_EXCEPTION) {
     auto KernelTraceInfoRecord =
         AMDGPUDevice.KernelLaunchTraces.getExclusiveAccessor();
@@ -3542,9 +3544,6 @@ void AMDGPUQueueTy::callbackError(hsa_status_t Status, hsa_queue_t *Source,
                                       AsyncInfoWrapperMatcher);
   }
 
-  auto *Device = reinterpret_cast<AMDGPUDeviceTy *>(Data);
-  Device->reportSanitizerError();
-
   auto Err = Plugin::check(Status, "Received error in queue %p: %s", Source);
   FATAL_MESSAGE(1, "%s", toString(std::move(Err)).data());
 }

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

@@ -987,7 +987,7 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
     auto AllocationTraceMap = AllocationTraces.getExclusiveAccessor();
     for (auto &It : *AllocationTraceMap) {
       if (It.first <= DevicePtr &&
-          advanceVoidPtr(It.first, It.second->Size) > DevicePtr)
+          utils::advancePtr(It.first, It.second->Size) > DevicePtr)
         return It.second;
     }
     return nullptr;

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

@@ -17,6 +17,7 @@
 #include "ErrorReporting.h"
 #include "GlobalHandler.h"
 #include "JIT.h"
+#include "Shared/EnvironmentVar.h"
 #include "Shared/Sanitizer.h"
 #include "Shared/Utils.h"
 #include "Utils/ELF.h"
@@ -939,24 +940,33 @@ GenericDeviceTy::loadBinary(GenericPluginTy &Plugin,
   if (auto Err = callGlobalConstructors(Plugin, *Image))
     return std::move(Err);
 
-  auto GetKernel = [&](StringRef Name) -> GenericKernelTy * {
-    auto KernelOrErr = constructKernel(Name.data());
-    if (Error Err = KernelOrErr.takeError()) {
-      REPORT("Failure to look up kernel: %s\n",
-             toString(std::move(Err)).data());
-      return nullptr;
+  auto *&SanitizerTrapInfo = SanitizerTrapInfos[Image];
+  GlobalTy TrapId("__sanitizer_trap_info_ptr", sizeof(SanitizerTrapInfo),
+                  &SanitizerTrapInfo);
+  GenericGlobalHandlerTy &GHandler = Plugin.getGlobalHandler();
+  if (GHandler.isSymbolInImage(*this, *Image, TrapId.getName())) {
+
+    SanitizerTrapInfo = reinterpret_cast<SanitizerTrapInfoTy *>(allocate(
+        sizeof(*SanitizerTrapInfo), &SanitizerTrapInfo, TARGET_ALLOC_HOST));
+    memset(SanitizerTrapInfo, '\0', sizeof(SanitizerTrapInfoTy));
+
+    if (auto Err = GHandler.writeGlobalToDevice(*this, *Image, TrapId))
+      return Err;
+    {
+      auto KernelOrErr = getKernel("__sanitizer_register", Image);
+      if (auto Err = KernelOrErr.takeError())
+        consumeError(std::move(Err));
+      else
+        addGPUSanNewFn(*KernelOrErr);
     }
-    GenericKernelTy &Kernel = *KernelOrErr;
-    if (auto Err = Kernel.init(*this, *Image)) {
-      REPORT("Failure to init kernel: %s\n", toString(std::move(Err)).data());
-      return nullptr;
+    {
+      auto KernelOrErr = getKernel("__sanitizer_unregister", Image);
+      if (auto Err = KernelOrErr.takeError())
+        consumeError(std::move(Err));
+      else
+        addGPUSanFreeFn(*KernelOrErr);
     }
-    return &Kernel;
-  };
-  if (GenericKernelTy *Kernel = GetKernel("__sanitizer_register"))
-    addGPUSanNewFn(*Kernel);
-  if (GenericKernelTy *Kernel = GetKernel("__sanitizer_unregister"))
-    addGPUSanFreeFn(*Kernel);
+  }
 
   // Return the pointer to the table of entries.
   return Image;
@@ -1036,16 +1046,6 @@ Error GenericDeviceTy::setupDeviceMemoryPool(GenericPluginTy &Plugin,
   if (auto Err = GHandler.writeGlobalToDevice(*this, Image, TrackerGlobal))
     return Err;
 
-  auto *&SanitizerTrapInfo = SanitizerTrapInfos[&Image];
-  SanitizerTrapInfo = reinterpret_cast<SanitizerTrapInfoTy *>(allocate(
-      sizeof(*SanitizerTrapInfo), &SanitizerTrapInfo, TARGET_ALLOC_HOST));
-  memset(SanitizerTrapInfo, '\0', sizeof(SanitizerTrapInfoTy));
-
-  GlobalTy TrapId("__sanitizer_trap_info_ptr", sizeof(SanitizerTrapInfo),
-                  &SanitizerTrapInfo);
-  if (auto Err = GHandler.writeGlobalToDevice(*this, Image, TrapId))
-    return Err;
-
   // Create the metainfo of the device environment global.
   GlobalTy DevEnvGlobal("__omp_rtl_device_memory_pool",
                         sizeof(DeviceMemoryPoolTy), &DeviceMemoryPool);
@@ -2404,32 +2404,17 @@ void GPUSanTy::checkAndReportError() {
             Green(), STI->PtrSlot, STI->PtrTag, STI->PtrKind, STI->AllocationId,
             STI->AllocationTag, STI->AllocationKind, Default());
 
-    AllocationInfoTy *AllocationInfo = nullptr;
-    if (STI->AllocationStart) {
-      auto AllocationTraceMap = Device.AllocationTraces.getExclusiveAccessor();
-      AllocationInfo = (*AllocationTraceMap)[STI->AllocationStart];
-    }
-    if (AllocationInfo) {
-      fprintf(stderr, "\nAllocated at\n");
-      fprintf(stderr, "%s", AllocationInfo->AllocationTrace.c_str());
-
-      if (!AllocationInfo->DeallocationTrace.empty()) {
-        fprintf(stderr, "\nDeallocated at\n");
-        fprintf(stderr, "%s", AllocationInfo->DeallocationTrace.c_str());
-      }
-    }
-
   };
 
   switch (STI->ErrorCode) {
   case SanitizerTrapInfoTy::None:
     llvm_unreachable("Unexpected exception");
   case SanitizerTrapInfoTy::ExceedsLength:
-    fprintf(stderr, "%sERROR: OffloadSanitizer %s %d\n%s", Red(),
+    fprintf(stderr, "%sERROR: OffloadSanitizer %s %lu\n%s", Red(),
             "exceeds length", STI->PtrSlot, Default());
     break;
   case SanitizerTrapInfoTy::ExceedsSlots:
-    fprintf(stderr, "%sERROR: OffloadSanitizer %s %d\n%s", Red(),
+    fprintf(stderr, "%sERROR: OffloadSanitizer %s %lu\n%s", Red(),
             "exceeds slots", STI->PtrSlot, Default());
     break;
   case SanitizerTrapInfoTy::PointerOutsideAllocation:
@@ -2449,7 +2434,7 @@ void GPUSanTy::checkAndReportError() {
     DiagnoseAccess("use-after-free");
     break;
   case SanitizerTrapInfoTy::MemoryLeak:
-    fprintf(stderr, "%sERROR: OffloadSanitizer memory leak at slot %d\n%s",
+    fprintf(stderr, "%sERROR: OffloadSanitizer memory leak at slot %lu\n%s",
             Red(), STI->PtrSlot, Default());
     break;
   case SanitizerTrapInfoTy::GarbagePointer:

offload/src/Kernels/Sanitizer.cpp

@@ -15,12 +15,12 @@ __device__ void __sanitizer_register_host(void *P, uint64_t Bytes,
                                           uint64_t Slot);
 __device__ void __sanitizer_unregister_host(void *P);
 
-[[clang::disable_sanitizer_instrumentation]] __global__ void
+[[gnu::weak, clang::disable_sanitizer_instrumentation]] __global__ void
 __sanitizer_register(void *P, uint64_t Bytes, uint64_t Slot) {
   __sanitizer_register_host(P, Bytes, Slot);
 }
 
-[[clang::disable_sanitizer_instrumentation]] __global__ void
+[[gnu::weak, clang::disable_sanitizer_instrumentation]] __global__ void
 __sanitizer_unregister(void *P) {
   __sanitizer_unregister_host(P);
 }

offload/src/exports

@@ -75,9 +75,6 @@ VERS1.0 {
     __llvmPushCallConfiguration;
     __llvmPopCallConfiguration;
     llvmLaunchKernel;
-    ompx_new_allocation_host;
-    ompx_free_allocation_host;
-    ompx_register_image_functions;
   local:
     *;
 };

offload/src/omptarget.cpp

@@ -215,11 +215,6 @@ void *targetAllocExplicit(size_t Size, int64_t DeviceNum, int Kind,
     return Rc;
   }
 
-  if (checkDeviceAndCtors(DeviceNum, nullptr)) {
-    DP("Not offloading to device %" PRId64 "\n", DeviceNum);
-    return Rc;
-  }
-
   auto DeviceOrErr = PM->getDevice(DeviceNum);
   if (!DeviceOrErr)
     FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());