diff --git a/sycl-jit/CMakeLists.txt b/sycl-jit/CMakeLists.txt index 874856a63d363..b790455ceeed0 100644 --- a/sycl-jit/CMakeLists.txt +++ b/sycl-jit/CMakeLists.txt @@ -9,19 +9,23 @@ set(SYCL_JIT_BASE_DIR ${CMAKE_CURRENT_SOURCE_DIR}) # directories, similar to how clang/CMakeLists.txt does it. set(LLVM_SPIRV_INCLUDE_DIRS "${LLVM_MAIN_SRC_DIR}/../llvm-spirv/include") -# Set library-wide warning options. -set(SYCL_JIT_WARNING_FLAGS -Wall -Wextra) +if (NOT WIN32 AND NOT CYGWIN) + # Set library-wide warning options. + set(SYCL_JIT_WARNING_FLAGS -Wall -Wextra) -option(SYCL_JIT_ENABLE_WERROR "Treat all warnings as errors in SYCL kernel JIT library" ON) -if(SYCL_JIT_ENABLE_WERROR) - list(APPEND SYCL_JIT_WARNING_FLAGS -Werror) -endif(SYCL_JIT_ENABLE_WERROR) + option(SYCL_JIT_ENABLE_WERROR "Treat all warnings as errors in SYCL kernel JIT library" ON) + if(SYCL_JIT_ENABLE_WERROR) + list(APPEND SYCL_JIT_WARNING_FLAGS -Werror) + endif(SYCL_JIT_ENABLE_WERROR) +endif() -if(WIN32) - message(WARNING "Kernel JIT not yet supported on Windows") -else(WIN32) - add_subdirectory(common) - add_subdirectory(jit-compiler) - add_subdirectory(passes) + +add_subdirectory(common) +add_subdirectory(jit-compiler) +add_subdirectory(passes) + +# Loadable plugins for opt aren't supported on Windows, +# so we can't execute the tests. +if (NOT WIN32 AND NOT CYGWIN) add_subdirectory(test) -endif(WIN32) +endif() diff --git a/sycl-jit/jit-compiler/CMakeLists.txt b/sycl-jit/jit-compiler/CMakeLists.txt index 09af2de6853ae..6dc5154486c6f 100644 --- a/sycl-jit/jit-compiler/CMakeLists.txt +++ b/sycl-jit/jit-compiler/CMakeLists.txt @@ -40,6 +40,10 @@ add_llvm_library(sycl-jit clangSerialization ) +if(WIN32) + target_link_libraries(sycl-jit PRIVATE Shlwapi) +endif() + target_compile_options(sycl-jit PRIVATE ${SYCL_JIT_WARNING_FLAGS}) # Mark LLVM and SPIR-V headers as system headers to ignore warnigns in them. diff --git a/sycl-jit/jit-compiler/include/KernelFusion.h b/sycl-jit/jit-compiler/include/KernelFusion.h index f149e05692627..d3575f33189aa 100644 --- a/sycl-jit/jit-compiler/include/KernelFusion.h +++ b/sycl-jit/jit-compiler/include/KernelFusion.h @@ -9,6 +9,12 @@ #ifndef SYCL_FUSION_JIT_COMPILER_KERNELFUSION_H #define SYCL_FUSION_JIT_COMPILER_KERNELFUSION_H +#ifdef _WIN32 +#define KF_EXPORT_SYMBOL __declspec(dllexport) +#else +#define KF_EXPORT_SYMBOL +#endif + #include "Kernel.h" #include "Options.h" #include "Parameter.h" @@ -55,25 +61,31 @@ extern "C" { #ifdef __clang__ #pragma clang diagnostic ignored "-Wreturn-type-c-linkage" #endif // __clang__ -JITResult fuseKernels(View KernelInformation, - const char *FusedKernelName, - View Identities, - BarrierFlags BarriersFlags, - View Internalization, - View JITConstants); -JITResult materializeSpecConstants(const char *KernelName, - jit_compiler::SYCLKernelBinaryInfo &BinInfo, - View SpecConstBlob); +#ifdef _MSC_VER +#pragma warning(push) +#pragma warning(disable : 4190) +#endif // _MSC_VER + +KF_EXPORT_SYMBOL JITResult +fuseKernels(View KernelInformation, const char *FusedKernelName, + View Identities, BarrierFlags BarriersFlags, + View Internalization, + View JITConstants); + +KF_EXPORT_SYMBOL JITResult materializeSpecConstants( + const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo, + View SpecConstBlob); -JITResult compileSYCL(InMemoryFile SourceFile, View IncludeFiles, - View UserArgs); +KF_EXPORT_SYMBOL JITResult compileSYCL(InMemoryFile SourceFile, + View IncludeFiles, + View UserArgs); /// Clear all previously set options. -void resetJITConfiguration(); +KF_EXPORT_SYMBOL void resetJITConfiguration(); /// Add an option to the configuration. -void addToJITConfiguration(OptionStorage &&Opt); +KF_EXPORT_SYMBOL void addToJITConfiguration(OptionStorage &&Opt); } // end of extern "C" diff --git a/sycl-jit/jit-compiler/lib/KernelFusion.cpp b/sycl-jit/jit-compiler/lib/KernelFusion.cpp index 81037438061ae..86317c23e78de 100644 --- a/sycl-jit/jit-compiler/lib/KernelFusion.cpp +++ b/sycl-jit/jit-compiler/lib/KernelFusion.cpp @@ -71,10 +71,9 @@ static bool isTargetFormatSupported(BinaryFormat TargetFormat) { } } -extern "C" JITResult -materializeSpecConstants(const char *KernelName, - jit_compiler::SYCLKernelBinaryInfo &BinInfo, - View SpecConstBlob) { +extern "C" KF_EXPORT_SYMBOL JITResult materializeSpecConstants( + const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo, + View SpecConstBlob) { auto &JITCtx = JITContext::getInstance(); TargetInfo TargetInfo = ConfigHelper::get(); @@ -115,12 +114,11 @@ materializeSpecConstants(const char *KernelName, return JITResult{MaterializerKernelInfo}; } -extern "C" JITResult fuseKernels(View KernelInformation, - const char *FusedKernelName, - View Identities, - BarrierFlags BarriersFlags, - View Internalization, - View Constants) { +extern "C" KF_EXPORT_SYMBOL JITResult +fuseKernels(View KernelInformation, const char *FusedKernelName, + View Identities, BarrierFlags BarriersFlags, + View Internalization, + View Constants) { std::vector KernelsToFuse; llvm::transform(KernelInformation, std::back_inserter(KernelsToFuse), @@ -236,9 +234,9 @@ extern "C" JITResult fuseKernels(View KernelInformation, return JITResult{FusedKernelInfo}; } -extern "C" JITResult compileSYCL(InMemoryFile SourceFile, - View IncludeFiles, - View UserArgs) { +extern "C" KF_EXPORT_SYMBOL JITResult +compileSYCL(InMemoryFile SourceFile, View IncludeFiles, + View UserArgs) { auto ModuleOrErr = compileDeviceCode(SourceFile, IncludeFiles, UserArgs); if (!ModuleOrErr) { return errorToFusionResult(ModuleOrErr.takeError(), @@ -261,8 +259,10 @@ extern "C" JITResult compileSYCL(InMemoryFile SourceFile, return JITResult{Kernel}; } -extern "C" void resetJITConfiguration() { ConfigHelper::reset(); } +extern "C" KF_EXPORT_SYMBOL void resetJITConfiguration() { + ConfigHelper::reset(); +} -extern "C" void addToJITConfiguration(OptionStorage &&Opt) { +extern "C" KF_EXPORT_SYMBOL void addToJITConfiguration(OptionStorage &&Opt) { ConfigHelper::getConfig().set(std::move(Opt)); } diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index 6054cc5927eae..f694c8cd57136 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -20,6 +20,49 @@ static char X; // Dummy symbol, used as an anchor for `dlinfo` below. #endif +#ifdef _WIN32 +#include // For std::filesystem::path ( C++17 only ) +#include // For PathRemoveFileSpec +#include // For GetModuleFileName, HMODULE, DWORD, MAX_PATH + +// cribbed from sycl/source/detail/os_util.cpp +using OSModuleHandle = intptr_t; +static constexpr OSModuleHandle ExeModuleHandle = -1; +static OSModuleHandle getOSModuleHandle(const void *VirtAddr) { + HMODULE PhModule; + DWORD Flag = GET_MODULE_HANDLE_EX_FLAG_FROM_ADDRESS | + GET_MODULE_HANDLE_EX_FLAG_UNCHANGED_REFCOUNT; + auto LpModuleAddr = reinterpret_cast(VirtAddr); + if (!GetModuleHandleExA(Flag, LpModuleAddr, &PhModule)) { + // Expect the caller to check for zero and take + // necessary action + return 0; + } + if (PhModule == GetModuleHandleA(nullptr)) + return ExeModuleHandle; + return reinterpret_cast(PhModule); +} + +// cribbed from sycl/source/detail/os_util.cpp +/// Returns an absolute path where the object was found. +std::wstring getCurrentDSODir() { + wchar_t Path[MAX_PATH]; + auto Handle = getOSModuleHandle(reinterpret_cast(&getCurrentDSODir)); + DWORD Ret = GetModuleFileName( + reinterpret_cast(ExeModuleHandle == Handle ? 0 : Handle), Path, + MAX_PATH); + assert(Ret < MAX_PATH && "Path is longer than MAX_PATH?"); + assert(Ret > 0 && "GetModuleFileName failed"); + (void)Ret; + + BOOL RetCode = PathRemoveFileSpec(Path); + assert(RetCode && "PathRemoveFileSpec failed"); + (void)RetCode; + + return Path; +} +#endif // _WIN32 + static constexpr auto InvalidDPCPPRoot = ""; static const std::string &getDPCPPRoot() { @@ -42,6 +85,10 @@ static const std::string &getDPCPPRoot() { } #endif // _GNU_SOURCE +#ifdef _WIN32 + DPCPPRoot = std::filesystem::path(getCurrentDSODir()).parent_path().string(); +#endif // _WIN32 + // TODO: Implemenent other means of determining the DPCPP root, e.g. // evaluating the `CMPLR_ROOT` env. diff --git a/sycl-jit/passes/CMakeLists.txt b/sycl-jit/passes/CMakeLists.txt index 29e83d225d81b..b6cb30bd809f3 100644 --- a/sycl-jit/passes/CMakeLists.txt +++ b/sycl-jit/passes/CMakeLists.txt @@ -1,49 +1,54 @@ -# Module library for usage as library/pass-plugin with LLVM opt. -add_llvm_library(SYCLKernelJIT MODULE - SYCLFusionPasses.cpp - kernel-fusion/Builtins.cpp - kernel-fusion/SYCLKernelFusion.cpp - kernel-fusion/SYCLSpecConstMaterializer.cpp - kernel-info/SYCLKernelInfo.cpp - internalization/Internalization.cpp - syclcp/SYCLCP.cpp - cleanup/Cleanup.cpp - debug/PassDebug.cpp - target/TargetFusionInfo.cpp - - DEPENDS - intrinsics_gen -) +# See llvm/examples/Bye/CmakeLists.txt as to why this kind of loadable plugin libraries +# isn't supported on Windows. +if (NOT WIN32 AND NOT CYGWIN) + # Module library for usage as library/pass-plugin with LLVM opt. + add_llvm_library(SYCLKernelJIT MODULE + SYCLFusionPasses.cpp + kernel-fusion/Builtins.cpp + kernel-fusion/SYCLKernelFusion.cpp + kernel-fusion/SYCLSpecConstMaterializer.cpp + kernel-info/SYCLKernelInfo.cpp + internalization/Internalization.cpp + syclcp/SYCLCP.cpp + cleanup/Cleanup.cpp + debug/PassDebug.cpp + target/TargetFusionInfo.cpp + + DEPENDS + intrinsics_gen + ) + + target_compile_options(SYCLKernelJIT PRIVATE ${SYCL_JIT_WARNING_FLAGS}) + + # Mark LLVM headers as system headers to ignore warnigns in them. This + # classification remains intact even if the same path is added as a normal + # include path in GCC and Clang. + target_include_directories(SYCLKernelJIT + SYSTEM PRIVATE + ${LLVM_MAIN_INCLUDE_DIR} + ) + target_include_directories(SYCLKernelJIT + PUBLIC + ${CMAKE_CURRENT_SOURCE_DIR} + PRIVATE + ${SYCL_JIT_BASE_DIR}/common/include + ) + + target_link_libraries(SYCLKernelJIT + PRIVATE + sycl-jit-common + ) + + add_dependencies(SYCLKernelJIT sycl-headers) + + if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) + target_compile_definitions(SYCLKernelJIT PRIVATE JIT_SUPPORT_PTX) + endif() + + if("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD) + target_compile_definitions(SYCLKernelJIT PRIVATE JIT_SUPPORT_AMDGCN) + endif() -target_compile_options(SYCLKernelJIT PRIVATE ${SYCL_JIT_WARNING_FLAGS}) - -# Mark LLVM headers as system headers to ignore warnigns in them. This -# classification remains intact even if the same path is added as a normal -# include path in GCC and Clang. -target_include_directories(SYCLKernelJIT - SYSTEM PRIVATE - ${LLVM_MAIN_INCLUDE_DIR} -) -target_include_directories(SYCLKernelJIT - PUBLIC - ${CMAKE_CURRENT_SOURCE_DIR} - PRIVATE - ${SYCL_JIT_BASE_DIR}/common/include -) - -target_link_libraries(SYCLKernelJIT - PRIVATE - sycl-jit-common -) - -add_dependencies(SYCLKernelJIT sycl-headers) - -if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelJIT PRIVATE JIT_SUPPORT_PTX) -endif() - -if("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelJIT PRIVATE JIT_SUPPORT_AMDGCN) endif() # Static library for linking with the jit_compiler diff --git a/sycl-jit/passes/target/TargetFusionInfo.cpp b/sycl-jit/passes/target/TargetFusionInfo.cpp index af6589a1609ab..eacd339595432 100644 --- a/sycl-jit/passes/target/TargetFusionInfo.cpp +++ b/sycl-jit/passes/target/TargetFusionInfo.cpp @@ -356,9 +356,12 @@ class SPIRVTargetFusionInfo : public TargetFusionInfoImpl { Name = Name.drop_front(Name.find(SPIRVBuiltinPrefix) + SPIRVBuiltinPrefix.size()); // Check that Name does not start with any name in UnsafeBuiltIns - const auto *Iter = - std::upper_bound(UnsafeBuiltIns.begin(), UnsafeBuiltIns.end(), Name); - return Iter == UnsafeBuiltIns.begin() || !Name.starts_with(*(Iter - 1)); + for (const StringRef &Unsafe : UnsafeBuiltIns) { + if (Name.starts_with(Unsafe)) { + return false; + } + } + return true; } unsigned getIndexSpaceBuiltinBitwidth() const override { return 64; } diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index 385c70a04679f..a29bfc6310e39 100644 --- a/sycl/CMakeLists.txt +++ b/sycl/CMakeLists.txt @@ -30,11 +30,6 @@ endif() # Option to enable JIT, this in turn makes kernel fusion and spec constant # materialization possible. option(SYCL_ENABLE_EXTENSION_JIT "Enable extension to JIT kernels" ON) -if(SYCL_ENABLE_EXTENSION_JIT AND WIN32) - message(WARNING "Extension to JIT kernels not yet supported on Windows") - set(SYCL_ENABLE_EXTENSION_JIT OFF CACHE - BOOL "Extension to JIT kernels not yet supported on Windows" FORCE) -endif() if (NOT XPTI_INCLUDES) set(XPTI_INCLUDES ${CMAKE_CURRENT_SOURCE_DIR}/../xpti/include) diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index c0c22954822b7..19f1915943f05 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -15,6 +15,7 @@ #include #include #include +#include #include #include @@ -30,7 +31,12 @@ static inline void printPerformanceWarning(const std::string &Message) { jit_compiler::jit_compiler() { auto checkJITLibrary = [this]() -> bool { +#ifdef _WIN32 + static const std::string dir = sycl::detail::OSUtil::getCurrentDSODir(); + static const std::string JITLibraryName = dir + "\\" + "sycl-jit.dll"; +#else static const std::string JITLibraryName = "libsycl-jit.so"; +#endif void *LibraryPtr = sycl::detail::ur::loadOsLibrary(JITLibraryName); if (LibraryPtr == nullptr) { @@ -625,6 +631,7 @@ ur_kernel_handle_t jit_compiler::materializeSpecConstants( QueueImplPtr Queue, const RTDeviceBinaryImage *BinImage, const std::string &KernelName, const std::vector &SpecConstBlob) { +#ifndef _WIN32 if (!BinImage) { throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), "No suitable IR available for materializing"); @@ -716,6 +723,13 @@ ur_kernel_handle_t jit_compiler::materializeSpecConstants( } return NewKernel; +#else // _WIN32 + (void)Queue; + (void)BinImage; + (void)KernelName; + (void)SpecConstBlob; + return nullptr; +#endif // _WIN32 } std::unique_ptr diff --git a/sycl/source/detail/os_util.cpp b/sycl/source/detail/os_util.cpp index 860ad71f9f7ea..c8ee3b8f33c11 100644 --- a/sycl/source/detail/os_util.cpp +++ b/sycl/source/detail/os_util.cpp @@ -148,8 +148,8 @@ std::string OSUtil::getDirName(const char *Path) { #elif defined(__SYCL_RT_OS_WINDOWS) /// Returns an absolute path where the object was found. -// ur_win_proxy_loader.dll uses this same logic. If it is changed -// significantly, it might be wise to change it there too. +// ur_win_proxy_loader.dll and sycl-jit.dll use this same logic. If it is +// changed significantly, it might be wise to change it there too. std::string OSUtil::getCurrentDSODir() { char Path[MAX_PATH]; Path[0] = '\0'; diff --git a/sycl/source/detail/scheduler/scheduler.cpp b/sycl/source/detail/scheduler/scheduler.cpp index ac1d8ca44c5dc..efbbb52acab73 100644 --- a/sycl/source/detail/scheduler/scheduler.cpp +++ b/sycl/source/detail/scheduler/scheduler.cpp @@ -583,16 +583,16 @@ ur_kernel_handle_t Scheduler::completeSpecConstMaterialization( [[maybe_unused]] const RTDeviceBinaryImage *BinImage, [[maybe_unused]] const std::string &KernelName, [[maybe_unused]] std::vector &SpecConstBlob) { -#if SYCL_EXT_JIT_ENABLE +#if SYCL_EXT_JIT_ENABLE && !_WIN32 return detail::jit_compiler::get_instance().materializeSpecConstants( Queue, BinImage, KernelName, SpecConstBlob); -#else // SYCL_EXT_JIT_ENABLE +#else // SYCL_EXT_JIT_ENABLE && !_WIN32 if (detail::SYCLConfig::get() > 0) { std::cerr << "WARNING: Materialization of spec constants not supported by " "this build\n"; } return nullptr; -#endif // SYCL_EXT_JIT_ENABLE +#endif // SYCL_EXT_JIT_ENABLE && !_WIN32 } EventImplPtr Scheduler::addCommandGraphUpdate( diff --git a/sycl/source/detail/windows_os_utils.hpp b/sycl/source/detail/windows_os_utils.hpp index 690fbba46371c..f9141f2a4c5e1 100644 --- a/sycl/source/detail/windows_os_utils.hpp +++ b/sycl/source/detail/windows_os_utils.hpp @@ -10,6 +10,8 @@ #include +// ur_proxy_loader.dll and sycl-jit.dll use this same logic. If it changed +// significantly, then it'd be wise to update those versions as well. using OSModuleHandle = intptr_t; constexpr OSModuleHandle ExeModuleHandle = -1; inline OSModuleHandle getOSModuleHandle(const void *VirtAddr) { diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index cc45096b8564c..1588d55dfd2e2 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -9,10 +9,6 @@ // REQUIRES: (opencl || level_zero) // UNSUPPORTED: accelerator -// UNSUPPORTED: windows -// UNSUPPORTED-TRACKER: CMPLRLLVM-63166 -// in CMakeLists). - // RUN: %{build} -o %t.out // RUN: %{run} %t.out // RUN: %{l0_leak_check} %{run} %t.out @@ -95,7 +91,7 @@ void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) { sycl::free(usmPtr, Queue); } -void test_build_and_run() { +int test_build_and_run() { namespace syclex = sycl::ext::oneapi::experimental; using source_kb = sycl::kernel_bundle; using exe_kb = sycl::kernel_bundle; @@ -110,7 +106,7 @@ void test_build_and_run() { "kernel bundle extension: " << q.get_device().get_info() << std::endl; - return; + return -1; } // Create from source. @@ -141,12 +137,14 @@ void test_build_and_run() { // Test the kernels. test_1(q, k, 37 + 5); // ff_cp seeds 37. AddEm will add 5 more. + + return 0; } int main() { #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER - test_build_and_run(); + return test_build_and_run(); #else static_assert(false, "Kernel Compiler feature test macro undefined"); #endif diff --git a/sycl/test/CMakeLists.txt b/sycl/test/CMakeLists.txt index be2332d4a4c8a..144e2204361ef 100644 --- a/sycl/test/CMakeLists.txt +++ b/sycl/test/CMakeLists.txt @@ -96,6 +96,7 @@ add_lit_testsuite(check-sycl-dumps "Running ABI dump tests only" EXCLUDE_FROM_CHECK_ALL ) -if(SYCL_ENABLE_EXTENSION_JIT) +if(NOT WIN32 AND SYCL_ENABLE_EXTENSION_JIT) + # lit-based testing of JIT passes isn't supported on Windows. add_dependencies(check-sycl check-sycl-jit) -endif(SYCL_ENABLE_EXTENSION_JIT) +endif()