1 //===- RocmRuntimeWrappers.cpp - MLIR ROCM runtime wrapper library --------===// 2 // 3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 // See https://llvm.org/LICENSE.txt for license information. 5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 // 7 //===----------------------------------------------------------------------===// 8 // 9 // Implements C wrappers around the ROCM library for easy linking in ORC jit. 10 // Also adds some debugging helpers that are helpful when writing MLIR code to 11 // run on GPUs. 12 // 13 //===----------------------------------------------------------------------===// 14 15 #include <cassert> 16 #include <numeric> 17 18 #include "mlir/ExecutionEngine/CRunnerUtils.h" 19 #include "llvm/ADT/ArrayRef.h" 20 21 #include "hip/hip_runtime.h" 22 23 #define HIP_REPORT_IF_ERROR(expr) \ 24 [](hipError_t result) { \ 25 if (!result) \ 26 return; \ 27 const char *name = hipGetErrorName(result); \ 28 if (!name) \ 29 name = "<unknown>"; \ 30 fprintf(stderr, "'%s' failed with '%s'\n", #expr, name); \ 31 }(expr) 32 33 thread_local static int32_t defaultDevice = 0; 34 35 // Sets the `Context` for the duration of the instance and restores the previous 36 // context on destruction. 37 class ScopedContext { 38 public: 39 ScopedContext() { 40 // Static reference to HIP primary context for device ordinal defaultDevice. 41 static hipCtx_t context = [] { 42 HIP_REPORT_IF_ERROR(hipInit(/*flags=*/0)); 43 hipDevice_t device; 44 HIP_REPORT_IF_ERROR(hipDeviceGet(&device, /*ordinal=*/defaultDevice)); 45 hipCtx_t ctx; 46 HIP_REPORT_IF_ERROR(hipDevicePrimaryCtxRetain(&ctx, device)); 47 return ctx; 48 }(); 49 50 HIP_REPORT_IF_ERROR(hipCtxPushCurrent(context)); 51 } 52 53 ~ScopedContext() { HIP_REPORT_IF_ERROR(hipCtxPopCurrent(nullptr)); } 54 }; 55 56 extern "C" hipModule_t mgpuModuleLoad(void *data) { 57 ScopedContext scopedContext; 58 hipModule_t module = nullptr; 59 HIP_REPORT_IF_ERROR(hipModuleLoadData(&module, data)); 60 return module; 61 } 62 63 extern "C" void mgpuModuleUnload(hipModule_t module) { 64 HIP_REPORT_IF_ERROR(hipModuleUnload(module)); 65 } 66 67 extern "C" hipFunction_t mgpuModuleGetFunction(hipModule_t module, 68 const char *name) { 69 hipFunction_t function = nullptr; 70 HIP_REPORT_IF_ERROR(hipModuleGetFunction(&function, module, name)); 71 return function; 72 } 73 74 // The wrapper uses intptr_t instead of ROCM's unsigned int to match 75 // the type of MLIR's index type. This avoids the need for casts in the 76 // generated MLIR code. 77 extern "C" void mgpuLaunchKernel(hipFunction_t function, intptr_t gridX, 78 intptr_t gridY, intptr_t gridZ, 79 intptr_t blockX, intptr_t blockY, 80 intptr_t blockZ, int32_t smem, 81 hipStream_t stream, void **params, 82 void **extra) { 83 ScopedContext scopedContext; 84 HIP_REPORT_IF_ERROR(hipModuleLaunchKernel(function, gridX, gridY, gridZ, 85 blockX, blockY, blockZ, smem, 86 stream, params, extra)); 87 } 88 89 extern "C" hipStream_t mgpuStreamCreate() { 90 ScopedContext scopedContext; 91 hipStream_t stream = nullptr; 92 HIP_REPORT_IF_ERROR(hipStreamCreate(&stream)); 93 return stream; 94 } 95 96 extern "C" void mgpuStreamDestroy(hipStream_t stream) { 97 HIP_REPORT_IF_ERROR(hipStreamDestroy(stream)); 98 } 99 100 extern "C" void mgpuStreamSynchronize(hipStream_t stream) { 101 return HIP_REPORT_IF_ERROR(hipStreamSynchronize(stream)); 102 } 103 104 extern "C" void mgpuStreamWaitEvent(hipStream_t stream, hipEvent_t event) { 105 HIP_REPORT_IF_ERROR(hipStreamWaitEvent(stream, event, /*flags=*/0)); 106 } 107 108 extern "C" hipEvent_t mgpuEventCreate() { 109 ScopedContext scopedContext; 110 hipEvent_t event = nullptr; 111 HIP_REPORT_IF_ERROR(hipEventCreateWithFlags(&event, hipEventDisableTiming)); 112 return event; 113 } 114 115 extern "C" void mgpuEventDestroy(hipEvent_t event) { 116 HIP_REPORT_IF_ERROR(hipEventDestroy(event)); 117 } 118 119 extern "C" void mgpuEventSynchronize(hipEvent_t event) { 120 HIP_REPORT_IF_ERROR(hipEventSynchronize(event)); 121 } 122 123 extern "C" void mgpuEventRecord(hipEvent_t event, hipStream_t stream) { 124 HIP_REPORT_IF_ERROR(hipEventRecord(event, stream)); 125 } 126 127 extern "C" void *mgpuMemAlloc(uint64_t sizeBytes, hipStream_t /*stream*/) { 128 ScopedContext scopedContext; 129 void *ptr; 130 HIP_REPORT_IF_ERROR(hipMalloc(&ptr, sizeBytes)); 131 return ptr; 132 } 133 134 extern "C" void mgpuMemFree(void *ptr, hipStream_t /*stream*/) { 135 HIP_REPORT_IF_ERROR(hipFree(ptr)); 136 } 137 138 extern "C" void mgpuMemcpy(void *dst, void *src, size_t sizeBytes, 139 hipStream_t stream) { 140 HIP_REPORT_IF_ERROR( 141 hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream)); 142 } 143 144 extern "C" void mgpuMemset32(void *dst, int value, size_t count, 145 hipStream_t stream) { 146 HIP_REPORT_IF_ERROR(hipMemsetD32Async(reinterpret_cast<hipDeviceptr_t>(dst), 147 value, count, stream)); 148 } 149 /// Helper functions for writing mlir example code 150 151 // Allows to register byte array with the ROCM runtime. Helpful until we have 152 // transfer functions implemented. 153 extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) { 154 ScopedContext scopedContext; 155 HIP_REPORT_IF_ERROR(hipHostRegister(ptr, sizeBytes, /*flags=*/0)); 156 } 157 158 // Allows to register a MemRef with the ROCm runtime. Helpful until we have 159 // transfer functions implemented. 160 extern "C" void 161 mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType<char, 1> *descriptor, 162 int64_t elementSizeBytes) { 163 164 llvm::SmallVector<int64_t, 4> denseStrides(rank); 165 llvm::ArrayRef<int64_t> sizes(descriptor->sizes, rank); 166 llvm::ArrayRef<int64_t> strides(sizes.end(), rank); 167 168 std::partial_sum(sizes.rbegin(), sizes.rend(), denseStrides.rbegin(), 169 std::multiplies<int64_t>()); 170 auto sizeBytes = denseStrides.front() * elementSizeBytes; 171 172 // Only densely packed tensors are currently supported. 173 std::rotate(denseStrides.begin(), denseStrides.begin() + 1, 174 denseStrides.end()); 175 denseStrides.back() = 1; 176 assert(strides == llvm::makeArrayRef(denseStrides)); 177 178 auto ptr = descriptor->data + descriptor->offset * elementSizeBytes; 179 mgpuMemHostRegister(ptr, sizeBytes); 180 } 181 182 template <typename T> 183 void mgpuMemGetDevicePointer(T *hostPtr, T **devicePtr) { 184 HIP_REPORT_IF_ERROR(hipSetDevice(0)); 185 HIP_REPORT_IF_ERROR( 186 hipHostGetDevicePointer((void **)devicePtr, hostPtr, /*flags=*/0)); 187 } 188 189 extern "C" StridedMemRefType<float, 1> 190 mgpuMemGetDeviceMemRef1dFloat(float *allocated, float *aligned, int64_t offset, 191 int64_t size, int64_t stride) { 192 float *devicePtr = nullptr; 193 mgpuMemGetDevicePointer(aligned, &devicePtr); 194 return {devicePtr, devicePtr, offset, {size}, {stride}}; 195 } 196 197 extern "C" StridedMemRefType<int32_t, 1> 198 mgpuMemGetDeviceMemRef1dInt32(int32_t *allocated, int32_t *aligned, 199 int64_t offset, int64_t size, int64_t stride) { 200 int32_t *devicePtr = nullptr; 201 mgpuMemGetDevicePointer(aligned, &devicePtr); 202 return {devicePtr, devicePtr, offset, {size}, {stride}}; 203 } 204 205 extern "C" void mgpuSetDefaultDevice(int32_t device) { 206 defaultDevice = device; 207 HIP_REPORT_IF_ERROR(hipSetDevice(device)); 208 } 209