From 10a8ec86a2264bfb8127d0744394c5a69a396294 Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak@amd.com>
Date: Tue, 27 Sep 2022 16:13:07 +0000
Subject: [PATCH] [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
---
 mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp | 27 ------------------------
 1 file changed, 27 deletions(-)

diff --git a/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp
index 34363cc..0d85531 100644
--- 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));
 }
 
-- 
2.7.4