[Mlir-commits] [mlir] 2d62212 - [mlir] Export CUDA and Vulkan runtime wrappers on Windows
Kern Handa
llvmlistbot at llvm.org
Sun Feb 21 22:59:03 PST 2021
Author: Kern Handa
Date: 2021-02-21T22:58:55-08:00
New Revision: 2d62212b06be4d712dcc461c430e12c73a561449
URL: https://github.com/llvm/llvm-project/commit/2d62212b06be4d712dcc461c430e12c73a561449
DIFF: https://github.com/llvm/llvm-project/commit/2d62212b06be4d712dcc461c430e12c73a561449.diff
LOG: [mlir] Export CUDA and Vulkan runtime wrappers on Windows
Reviewed By: mehdi_amini
Differential Revision: https://reviews.llvm.org/D97140
Added:
Modified:
mlir/include/mlir/ExecutionEngine/RunnerUtils.h
mlir/tools/mlir-cuda-runner/cuda-runtime-wrappers.cpp
mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp
Removed:
################################################################################
diff --git a/mlir/include/mlir/ExecutionEngine/RunnerUtils.h b/mlir/include/mlir/ExecutionEngine/RunnerUtils.h
index c8543874b2fd..976cad333778 100644
--- a/mlir/include/mlir/ExecutionEngine/RunnerUtils.h
+++ b/mlir/include/mlir/ExecutionEngine/RunnerUtils.h
@@ -330,6 +330,8 @@ int64_t verifyMemRef(UnrankedMemRefType<T> &actual,
extern "C" MLIR_RUNNERUTILS_EXPORT void
_mlir_ciface_print_memref_i8(UnrankedMemRefType<int8_t> *M);
extern "C" MLIR_RUNNERUTILS_EXPORT void
+_mlir_ciface_print_memref_i32(UnrankedMemRefType<int32_t> *M);
+extern "C" MLIR_RUNNERUTILS_EXPORT void
_mlir_ciface_print_memref_f32(UnrankedMemRefType<float> *M);
extern "C" MLIR_RUNNERUTILS_EXPORT void
_mlir_ciface_print_memref_f64(UnrankedMemRefType<double> *M);
diff --git a/mlir/tools/mlir-cuda-runner/cuda-runtime-wrappers.cpp b/mlir/tools/mlir-cuda-runner/cuda-runtime-wrappers.cpp
index b8554bbc256b..3cc2727bd45f 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 @@ class ScopedContext {
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<char, 1> *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 ebc385da3248..c7ee02cdc9d8 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 @@ class VulkanRuntimeManager {
} // namespace
-template <typename T, int N> struct MemRefDescriptor {
+template <typename T, int N>
+struct MemRefDescriptor {
T *allocated;
T *aligned;
int64_t offset;
More information about the Mlir-commits
mailing list