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:
ScopedContext()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
~ScopedContext()53 ~ScopedContext() { HIP_REPORT_IF_ERROR(hipCtxPopCurrent(nullptr)); }
54 };
55
mgpuModuleLoad(void * data)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
mgpuModuleUnload(hipModule_t module)63 extern "C" void mgpuModuleUnload(hipModule_t module) {
64 HIP_REPORT_IF_ERROR(hipModuleUnload(module));
65 }
66
mgpuModuleGetFunction(hipModule_t module,const char * name)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.
mgpuLaunchKernel(hipFunction_t function,intptr_t gridX,intptr_t gridY,intptr_t gridZ,intptr_t blockX,intptr_t blockY,intptr_t blockZ,int32_t smem,hipStream_t stream,void ** params,void ** extra)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
mgpuStreamCreate()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
mgpuStreamDestroy(hipStream_t stream)96 extern "C" void mgpuStreamDestroy(hipStream_t stream) {
97 HIP_REPORT_IF_ERROR(hipStreamDestroy(stream));
98 }
99
mgpuStreamSynchronize(hipStream_t stream)100 extern "C" void mgpuStreamSynchronize(hipStream_t stream) {
101 return HIP_REPORT_IF_ERROR(hipStreamSynchronize(stream));
102 }
103
mgpuStreamWaitEvent(hipStream_t stream,hipEvent_t event)104 extern "C" void mgpuStreamWaitEvent(hipStream_t stream, hipEvent_t event) {
105 HIP_REPORT_IF_ERROR(hipStreamWaitEvent(stream, event, /*flags=*/0));
106 }
107
mgpuEventCreate()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
mgpuEventDestroy(hipEvent_t event)115 extern "C" void mgpuEventDestroy(hipEvent_t event) {
116 HIP_REPORT_IF_ERROR(hipEventDestroy(event));
117 }
118
mgpuEventSynchronize(hipEvent_t event)119 extern "C" void mgpuEventSynchronize(hipEvent_t event) {
120 HIP_REPORT_IF_ERROR(hipEventSynchronize(event));
121 }
122
mgpuEventRecord(hipEvent_t event,hipStream_t stream)123 extern "C" void mgpuEventRecord(hipEvent_t event, hipStream_t stream) {
124 HIP_REPORT_IF_ERROR(hipEventRecord(event, stream));
125 }
126
mgpuMemAlloc(uint64_t sizeBytes,hipStream_t)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
mgpuMemFree(void * ptr,hipStream_t)134 extern "C" void mgpuMemFree(void *ptr, hipStream_t /*stream*/) {
135 HIP_REPORT_IF_ERROR(hipFree(ptr));
136 }
137
mgpuMemcpy(void * dst,void * src,size_t sizeBytes,hipStream_t stream)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
mgpuMemset32(void * dst,int value,size_t count,hipStream_t stream)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.
mgpuMemHostRegister(void * ptr,uint64_t sizeBytes)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
mgpuMemHostRegisterMemRef(int64_t rank,StridedMemRefType<char,1> * descriptor,int64_t elementSizeBytes)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>
mgpuMemGetDevicePointer(T * hostPtr,T ** devicePtr)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>
mgpuMemGetDeviceMemRef1dFloat(float * allocated,float * aligned,int64_t offset,int64_t size,int64_t stride)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>
mgpuMemGetDeviceMemRef1dInt32(int32_t * allocated,int32_t * aligned,int64_t offset,int64_t size,int64_t stride)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
mgpuSetDefaultDevice(int32_t device)205 extern "C" void mgpuSetDefaultDevice(int32_t device) {
206 defaultDevice = device;
207 HIP_REPORT_IF_ERROR(hipSetDevice(device));
208 }
209