[SYCL] SYCL_JIT support on Windows for kernel_compiler

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()

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.

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<SYCLKernelInfo> KernelInformation,
-                      const char *FusedKernelName,
-                      View<ParameterIdentity> Identities,
-                      BarrierFlags BarriersFlags,
-                      View<ParameterInternalization> Internalization,
-                      View<jit_compiler::JITConstant> JITConstants);
 
-JITResult materializeSpecConstants(const char *KernelName,
-                                   jit_compiler::SYCLKernelBinaryInfo &BinInfo,
-                                   View<unsigned char> SpecConstBlob);
+#ifdef _MSC_VER
+#pragma warning(push)
+#pragma warning(disable : 4190)
+#endif // _MSC_VER
+
+KF_EXPORT_SYMBOL JITResult
+fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
+            View<ParameterIdentity> Identities, BarrierFlags BarriersFlags,
+            View<ParameterInternalization> Internalization,
+            View<jit_compiler::JITConstant> JITConstants);
+
+KF_EXPORT_SYMBOL JITResult materializeSpecConstants(
+    const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo,
+    View<unsigned char> SpecConstBlob);
 
-JITResult compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
-                      View<const char *> UserArgs);
+KF_EXPORT_SYMBOL JITResult compileSYCL(InMemoryFile SourceFile,
+                                       View<InMemoryFile> IncludeFiles,
+                                       View<const char *> 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"
 

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<unsigned char> SpecConstBlob) {
+extern "C" KF_EXPORT_SYMBOL JITResult materializeSpecConstants(
+    const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo,
+    View<unsigned char> SpecConstBlob) {
   auto &JITCtx = JITContext::getInstance();
 
   TargetInfo TargetInfo = ConfigHelper::get<option::JITTargetInfo>();
@@ -115,12 +114,11 @@ materializeSpecConstants(const char *KernelName,
   return JITResult{MaterializerKernelInfo};
 }
 
-extern "C" JITResult fuseKernels(View<SYCLKernelInfo> KernelInformation,
-                                 const char *FusedKernelName,
-                                 View<ParameterIdentity> Identities,
-                                 BarrierFlags BarriersFlags,
-                                 View<ParameterInternalization> Internalization,
-                                 View<jit_compiler::JITConstant> Constants) {
+extern "C" KF_EXPORT_SYMBOL JITResult
+fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
+            View<ParameterIdentity> Identities, BarrierFlags BarriersFlags,
+            View<ParameterInternalization> Internalization,
+            View<jit_compiler::JITConstant> Constants) {
 
   std::vector<std::string> KernelsToFuse;
   llvm::transform(KernelInformation, std::back_inserter(KernelsToFuse),
@@ -236,9 +234,9 @@ extern "C" JITResult fuseKernels(View<SYCLKernelInfo> KernelInformation,
   return JITResult{FusedKernelInfo};
 }
 
-extern "C" JITResult compileSYCL(InMemoryFile SourceFile,
-                                 View<InMemoryFile> IncludeFiles,
-                                 View<const char *> UserArgs) {
+extern "C" KF_EXPORT_SYMBOL JITResult
+compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
+            View<const char *> 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));
 }

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 <filesystem> // For std::filesystem::path ( C++17 only )
+#include <shlwapi.h>  // For PathRemoveFileSpec
+#include <windows.h>  // 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<LPCSTR>(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<OSModuleHandle>(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<void *>(&getCurrentDSODir));
+  DWORD Ret = GetModuleFileName(
+      reinterpret_cast<HMODULE>(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 = "<invalid>";
 
 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.
 

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

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; }

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)

sycl/source/detail/jit_compiler.cpp

@@ -15,6 +15,7 @@
 #include <detail/kernel_impl.hpp>
 #include <detail/queue_impl.hpp>
 #include <detail/sycl_mem_obj_t.hpp>
+#include <sycl/detail/os_util.hpp>
 #include <sycl/detail/ur.hpp>
 #include <sycl/kernel_bundle.hpp>
 
@@ -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<unsigned char> &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<detail::CG>