1 //===- rocm-runtime-wrappers.cpp - MLIR ROCM runner 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 #include "llvm/Support/raw_ostream.h"
21
22 #include "hip/hip_runtime.h"
23
24 #define HIP_REPORT_IF_ERROR(expr) \
25 [](hipError_t result) { \
26 if (!result) \
27 return; \
28 const char *name = hipGetErrorName(result); \
29 if (!name) \
30 name = "<unknown>"; \
31 llvm::errs() << "'" << #expr << "' failed with '" << name << "'\n"; \
32 }(expr)
33
34 // Static initialization of HIP context for device ordinal 0.
__anond51d5d1e0102null35 static auto InitializeCtx = [] {
36 HIP_REPORT_IF_ERROR(hipInit(/*flags=*/0));
37 hipDevice_t device;
38 HIP_REPORT_IF_ERROR(hipDeviceGet(&device, /*ordinal=*/0));
39 hipContext_t context;
40 HIP_REPORT_IF_ERROR(hipCtxCreate(&context, /*flags=*/0, device));
41 return 0;
42 }();
43
mgpuModuleLoad(void * data)44 extern "C" hipModule_t mgpuModuleLoad(void *data) {
45 hipModule_t module = nullptr;
46 HIP_REPORT_IF_ERROR(hipModuleLoadData(&module, data));
47 return module;
48 }
49
mgpuModuleUnload(hipModule_t module)50 extern "C" void mgpuModuleUnload(hipModule_t module) {
51 HIP_REPORT_IF_ERROR(hipModuleUnload(module));
52 }
53
mgpuModuleGetFunction(hipModule_t module,const char * name)54 extern "C" hipFunction_t mgpuModuleGetFunction(hipModule_t module,
55 const char *name) {
56 hipFunction_t function = nullptr;
57 HIP_REPORT_IF_ERROR(hipModuleGetFunction(&function, module, name));
58 return function;
59 }
60
61 // The wrapper uses intptr_t instead of ROCM's unsigned int to match
62 // the type of MLIR's index type. This avoids the need for casts in the
63 // 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)64 extern "C" void mgpuLaunchKernel(hipFunction_t function, intptr_t gridX,
65 intptr_t gridY, intptr_t gridZ,
66 intptr_t blockX, intptr_t blockY,
67 intptr_t blockZ, int32_t smem,
68 hipStream_t stream, void **params,
69 void **extra) {
70 HIP_REPORT_IF_ERROR(hipModuleLaunchKernel(function, gridX, gridY, gridZ,
71 blockX, blockY, blockZ, smem,
72 stream, params, extra));
73 }
74
mgpuStreamCreate()75 extern "C" hipStream_t mgpuStreamCreate() {
76 hipStream_t stream = nullptr;
77 HIP_REPORT_IF_ERROR(hipStreamCreate(&stream));
78 return stream;
79 }
80
mgpuStreamDestroy(hipStream_t stream)81 extern "C" void mgpuStreamDestroy(hipStream_t stream) {
82 HIP_REPORT_IF_ERROR(hipStreamDestroy(stream));
83 }
84
mgpuStreamSynchronize(hipStream_t stream)85 extern "C" void mgpuStreamSynchronize(hipStream_t stream) {
86 return HIP_REPORT_IF_ERROR(hipStreamSynchronize(stream));
87 }
88
mgpuStreamWaitEvent(hipStream_t stream,hipEvent_t event)89 extern "C" void mgpuStreamWaitEvent(hipStream_t stream, hipEvent_t event) {
90 HIP_REPORT_IF_ERROR(hipStreamWaitEvent(stream, event, /*flags=*/0));
91 }
92
mgpuEventCreate()93 extern "C" hipEvent_t mgpuEventCreate() {
94 hipEvent_t event = nullptr;
95 HIP_REPORT_IF_ERROR(hipEventCreateWithFlags(&event, hipEventDisableTiming));
96 return event;
97 }
98
mgpuEventDestroy(hipEvent_t event)99 extern "C" void mgpuEventDestroy(hipEvent_t event) {
100 HIP_REPORT_IF_ERROR(hipEventDestroy(event));
101 }
102
mgpuEventSynchronize(hipEvent_t event)103 extern "C" void mgpuEventSynchronize(hipEvent_t event) {
104 HIP_REPORT_IF_ERROR(hipEventSynchronize(event));
105 }
106
mgpuEventRecord(hipEvent_t event,hipStream_t stream)107 extern "C" void mgpuEventRecord(hipEvent_t event, hipStream_t stream) {
108 HIP_REPORT_IF_ERROR(hipEventRecord(event, stream));
109 }
110
mgpuMemAlloc(uint64_t sizeBytes,hipStream_t)111 extern "C" void *mgpuMemAlloc(uint64_t sizeBytes, hipStream_t /*stream*/) {
112 void *ptr;
113 HIP_REPORT_IF_ERROR(hipMemAlloc(&ptr, sizeBytes));
114 return ptr;
115 }
116
mgpuMemFree(void * ptr,hipStream_t)117 extern "C" void mgpuMemFree(void *ptr, hipStream_t /*stream*/) {
118 HIP_REPORT_IF_ERROR(hipMemFree(ptr));
119 }
120
121 /// Helper functions for writing mlir example code
122
123 // Allows to register byte array with the ROCM runtime. Helpful until we have
124 // transfer functions implemented.
mgpuMemHostRegister(void * ptr,uint64_t sizeBytes)125 extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) {
126 HIP_REPORT_IF_ERROR(hipHostRegister(ptr, sizeBytes, /*flags=*/0));
127 }
128
129 // Allows to register a MemRef with the ROCm runtime. Helpful until we have
130 // transfer functions implemented.
131 extern "C" void
mgpuMemHostRegisterMemRef(int64_t rank,StridedMemRefType<char,1> * descriptor,int64_t elementSizeBytes)132 mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType<char, 1> *descriptor,
133 int64_t elementSizeBytes) {
134
135 llvm::SmallVector<int64_t, 4> denseStrides(rank);
136 llvm::ArrayRef<int64_t> sizes(descriptor->sizes, rank);
137 llvm::ArrayRef<int64_t> strides(sizes.end(), rank);
138
139 std::partial_sum(sizes.rbegin(), sizes.rend(), denseStrides.rbegin(),
140 std::multiplies<int64_t>());
141 auto sizeBytes = denseStrides.front() * elementSizeBytes;
142
143 // Only densely packed tensors are currently supported.
144 std::rotate(denseStrides.begin(), denseStrides.begin() + 1,
145 denseStrides.end());
146 denseStrides.back() = 1;
147 assert(strides == llvm::makeArrayRef(denseStrides));
148
149 auto ptr = descriptor->data + descriptor->offset * elementSizeBytes;
150 mgpuMemHostRegister(ptr, sizeBytes);
151 }
152
153 template <typename T>
mgpuMemGetDevicePointer(T * hostPtr,T ** devicePtr)154 void mgpuMemGetDevicePointer(T *hostPtr, T **devicePtr) {
155 HIP_REPORT_IF_ERROR(hipSetDevice(0));
156 HIP_REPORT_IF_ERROR(
157 hipHostGetDevicePointer((void **)devicePtr, hostPtr, /*flags=*/0));
158 }
159
160 extern "C" StridedMemRefType<float, 1>
mgpuMemGetDeviceMemRef1dFloat(float * allocated,float * aligned,int64_t offset,int64_t size,int64_t stride)161 mgpuMemGetDeviceMemRef1dFloat(float *allocated, float *aligned, int64_t offset,
162 int64_t size, int64_t stride) {
163 float *devicePtr = nullptr;
164 mgpuMemGetDevicePointer(aligned, &devicePtr);
165 return {devicePtr, devicePtr, offset, {size}, {stride}};
166 }
167
168 extern "C" StridedMemRefType<int32_t, 1>
mgpuMemGetDeviceMemRef1dInt32(int32_t * allocated,int32_t * aligned,int64_t offset,int64_t size,int64_t stride)169 mgpuMemGetDeviceMemRef1dInt32(int32_t *allocated, int32_t *aligned,
170 int64_t offset, int64_t size, int64_t stride) {
171 int32_t *devicePtr = nullptr;
172 mgpuMemGetDevicePointer(aligned, &devicePtr);
173 return {devicePtr, devicePtr, offset, {size}, {stride}};
174 }
175