[mlir][ExecutionEngine] Remove ScopedContext from ROCm wrappers

The push/pop context APIs are deprecated in HIP, and keeping the default device set is handled in IHP using hipSetDevice(). Reviewed By: ThomasRaoux Differential Revision: https://reviews.llvm.org/D134747
llvm · Sep 27, 2022 · 10a8ec8 · 10a8ec8
1 parent ed8409d
commit 10a8ec8
Showing 1 changed file with 0 additions and 27 deletions.
diff --git a/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp
@@ -32,29 +32,7 @@
 
 thread_local static int32_t defaultDevice = 0;
 
-// Sets the `Context` for the duration of the instance and restores the previous
-// context on destruction.
-class ScopedContext {
-public:
-  ScopedContext() {
-    // Static reference to HIP primary context for device ordinal defaultDevice.
-    static hipCtx_t context = [] {
-      HIP_REPORT_IF_ERROR(hipInit(/*flags=*/0));
-      hipDevice_t device;
-      HIP_REPORT_IF_ERROR(hipDeviceGet(&device, /*ordinal=*/defaultDevice));
-      hipCtx_t ctx;
-      HIP_REPORT_IF_ERROR(hipDevicePrimaryCtxRetain(&ctx, device));
-      return ctx;
-    }();
-
-    HIP_REPORT_IF_ERROR(hipCtxPushCurrent(context));
-  }
-
-  ~ScopedContext() { HIP_REPORT_IF_ERROR(hipCtxPopCurrent(nullptr)); }
-};
-
 extern "C" hipModule_t mgpuModuleLoad(void *data) {
-  ScopedContext scopedContext;
   hipModule_t module = nullptr;
   HIP_REPORT_IF_ERROR(hipModuleLoadData(&module, data));
   return module;
@@ -80,14 +58,12 @@ extern "C" void mgpuLaunchKernel(hipFunction_t function, intptr_t gridX,
                                  intptr_t blockZ, int32_t smem,
                                  hipStream_t stream, void **params,
                                  void **extra) {
-  ScopedContext scopedContext;
   HIP_REPORT_IF_ERROR(hipModuleLaunchKernel(function, gridX, gridY, gridZ,
                                             blockX, blockY, blockZ, smem,
                                             stream, params, extra));
 }
 
 extern "C" hipStream_t mgpuStreamCreate() {
-  ScopedContext scopedContext;
   hipStream_t stream = nullptr;
   HIP_REPORT_IF_ERROR(hipStreamCreate(&stream));
   return stream;
@@ -106,7 +82,6 @@ extern "C" void mgpuStreamWaitEvent(hipStream_t stream, hipEvent_t event) {
 }
 
 extern "C" hipEvent_t mgpuEventCreate() {
-  ScopedContext scopedContext;
   hipEvent_t event = nullptr;
   HIP_REPORT_IF_ERROR(hipEventCreateWithFlags(&event, hipEventDisableTiming));
   return event;
@@ -125,7 +100,6 @@ extern "C" void mgpuEventRecord(hipEvent_t event, hipStream_t stream) {
 }
 
 extern "C" void *mgpuMemAlloc(uint64_t sizeBytes, hipStream_t /*stream*/) {
-  ScopedContext scopedContext;
   void *ptr;
   HIP_REPORT_IF_ERROR(hipMalloc(&ptr, sizeBytes));
   return ptr;
@@ -151,7 +125,6 @@ extern "C" void mgpuMemset32(void *dst, int value, size_t count,
 // Allows to register byte array with the ROCM runtime. Helpful until we have
 // transfer functions implemented.
 extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) {
-  ScopedContext scopedContext;
   HIP_REPORT_IF_ERROR(hipHostRegister(ptr, sizeBytes, /*flags=*/0));
 }