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 // Sets the `Context` for the duration of the instance and restores the previous
34 // context on destruction.
35 class ScopedContext {
36 public:
37   ScopedContext() {
38     // Static reference to HIP primary context for device ordinal 0.
39     static hipCtx_t context = [] {
40       HIP_REPORT_IF_ERROR(hipInit(/*flags=*/0));
41       hipDevice_t device;
42       HIP_REPORT_IF_ERROR(hipDeviceGet(&device, /*ordinal=*/0));
43       hipCtx_t ctx;
44       HIP_REPORT_IF_ERROR(hipDevicePrimaryCtxRetain(&ctx, device));
45       return ctx;
46     }();
47 
48     HIP_REPORT_IF_ERROR(hipCtxPushCurrent(context));
49   }
50 
51   ~ScopedContext() { HIP_REPORT_IF_ERROR(hipCtxPopCurrent(nullptr)); }
52 };
53 
54 extern "C" hipModule_t mgpuModuleLoad(void *data) {
55   ScopedContext scopedContext;
56   hipModule_t module = nullptr;
57   HIP_REPORT_IF_ERROR(hipModuleLoadData(&module, data));
58   return module;
59 }
60 
61 extern "C" void mgpuModuleUnload(hipModule_t module) {
62   HIP_REPORT_IF_ERROR(hipModuleUnload(module));
63 }
64 
65 extern "C" hipFunction_t mgpuModuleGetFunction(hipModule_t module,
66                                                const char *name) {
67   hipFunction_t function = nullptr;
68   HIP_REPORT_IF_ERROR(hipModuleGetFunction(&function, module, name));
69   return function;
70 }
71 
72 // The wrapper uses intptr_t instead of ROCM's unsigned int to match
73 // the type of MLIR's index type. This avoids the need for casts in the
74 // generated MLIR code.
75 extern "C" void mgpuLaunchKernel(hipFunction_t function, intptr_t gridX,
76                                  intptr_t gridY, intptr_t gridZ,
77                                  intptr_t blockX, intptr_t blockY,
78                                  intptr_t blockZ, int32_t smem,
79                                  hipStream_t stream, void **params,
80                                  void **extra) {
81   ScopedContext scopedContext;
82   HIP_REPORT_IF_ERROR(hipModuleLaunchKernel(function, gridX, gridY, gridZ,
83                                             blockX, blockY, blockZ, smem,
84                                             stream, params, extra));
85 }
86 
87 extern "C" hipStream_t mgpuStreamCreate() {
88   ScopedContext scopedContext;
89   hipStream_t stream = nullptr;
90   HIP_REPORT_IF_ERROR(hipStreamCreate(&stream));
91   return stream;
92 }
93 
94 extern "C" void mgpuStreamDestroy(hipStream_t stream) {
95   HIP_REPORT_IF_ERROR(hipStreamDestroy(stream));
96 }
97 
98 extern "C" void mgpuStreamSynchronize(hipStream_t stream) {
99   return HIP_REPORT_IF_ERROR(hipStreamSynchronize(stream));
100 }
101 
102 extern "C" void mgpuStreamWaitEvent(hipStream_t stream, hipEvent_t event) {
103   HIP_REPORT_IF_ERROR(hipStreamWaitEvent(stream, event, /*flags=*/0));
104 }
105 
106 extern "C" hipEvent_t mgpuEventCreate() {
107   ScopedContext scopedContext;
108   hipEvent_t event = nullptr;
109   HIP_REPORT_IF_ERROR(hipEventCreateWithFlags(&event, hipEventDisableTiming));
110   return event;
111 }
112 
113 extern "C" void mgpuEventDestroy(hipEvent_t event) {
114   HIP_REPORT_IF_ERROR(hipEventDestroy(event));
115 }
116 
117 extern "C" void mgpuEventSynchronize(hipEvent_t event) {
118   HIP_REPORT_IF_ERROR(hipEventSynchronize(event));
119 }
120 
121 extern "C" void mgpuEventRecord(hipEvent_t event, hipStream_t stream) {
122   HIP_REPORT_IF_ERROR(hipEventRecord(event, stream));
123 }
124 
125 extern "C" void *mgpuMemAlloc(uint64_t sizeBytes, hipStream_t /*stream*/) {
126   ScopedContext scopedContext;
127   void *ptr;
128   HIP_REPORT_IF_ERROR(hipMalloc(&ptr, sizeBytes));
129   return ptr;
130 }
131 
132 extern "C" void mgpuMemFree(void *ptr, hipStream_t /*stream*/) {
133   HIP_REPORT_IF_ERROR(hipFree(ptr));
134 }
135 
136 extern "C" void mgpuMemcpy(void *dst, void *src, size_t sizeBytes,
137                            hipStream_t stream) {
138   HIP_REPORT_IF_ERROR(
139       hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream));
140 }
141 
142 extern "C" void mgpuMemset32(void *dst, int value, size_t count,
143                              hipStream_t stream) {
144   HIP_REPORT_IF_ERROR(hipMemsetD32Async(reinterpret_cast<hipDeviceptr_t>(dst),
145                                         value, count, stream));
146 }
147 /// Helper functions for writing mlir example code
148 
149 // Allows to register byte array with the ROCM runtime. Helpful until we have
150 // transfer functions implemented.
151 extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) {
152   ScopedContext scopedContext;
153   HIP_REPORT_IF_ERROR(hipHostRegister(ptr, sizeBytes, /*flags=*/0));
154 }
155 
156 // Allows to register a MemRef with the ROCm runtime. Helpful until we have
157 // transfer functions implemented.
158 extern "C" void
159 mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType<char, 1> *descriptor,
160                           int64_t elementSizeBytes) {
161 
162   llvm::SmallVector<int64_t, 4> denseStrides(rank);
163   llvm::ArrayRef<int64_t> sizes(descriptor->sizes, rank);
164   llvm::ArrayRef<int64_t> strides(sizes.end(), rank);
165 
166   std::partial_sum(sizes.rbegin(), sizes.rend(), denseStrides.rbegin(),
167                    std::multiplies<int64_t>());
168   auto sizeBytes = denseStrides.front() * elementSizeBytes;
169 
170   // Only densely packed tensors are currently supported.
171   std::rotate(denseStrides.begin(), denseStrides.begin() + 1,
172               denseStrides.end());
173   denseStrides.back() = 1;
174   assert(strides == llvm::makeArrayRef(denseStrides));
175 
176   auto ptr = descriptor->data + descriptor->offset * elementSizeBytes;
177   mgpuMemHostRegister(ptr, sizeBytes);
178 }
179 
180 template <typename T>
181 void mgpuMemGetDevicePointer(T *hostPtr, T **devicePtr) {
182   HIP_REPORT_IF_ERROR(hipSetDevice(0));
183   HIP_REPORT_IF_ERROR(
184       hipHostGetDevicePointer((void **)devicePtr, hostPtr, /*flags=*/0));
185 }
186 
187 extern "C" StridedMemRefType<float, 1>
188 mgpuMemGetDeviceMemRef1dFloat(float *allocated, float *aligned, int64_t offset,
189                               int64_t size, int64_t stride) {
190   float *devicePtr = nullptr;
191   mgpuMemGetDevicePointer(aligned, &devicePtr);
192   return {devicePtr, devicePtr, offset, {size}, {stride}};
193 }
194 
195 extern "C" StridedMemRefType<int32_t, 1>
196 mgpuMemGetDeviceMemRef1dInt32(int32_t *allocated, int32_t *aligned,
197                               int64_t offset, int64_t size, int64_t stride) {
198   int32_t *devicePtr = nullptr;
199   mgpuMemGetDevicePointer(aligned, &devicePtr);
200   return {devicePtr, devicePtr, offset, {size}, {stride}};
201 }
202