diff options
author | Krzysztof Drewniak <Krzysztof.Drewniak@amd.com> | 2022-09-27 19:13:07 +0300 |
---|---|---|
committer | Krzysztof Drewniak <Krzysztof.Drewniak@amd.com> | 2022-09-27 19:56:12 +0300 |
commit | 10a8ec86a2264bfb8127d0744394c5a69a396294 (patch) | |
tree | 97d33a08df8f447932f0ff677379832895907b8e | |
parent | ed8409dfa0a92d80c021f13ca271737492522cc7 (diff) |
[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
-rw-r--r-- | mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp | 27 |
1 files changed, 0 insertions, 27 deletions
diff --git a/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp index 34363ccc6141..0d85531373d4 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)); } |