forked from OSchip/llvm-project
198 lines
7.4 KiB
C++
198 lines
7.4 KiB
C++
//===- CudaRuntimeWrappers.cpp - MLIR CUDA API wrapper library ------------===//
|
|
//
|
|
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
|
|
// See https://llvm.org/LICENSE.txt for license information.
|
|
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
//
|
|
// Implements C wrappers around the CUDA library for easy linking in ORC jit.
|
|
// Also adds some debugging helpers that are helpful when writing MLIR code to
|
|
// run on GPUs.
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
#include "mlir/ExecutionEngine/CRunnerUtils.h"
|
|
|
|
#include <stdio.h>
|
|
|
|
#include "cuda.h"
|
|
|
|
#ifdef _WIN32
|
|
#define MLIR_CUDA_WRAPPERS_EXPORT __declspec(dllexport)
|
|
#else
|
|
#define MLIR_CUDA_WRAPPERS_EXPORT
|
|
#endif // _WIN32
|
|
|
|
#define CUDA_REPORT_IF_ERROR(expr) \
|
|
[](CUresult result) { \
|
|
if (!result) \
|
|
return; \
|
|
const char *name = nullptr; \
|
|
cuGetErrorName(result, &name); \
|
|
if (!name) \
|
|
name = "<unknown>"; \
|
|
fprintf(stderr, "'%s' failed with '%s'\n", #expr, name); \
|
|
}(expr)
|
|
|
|
thread_local static int32_t defaultDevice = 0;
|
|
|
|
// Make the primary context of the current default device current for the
|
|
// duration
|
|
// of the instance and restore the previous context on destruction.
|
|
class ScopedContext {
|
|
public:
|
|
ScopedContext() {
|
|
// Static reference to CUDA primary context for device ordinal
|
|
// defaultDevice.
|
|
static CUcontext context = [] {
|
|
CUDA_REPORT_IF_ERROR(cuInit(/*flags=*/0));
|
|
CUdevice device;
|
|
CUDA_REPORT_IF_ERROR(cuDeviceGet(&device, /*ordinal=*/defaultDevice));
|
|
CUcontext ctx;
|
|
// Note: this does not affect the current context.
|
|
CUDA_REPORT_IF_ERROR(cuDevicePrimaryCtxRetain(&ctx, device));
|
|
return ctx;
|
|
}();
|
|
|
|
CUDA_REPORT_IF_ERROR(cuCtxPushCurrent(context));
|
|
}
|
|
|
|
~ScopedContext() { CUDA_REPORT_IF_ERROR(cuCtxPopCurrent(nullptr)); }
|
|
};
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT CUmodule mgpuModuleLoad(void *data) {
|
|
ScopedContext scopedContext;
|
|
CUmodule module = nullptr;
|
|
CUDA_REPORT_IF_ERROR(cuModuleLoadData(&module, data));
|
|
return module;
|
|
}
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuModuleUnload(CUmodule module) {
|
|
CUDA_REPORT_IF_ERROR(cuModuleUnload(module));
|
|
}
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT CUfunction
|
|
mgpuModuleGetFunction(CUmodule module, const char *name) {
|
|
CUfunction function = nullptr;
|
|
CUDA_REPORT_IF_ERROR(cuModuleGetFunction(&function, module, name));
|
|
return function;
|
|
}
|
|
|
|
// The wrapper uses intptr_t instead of CUDA's unsigned int to match
|
|
// the type of MLIR's index type. This avoids the need for casts in the
|
|
// generated MLIR code.
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT void
|
|
mgpuLaunchKernel(CUfunction function, intptr_t gridX, intptr_t gridY,
|
|
intptr_t gridZ, intptr_t blockX, intptr_t blockY,
|
|
intptr_t blockZ, int32_t smem, CUstream stream, void **params,
|
|
void **extra) {
|
|
ScopedContext scopedContext;
|
|
CUDA_REPORT_IF_ERROR(cuLaunchKernel(function, gridX, gridY, gridZ, blockX,
|
|
blockY, blockZ, smem, stream, params,
|
|
extra));
|
|
}
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT CUstream mgpuStreamCreate() {
|
|
ScopedContext scopedContext;
|
|
CUstream stream = nullptr;
|
|
CUDA_REPORT_IF_ERROR(cuStreamCreate(&stream, CU_STREAM_NON_BLOCKING));
|
|
return stream;
|
|
}
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuStreamDestroy(CUstream stream) {
|
|
CUDA_REPORT_IF_ERROR(cuStreamDestroy(stream));
|
|
}
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT void
|
|
mgpuStreamSynchronize(CUstream stream) {
|
|
CUDA_REPORT_IF_ERROR(cuStreamSynchronize(stream));
|
|
}
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuStreamWaitEvent(CUstream stream,
|
|
CUevent event) {
|
|
CUDA_REPORT_IF_ERROR(cuStreamWaitEvent(stream, event, /*flags=*/0));
|
|
}
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT CUevent mgpuEventCreate() {
|
|
ScopedContext scopedContext;
|
|
CUevent event = nullptr;
|
|
CUDA_REPORT_IF_ERROR(cuEventCreate(&event, CU_EVENT_DISABLE_TIMING));
|
|
return event;
|
|
}
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuEventDestroy(CUevent event) {
|
|
CUDA_REPORT_IF_ERROR(cuEventDestroy(event));
|
|
}
|
|
|
|
extern MLIR_CUDA_WRAPPERS_EXPORT "C" void mgpuEventSynchronize(CUevent event) {
|
|
CUDA_REPORT_IF_ERROR(cuEventSynchronize(event));
|
|
}
|
|
|
|
extern MLIR_CUDA_WRAPPERS_EXPORT "C" void mgpuEventRecord(CUevent event,
|
|
CUstream stream) {
|
|
CUDA_REPORT_IF_ERROR(cuEventRecord(event, stream));
|
|
}
|
|
|
|
extern "C" void *mgpuMemAlloc(uint64_t sizeBytes, CUstream /*stream*/) {
|
|
ScopedContext scopedContext;
|
|
CUdeviceptr ptr;
|
|
CUDA_REPORT_IF_ERROR(cuMemAlloc(&ptr, sizeBytes));
|
|
return reinterpret_cast<void *>(ptr);
|
|
}
|
|
|
|
extern "C" void mgpuMemFree(void *ptr, CUstream /*stream*/) {
|
|
CUDA_REPORT_IF_ERROR(cuMemFree(reinterpret_cast<CUdeviceptr>(ptr)));
|
|
}
|
|
|
|
extern "C" void mgpuMemcpy(void *dst, void *src, size_t sizeBytes,
|
|
CUstream stream) {
|
|
CUDA_REPORT_IF_ERROR(cuMemcpyAsync(reinterpret_cast<CUdeviceptr>(dst),
|
|
reinterpret_cast<CUdeviceptr>(src),
|
|
sizeBytes, stream));
|
|
}
|
|
|
|
extern "C" void mgpuMemset32(void *dst, unsigned int value, size_t count,
|
|
CUstream stream) {
|
|
CUDA_REPORT_IF_ERROR(cuMemsetD32Async(reinterpret_cast<CUdeviceptr>(dst),
|
|
value, count, stream));
|
|
}
|
|
|
|
/// Helper functions for writing mlir example code
|
|
|
|
// Allows to register byte array with the CUDA runtime. Helpful until we have
|
|
// transfer functions implemented.
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT void
|
|
mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) {
|
|
ScopedContext scopedContext;
|
|
CUDA_REPORT_IF_ERROR(cuMemHostRegister(ptr, sizeBytes, /*flags=*/0));
|
|
}
|
|
|
|
/// Registers a memref with the CUDA runtime. `descriptor` is a pointer to a
|
|
/// ranked memref descriptor struct of rank `rank`. Helpful until we have
|
|
/// transfer functions implemented.
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT void
|
|
mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType<char, 1> *descriptor,
|
|
int64_t elementSizeBytes) {
|
|
// Only densely packed tensors are currently supported.
|
|
int64_t *denseStrides = (int64_t *)alloca(rank * sizeof(int64_t));
|
|
int64_t *sizes = descriptor->sizes;
|
|
for (int64_t i = rank - 1, runningStride = 1; i >= 0; i--) {
|
|
denseStrides[i] = runningStride;
|
|
runningStride *= sizes[i];
|
|
}
|
|
uint64_t sizeBytes = sizes[0] * denseStrides[0] * elementSizeBytes;
|
|
int64_t *strides = &sizes[rank];
|
|
(void)strides;
|
|
for (unsigned i = 0; i < rank; ++i)
|
|
assert(strides[i] == denseStrides[i] &&
|
|
"Mismatch in computed dense strides");
|
|
|
|
auto *ptr = descriptor->data + descriptor->offset * elementSizeBytes;
|
|
mgpuMemHostRegister(ptr, sizeBytes);
|
|
}
|
|
|
|
extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuSetDefaultDevice(int32_t device) {
|
|
defaultDevice = device;
|
|
}
|