From 2d62212b06be4d712dcc461c430e12c73a561449 Mon Sep 17 00:00:00 2001 From: Kern Handa Date: Sun, 21 Feb 2021 01:13:48 -0800 Subject: [PATCH] [mlir] Export CUDA and Vulkan runtime wrappers on Windows Reviewed By: mehdi_amini Differential Revision: https://reviews.llvm.org/D97140 --- mlir/include/mlir/ExecutionEngine/RunnerUtils.h | 2 + .../mlir-cuda-runner/cuda-runtime-wrappers.cpp | 47 +++++++++++++--------- .../mlir-vulkan-runner/vulkan-runtime-wrappers.cpp | 8 +++- 3 files changed, 38 insertions(+), 19 deletions(-) diff --git a/mlir/include/mlir/ExecutionEngine/RunnerUtils.h b/mlir/include/mlir/ExecutionEngine/RunnerUtils.h index c854387..976cad3 100644 --- a/mlir/include/mlir/ExecutionEngine/RunnerUtils.h +++ b/mlir/include/mlir/ExecutionEngine/RunnerUtils.h @@ -330,6 +330,8 @@ int64_t verifyMemRef(UnrankedMemRefType &actual, extern "C" MLIR_RUNNERUTILS_EXPORT void _mlir_ciface_print_memref_i8(UnrankedMemRefType *M); extern "C" MLIR_RUNNERUTILS_EXPORT void +_mlir_ciface_print_memref_i32(UnrankedMemRefType *M); +extern "C" MLIR_RUNNERUTILS_EXPORT void _mlir_ciface_print_memref_f32(UnrankedMemRefType *M); extern "C" MLIR_RUNNERUTILS_EXPORT void _mlir_ciface_print_memref_f64(UnrankedMemRefType *M); diff --git a/mlir/tools/mlir-cuda-runner/cuda-runtime-wrappers.cpp b/mlir/tools/mlir-cuda-runner/cuda-runtime-wrappers.cpp index b8554bb..3cc2727 100644 --- a/mlir/tools/mlir-cuda-runner/cuda-runtime-wrappers.cpp +++ b/mlir/tools/mlir-cuda-runner/cuda-runtime-wrappers.cpp @@ -20,6 +20,12 @@ #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) \ @@ -56,18 +62,19 @@ private: CUcontext previous; }; -extern "C" CUmodule mgpuModuleLoad(void *data) { +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" void mgpuModuleUnload(CUmodule module) { +extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuModuleUnload(CUmodule module) { CUDA_REPORT_IF_ERROR(cuModuleUnload(module)); } -extern "C" CUfunction mgpuModuleGetFunction(CUmodule module, const char *name) { +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; @@ -76,52 +83,55 @@ extern "C" CUfunction mgpuModuleGetFunction(CUmodule module, const char *name) { // 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" 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) { +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" CUstream mgpuStreamCreate() { +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" void mgpuStreamDestroy(CUstream stream) { +extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuStreamDestroy(CUstream stream) { CUDA_REPORT_IF_ERROR(cuStreamDestroy(stream)); } -extern "C" void mgpuStreamSynchronize(CUstream stream) { +extern "C" MLIR_CUDA_WRAPPERS_EXPORT void +mgpuStreamSynchronize(CUstream stream) { CUDA_REPORT_IF_ERROR(cuStreamSynchronize(stream)); } -extern "C" void mgpuStreamWaitEvent(CUstream stream, CUevent event) { +extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuStreamWaitEvent(CUstream stream, + CUevent event) { CUDA_REPORT_IF_ERROR(cuStreamWaitEvent(stream, event, /*flags=*/0)); } -extern "C" CUevent mgpuEventCreate() { +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" void mgpuEventDestroy(CUevent event) { +extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuEventDestroy(CUevent event) { CUDA_REPORT_IF_ERROR(cuEventDestroy(event)); } -extern "C" void mgpuEventSynchronize(CUevent event) { +extern MLIR_CUDA_WRAPPERS_EXPORT "C" void mgpuEventSynchronize(CUevent event) { CUDA_REPORT_IF_ERROR(cuEventSynchronize(event)); } -extern "C" void mgpuEventRecord(CUevent event, CUstream stream) { +extern MLIR_CUDA_WRAPPERS_EXPORT "C" void mgpuEventRecord(CUevent event, + CUstream stream) { CUDA_REPORT_IF_ERROR(cuEventRecord(event, stream)); } @@ -147,14 +157,15 @@ extern "C" void mgpuMemcpy(void *dst, void *src, uint64_t sizeBytes, // Allows to register byte array with the CUDA runtime. Helpful until we have // transfer functions implemented. -extern "C" void mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) { +extern "C" MLIR_CUDA_WRAPPERS_EXPORT void +mgpuMemHostRegister(void *ptr, uint64_t sizeBytes) { ScopedContext scopedContext; CUDA_REPORT_IF_ERROR(cuMemHostRegister(ptr, sizeBytes, /*flags=*/0)); } // Allows to register a MemRef with the CUDA runtime. Helpful until we have // transfer functions implemented. -extern "C" void +extern "C" MLIR_CUDA_WRAPPERS_EXPORT void mgpuMemHostRegisterMemRef(int64_t rank, StridedMemRefType *descriptor, int64_t elementSizeBytes) { diff --git a/mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp b/mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp index ebc385d..c7ee02c 100644 --- a/mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp +++ b/mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp @@ -17,7 +17,12 @@ #include "VulkanRuntime.h" // Explicitly export entry points to the vulkan-runtime-wrapper. + +#ifdef _WIN32 +#define VULKAN_WRAPPER_SYMBOL_EXPORT __declspec(dllexport) +#else #define VULKAN_WRAPPER_SYMBOL_EXPORT __attribute__((visibility("default"))) +#endif // _WIN32 namespace { @@ -65,7 +70,8 @@ private: } // namespace -template struct MemRefDescriptor { +template +struct MemRefDescriptor { T *allocated; T *aligned; int64_t offset; -- 2.7.4