| //===- 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) |
| |
| // Make the primary context of device 0 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 0. |
| static CUcontext context = [] { |
| CUDA_REPORT_IF_ERROR(cuInit(/*flags=*/0)); |
| CUdevice device; |
| CUDA_REPORT_IF_ERROR(cuDeviceGet(&device, /*ordinal=*/0)); |
| 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); |
| } |