diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS index aed7bc010e48a..085ad2fca16d4 100644 --- a/.github/CODEOWNERS +++ b/.github/CODEOWNERS @@ -103,7 +103,7 @@ buildbot/ @intel/dpcpp-devops-reviewers devops/ @intel/dpcpp-devops-reviewers # Kernel fusion JIT compiler -sycl-fusion/ @intel/dpcpp-kernel-fusion-reviewers +sycl-jit/ @intel/dpcpp-kernel-fusion-reviewers sycl/doc/design/KernelFusionJIT.md @intel/dpcpp-kernel-fusion-reviewers sycl/doc/extensions/experimental/sycl_ext_codeplay_kernel_fusion.asciidoc @intel/dpcpp-kernel-fusion-reviewers sycl/include/sycl/ext/codeplay/experimental/fusion_properties.hpp @intel/dpcpp-kernel-fusion-reviewers diff --git a/.github/workflows/sycl-detect-changes.yml b/.github/workflows/sycl-detect-changes.yml index 585a2a0cff971..396cd424ff4ac 100644 --- a/.github/workflows/sycl-detect-changes.yml +++ b/.github/workflows/sycl-detect-changes.yml @@ -30,9 +30,9 @@ jobs: clang: &clang - *llvm - 'clang/**' - sycl_fusion: &sycl-fusion + sycl_jit: &sycl-jit - *llvm - - 'sycl-fusion/**' + - 'sycl-jit/**' xptifw: &xptifw - 'xptifw/**' libclc: &libclc @@ -41,7 +41,7 @@ jobs: - 'libclc/**' sycl: &sycl - *clang - - *sycl-fusion + - *sycl-jit - *llvm_spirv - *xptifw - *libclc @@ -84,7 +84,7 @@ jobs: return '${{ steps.changes.outputs.changes }}'; } // Treat everything as changed for huge PRs. - return ["llvm", "llvm_spirv", "clang", "sycl_fusion", "xptifw", "libclc", "sycl", "ci", "esimd"]; + return ["llvm", "llvm_spirv", "clang", "sycl_jit", "xptifw", "libclc", "sycl", "ci", "esimd"]; - run: echo '${{ steps.result.outputs.result }}' diff --git a/buildbot/configure.py b/buildbot/configure.py index fc89f8b7b00bf..692a64fd31256 100644 --- a/buildbot/configure.py +++ b/buildbot/configure.py @@ -31,10 +31,10 @@ def do_configure(args): libclc_amd_target_names = ";amdgcn--amdhsa" libclc_nvidia_target_names = ";nvptx64--nvidiacl" - sycl_enable_fusion = "OFF" - if not args.disable_fusion: - llvm_external_projects += ";sycl-fusion" - sycl_enable_fusion = "ON" + sycl_enable_jit = "OFF" + if not args.disable_jit: + llvm_external_projects += ";sycl-jit" + sycl_enable_jit = "ON" if args.llvm_external_projects: llvm_external_projects += ";" + args.llvm_external_projects.replace(",", ";") @@ -45,7 +45,7 @@ def do_configure(args): xpti_dir = os.path.join(abs_src_dir, "xpti") xptifw_dir = os.path.join(abs_src_dir, "xptifw") libdevice_dir = os.path.join(abs_src_dir, "libdevice") - fusion_dir = os.path.join(abs_src_dir, "sycl-fusion") + jit_dir = os.path.join(abs_src_dir, "sycl-jit") llvm_targets_to_build = args.host_target llvm_enable_projects = "clang;" + llvm_external_projects libclc_build_native = "OFF" @@ -174,7 +174,7 @@ def do_configure(args): "-DXPTI_SOURCE_DIR={}".format(xpti_dir), "-DLLVM_EXTERNAL_XPTIFW_SOURCE_DIR={}".format(xptifw_dir), "-DLLVM_EXTERNAL_LIBDEVICE_SOURCE_DIR={}".format(libdevice_dir), - "-DLLVM_EXTERNAL_SYCL_FUSION_SOURCE_DIR={}".format(fusion_dir), + "-DLLVM_EXTERNAL_SYCL_JIT_SOURCE_DIR={}".format(jit_dir), "-DLLVM_ENABLE_PROJECTS={}".format(llvm_enable_projects), "-DSYCL_BUILD_PI_HIP_PLATFORM={}".format(sycl_build_pi_hip_platform), "-DLLVM_BUILD_TOOLS=ON", @@ -189,7 +189,7 @@ def do_configure(args): "-DXPTI_ENABLE_WERROR={}".format(xpti_enable_werror), "-DSYCL_CLANG_EXTRA_FLAGS={}".format(sycl_clang_extra_flags), "-DSYCL_ENABLE_PLUGINS={}".format(";".join(set(sycl_enabled_plugins))), - "-DSYCL_ENABLE_KERNEL_FUSION={}".format(sycl_enable_fusion), + "-DSYCL_ENABLE_EXTENSION_JIT={}".format(sycl_enable_jit), "-DSYCL_ENABLE_MAJOR_RELEASE_PREVIEW_LIB={}".format(sycl_preview_lib), "-DBUG_REPORT_URL=https://github.com/intel/llvm/issues", ] @@ -379,9 +379,9 @@ def main(): help="Disable building of the SYCL runtime major release preview library", ) parser.add_argument( - "--disable-fusion", + "--disable-jit", action="store_true", - help="Disable the kernel fusion JIT compiler", + help="Disable the kernel JIT compiler for AMD and Nvidia", ) parser.add_argument( "--add_security_flags", diff --git a/sycl-fusion/CMakeLists.txt b/sycl-jit/CMakeLists.txt similarity index 67% rename from sycl-fusion/CMakeLists.txt rename to sycl-jit/CMakeLists.txt index 313635862acb3..874856a63d363 100644 --- a/sycl-fusion/CMakeLists.txt +++ b/sycl-jit/CMakeLists.txt @@ -10,15 +10,15 @@ set(SYCL_JIT_BASE_DIR ${CMAKE_CURRENT_SOURCE_DIR}) set(LLVM_SPIRV_INCLUDE_DIRS "${LLVM_MAIN_SRC_DIR}/../llvm-spirv/include") # Set library-wide warning options. -set(SYCL_FUSION_WARNING_FLAGS -Wall -Wextra) +set(SYCL_JIT_WARNING_FLAGS -Wall -Wextra) -option(SYCL_FUSION_ENABLE_WERROR "Treat all warnings as errors in SYCL kernel fusion library" ON) -if(SYCL_FUSION_ENABLE_WERROR) - list(APPEND SYCL_FUSION_WARNING_FLAGS -Werror) -endif(SYCL_FUSION_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) if(WIN32) - message(WARNING "Kernel fusion not yet supported on Windows") + message(WARNING "Kernel JIT not yet supported on Windows") else(WIN32) add_subdirectory(common) add_subdirectory(jit-compiler) diff --git a/sycl-fusion/README.md b/sycl-jit/README.md similarity index 100% rename from sycl-fusion/README.md rename to sycl-jit/README.md diff --git a/sycl-fusion/common/CMakeLists.txt b/sycl-jit/common/CMakeLists.txt similarity index 65% rename from sycl-fusion/common/CMakeLists.txt rename to sycl-jit/common/CMakeLists.txt index cb942fa0674da..1e51b3e336d8b 100644 --- a/sycl-fusion/common/CMakeLists.txt +++ b/sycl-jit/common/CMakeLists.txt @@ -1,26 +1,26 @@ -add_llvm_library(sycl-fusion-common +add_llvm_library(sycl-jit-common lib/NDRangesHelper.cpp LINK_COMPONENTS Support ) -target_compile_options(sycl-fusion-common PRIVATE ${SYCL_FUSION_WARNING_FLAGS}) +target_compile_options(sycl-jit-common 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(sycl-fusion-common +target_include_directories(sycl-jit-common SYSTEM PRIVATE ${LLVM_MAIN_INCLUDE_DIR} ) -target_include_directories(sycl-fusion-common +target_include_directories(sycl-jit-common PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/include ${CMAKE_CURRENT_SOURCE_DIR}/lib ) -add_dependencies(sycl-fusion-common sycl-headers) +add_dependencies(sycl-jit-common sycl-headers) if (BUILD_SHARED_LIBS) if(NOT MSVC AND NOT APPLE) @@ -28,8 +28,8 @@ if (BUILD_SHARED_LIBS) # are exported and confuse the drivers. set(linker_script "${CMAKE_CURRENT_SOURCE_DIR}/ld-version-script.txt") target_link_libraries( - sycl-fusion-common PRIVATE "-Wl,--version-script=${linker_script}") - set_target_properties(sycl-fusion-common + sycl-jit-common PRIVATE "-Wl,--version-script=${linker_script}") + set_target_properties(sycl-jit-common PROPERTIES LINK_DEPENDS ${linker_script}) diff --git a/sycl-fusion/common/include/DynArray.h b/sycl-jit/common/include/DynArray.h similarity index 100% rename from sycl-fusion/common/include/DynArray.h rename to sycl-jit/common/include/DynArray.h diff --git a/sycl-fusion/common/include/Kernel.h b/sycl-jit/common/include/Kernel.h similarity index 100% rename from sycl-fusion/common/include/Kernel.h rename to sycl-jit/common/include/Kernel.h diff --git a/sycl-fusion/common/include/View.h b/sycl-jit/common/include/View.h similarity index 100% rename from sycl-fusion/common/include/View.h rename to sycl-jit/common/include/View.h diff --git a/sycl-fusion/common/ld-version-script.txt b/sycl-jit/common/ld-version-script.txt similarity index 100% rename from sycl-fusion/common/ld-version-script.txt rename to sycl-jit/common/ld-version-script.txt diff --git a/sycl-fusion/common/lib/ModuleInfo.h b/sycl-jit/common/lib/ModuleInfo.h similarity index 100% rename from sycl-fusion/common/lib/ModuleInfo.h rename to sycl-jit/common/lib/ModuleInfo.h diff --git a/sycl-fusion/common/lib/NDRangesHelper.cpp b/sycl-jit/common/lib/NDRangesHelper.cpp similarity index 100% rename from sycl-fusion/common/lib/NDRangesHelper.cpp rename to sycl-jit/common/lib/NDRangesHelper.cpp diff --git a/sycl-fusion/common/lib/NDRangesHelper.h b/sycl-jit/common/lib/NDRangesHelper.h similarity index 100% rename from sycl-fusion/common/lib/NDRangesHelper.h rename to sycl-jit/common/lib/NDRangesHelper.h diff --git a/sycl-fusion/jit-compiler/CMakeLists.txt b/sycl-jit/jit-compiler/CMakeLists.txt similarity index 70% rename from sycl-fusion/jit-compiler/CMakeLists.txt rename to sycl-jit/jit-compiler/CMakeLists.txt index 5e60de84c0709..f49833b5ac7f8 100644 --- a/sycl-fusion/jit-compiler/CMakeLists.txt +++ b/sycl-jit/jit-compiler/CMakeLists.txt @@ -1,5 +1,5 @@ -add_llvm_library(sycl-fusion +add_llvm_library(sycl-jit lib/KernelFusion.cpp lib/translation/KernelTranslation.cpp lib/translation/SPIRVLLVMTranslation.cpp @@ -31,17 +31,17 @@ add_llvm_library(sycl-fusion ${LLVM_TARGETS_TO_BUILD} ) -target_compile_options(sycl-fusion PRIVATE ${SYCL_FUSION_WARNING_FLAGS}) +target_compile_options(sycl-jit PRIVATE ${SYCL_JIT_WARNING_FLAGS}) # Mark LLVM and SPIR-V headers as system headers to ignore warnigns in them. # This classification remains intact even if the same paths are added as normal # include paths in GCC and Clang. -target_include_directories(sycl-fusion +target_include_directories(sycl-jit SYSTEM PRIVATE ${LLVM_MAIN_INCLUDE_DIR} ${LLVM_SPIRV_INCLUDE_DIRS} ) -target_include_directories(sycl-fusion +target_include_directories(sycl-jit PUBLIC $ $ @@ -52,22 +52,22 @@ target_include_directories(sycl-fusion find_package(Threads REQUIRED) -target_link_libraries(sycl-fusion +target_link_libraries(sycl-jit PRIVATE - sycl-fusion-common + sycl-jit-common LLVMSPIRVLib - SYCLKernelFusionPasses + SYCLKernelJITPasses ${CMAKE_THREAD_LIBS_INIT} ) -add_dependencies(sycl-fusion sycl-headers) +add_dependencies(sycl-jit sycl-headers) if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(sycl-fusion PRIVATE FUSION_JIT_SUPPORT_PTX) + target_compile_definitions(sycl-jit PRIVATE JIT_SUPPORT_PTX) endif() if("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(sycl-fusion PRIVATE FUSION_JIT_SUPPORT_AMDGCN) + target_compile_definitions(sycl-jit PRIVATE JIT_SUPPORT_AMDGCN) endif() if(NOT MSVC AND NOT APPLE) @@ -75,6 +75,6 @@ if(NOT MSVC AND NOT APPLE) # are exported and confuse the drivers. set(linker_script "${CMAKE_CURRENT_SOURCE_DIR}/ld-version-script.txt") target_link_libraries( - sycl-fusion PRIVATE "-Wl,--version-script=${linker_script}") - set_target_properties(sycl-fusion PROPERTIES LINK_DEPENDS ${linker_script}) + sycl-jit PRIVATE "-Wl,--version-script=${linker_script}") + set_target_properties(sycl-jit PROPERTIES LINK_DEPENDS ${linker_script}) endif() diff --git a/sycl-fusion/jit-compiler/include/KernelFusion.h b/sycl-jit/jit-compiler/include/KernelFusion.h similarity index 100% rename from sycl-fusion/jit-compiler/include/KernelFusion.h rename to sycl-jit/jit-compiler/include/KernelFusion.h diff --git a/sycl-fusion/jit-compiler/include/Options.h b/sycl-jit/jit-compiler/include/Options.h similarity index 100% rename from sycl-fusion/jit-compiler/include/Options.h rename to sycl-jit/jit-compiler/include/Options.h diff --git a/sycl-fusion/jit-compiler/include/Parameter.h b/sycl-jit/jit-compiler/include/Parameter.h similarity index 100% rename from sycl-fusion/jit-compiler/include/Parameter.h rename to sycl-jit/jit-compiler/include/Parameter.h diff --git a/sycl-fusion/jit-compiler/ld-version-script.txt b/sycl-jit/jit-compiler/ld-version-script.txt similarity index 100% rename from sycl-fusion/jit-compiler/ld-version-script.txt rename to sycl-jit/jit-compiler/ld-version-script.txt diff --git a/sycl-fusion/jit-compiler/lib/KernelFusion.cpp b/sycl-jit/jit-compiler/lib/KernelFusion.cpp similarity index 97% rename from sycl-fusion/jit-compiler/lib/KernelFusion.cpp rename to sycl-jit/jit-compiler/lib/KernelFusion.cpp index 2277c99374919..27886566b5199 100644 --- a/sycl-fusion/jit-compiler/lib/KernelFusion.cpp +++ b/sycl-jit/jit-compiler/lib/KernelFusion.cpp @@ -52,18 +52,18 @@ static bool isTargetFormatSupported(BinaryFormat TargetFormat) { case BinaryFormat::SPIRV: return true; case BinaryFormat::PTX: { -#ifdef FUSION_JIT_SUPPORT_PTX +#ifdef JIT_SUPPORT_PTX return true; -#else // FUSION_JIT_SUPPORT_PTX +#else // JIT_SUPPORT_PTX return false; -#endif // FUSION_JIT_SUPPORT_PTX +#endif // JIT_SUPPORT_PTX } case BinaryFormat::AMDGCN: { -#ifdef FUSION_JIT_SUPPORT_AMDGCN +#ifdef JIT_SUPPORT_AMDGCN return true; -#else // FUSION_JIT_SUPPORT_AMDGCN +#else // JIT_SUPPORT_AMDGCN return false; -#endif // FUSION_JIT_SUPPORT_AMDGCN +#endif // JIT_SUPPORT_AMDGCN } default: return false; diff --git a/sycl-fusion/jit-compiler/lib/fusion/FusionHelper.cpp b/sycl-jit/jit-compiler/lib/fusion/FusionHelper.cpp similarity index 100% rename from sycl-fusion/jit-compiler/lib/fusion/FusionHelper.cpp rename to sycl-jit/jit-compiler/lib/fusion/FusionHelper.cpp diff --git a/sycl-fusion/jit-compiler/lib/fusion/FusionHelper.h b/sycl-jit/jit-compiler/lib/fusion/FusionHelper.h similarity index 100% rename from sycl-fusion/jit-compiler/lib/fusion/FusionHelper.h rename to sycl-jit/jit-compiler/lib/fusion/FusionHelper.h diff --git a/sycl-fusion/jit-compiler/lib/fusion/FusionPipeline.cpp b/sycl-jit/jit-compiler/lib/fusion/FusionPipeline.cpp similarity index 100% rename from sycl-fusion/jit-compiler/lib/fusion/FusionPipeline.cpp rename to sycl-jit/jit-compiler/lib/fusion/FusionPipeline.cpp diff --git a/sycl-fusion/jit-compiler/lib/fusion/FusionPipeline.h b/sycl-jit/jit-compiler/lib/fusion/FusionPipeline.h similarity index 100% rename from sycl-fusion/jit-compiler/lib/fusion/FusionPipeline.h rename to sycl-jit/jit-compiler/lib/fusion/FusionPipeline.h diff --git a/sycl-fusion/jit-compiler/lib/fusion/Hashing.h b/sycl-jit/jit-compiler/lib/fusion/Hashing.h similarity index 100% rename from sycl-fusion/jit-compiler/lib/fusion/Hashing.h rename to sycl-jit/jit-compiler/lib/fusion/Hashing.h diff --git a/sycl-fusion/jit-compiler/lib/fusion/JITContext.cpp b/sycl-jit/jit-compiler/lib/fusion/JITContext.cpp similarity index 100% rename from sycl-fusion/jit-compiler/lib/fusion/JITContext.cpp rename to sycl-jit/jit-compiler/lib/fusion/JITContext.cpp diff --git a/sycl-fusion/jit-compiler/lib/fusion/JITContext.h b/sycl-jit/jit-compiler/lib/fusion/JITContext.h similarity index 100% rename from sycl-fusion/jit-compiler/lib/fusion/JITContext.h rename to sycl-jit/jit-compiler/lib/fusion/JITContext.h diff --git a/sycl-fusion/jit-compiler/lib/fusion/ModuleHelper.cpp b/sycl-jit/jit-compiler/lib/fusion/ModuleHelper.cpp similarity index 100% rename from sycl-fusion/jit-compiler/lib/fusion/ModuleHelper.cpp rename to sycl-jit/jit-compiler/lib/fusion/ModuleHelper.cpp diff --git a/sycl-fusion/jit-compiler/lib/fusion/ModuleHelper.h b/sycl-jit/jit-compiler/lib/fusion/ModuleHelper.h similarity index 100% rename from sycl-fusion/jit-compiler/lib/fusion/ModuleHelper.h rename to sycl-jit/jit-compiler/lib/fusion/ModuleHelper.h diff --git a/sycl-fusion/jit-compiler/lib/helper/ConfigHelper.cpp b/sycl-jit/jit-compiler/lib/helper/ConfigHelper.cpp similarity index 100% rename from sycl-fusion/jit-compiler/lib/helper/ConfigHelper.cpp rename to sycl-jit/jit-compiler/lib/helper/ConfigHelper.cpp diff --git a/sycl-fusion/jit-compiler/lib/helper/ConfigHelper.h b/sycl-jit/jit-compiler/lib/helper/ConfigHelper.h similarity index 100% rename from sycl-fusion/jit-compiler/lib/helper/ConfigHelper.h rename to sycl-jit/jit-compiler/lib/helper/ConfigHelper.h diff --git a/sycl-fusion/jit-compiler/lib/helper/ErrorHandling.h b/sycl-jit/jit-compiler/lib/helper/ErrorHandling.h similarity index 100% rename from sycl-fusion/jit-compiler/lib/helper/ErrorHandling.h rename to sycl-jit/jit-compiler/lib/helper/ErrorHandling.h diff --git a/sycl-fusion/jit-compiler/lib/translation/KernelTranslation.cpp b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp similarity index 98% rename from sycl-fusion/jit-compiler/lib/translation/KernelTranslation.cpp rename to sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp index 2065be52f5ae2..2980d19e9f4f7 100644 --- a/sycl-fusion/jit-compiler/lib/translation/KernelTranslation.cpp +++ b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp @@ -231,13 +231,13 @@ llvm::Expected KernelTranslator::translateToPTX( SYCLKernelInfo &KernelInfo, llvm::Module &Mod, JITContext &JITCtx, [[maybe_unused]] const std::string &TargetCPU, [[maybe_unused]] const std::string &TargetFeatures) { -#ifndef FUSION_JIT_SUPPORT_PTX +#ifndef JIT_SUPPORT_PTX (void)KernelInfo; (void)Mod; (void)JITCtx; return createStringError(inconvertibleErrorCode(), "PTX translation not supported in this build"); -#else // FUSION_JIT_SUPPORT_PTX +#else // JIT_SUPPORT_PTX LLVMInitializeNVPTXTargetInfo(); LLVMInitializeNVPTXTarget(); LLVMInitializeNVPTXAsmPrinter(); @@ -306,20 +306,20 @@ llvm::Expected KernelTranslator::translateToPTX( } return &JITCtx.emplaceKernelBinary(std::move(PTXASM), BinaryFormat::PTX); -#endif // FUSION_JIT_SUPPORT_PTX +#endif // JIT_SUPPORT_PTX } llvm::Expected KernelTranslator::translateToAMDGCN( SYCLKernelInfo &KernelInfo, llvm::Module &Mod, JITContext &JITCtx, [[maybe_unused]] const std::string &TargetCPU, [[maybe_unused]] const std::string &TargetFeatures) { -#ifndef FUSION_JIT_SUPPORT_AMDGCN +#ifndef JIT_SUPPORT_AMDGCN (void)KernelInfo; (void)Mod; (void)JITCtx; return createStringError(inconvertibleErrorCode(), "AMDGPU translation not supported in this build"); -#else // FUSION_JIT_SUPPORT_AMDGCN +#else // JIT_SUPPORT_AMDGCN LLVMInitializeAMDGPUTargetInfo(); LLVMInitializeAMDGPUTarget(); @@ -383,5 +383,5 @@ llvm::Expected KernelTranslator::translateToAMDGCN( } return &JITCtx.emplaceKernelBinary(std::move(AMDObj), BinaryFormat::AMDGCN); -#endif // FUSION_JIT_SUPPORT_AMDGCN +#endif // JIT_SUPPORT_AMDGCN } diff --git a/sycl-fusion/jit-compiler/lib/translation/KernelTranslation.h b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.h similarity index 100% rename from sycl-fusion/jit-compiler/lib/translation/KernelTranslation.h rename to sycl-jit/jit-compiler/lib/translation/KernelTranslation.h diff --git a/sycl-fusion/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp b/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp similarity index 100% rename from sycl-fusion/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp rename to sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp diff --git a/sycl-fusion/jit-compiler/lib/translation/SPIRVLLVMTranslation.h b/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.h similarity index 100% rename from sycl-fusion/jit-compiler/lib/translation/SPIRVLLVMTranslation.h rename to sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.h diff --git a/sycl-fusion/passes/CMakeLists.txt b/sycl-jit/passes/CMakeLists.txt similarity index 64% rename from sycl-fusion/passes/CMakeLists.txt rename to sycl-jit/passes/CMakeLists.txt index 4e4d23379b07c..29e83d225d81b 100644 --- a/sycl-fusion/passes/CMakeLists.txt +++ b/sycl-jit/passes/CMakeLists.txt @@ -1,5 +1,5 @@ # Module library for usage as library/pass-plugin with LLVM opt. -add_llvm_library(SYCLKernelFusion MODULE +add_llvm_library(SYCLKernelJIT MODULE SYCLFusionPasses.cpp kernel-fusion/Builtins.cpp kernel-fusion/SYCLKernelFusion.cpp @@ -15,39 +15,39 @@ add_llvm_library(SYCLKernelFusion MODULE intrinsics_gen ) -target_compile_options(SYCLKernelFusion PRIVATE ${SYCL_FUSION_WARNING_FLAGS}) +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(SYCLKernelFusion +target_include_directories(SYCLKernelJIT SYSTEM PRIVATE ${LLVM_MAIN_INCLUDE_DIR} ) -target_include_directories(SYCLKernelFusion +target_include_directories(SYCLKernelJIT PUBLIC ${CMAKE_CURRENT_SOURCE_DIR} PRIVATE ${SYCL_JIT_BASE_DIR}/common/include ) -target_link_libraries(SYCLKernelFusion +target_link_libraries(SYCLKernelJIT PRIVATE - sycl-fusion-common + sycl-jit-common ) -add_dependencies(SYCLKernelFusion sycl-headers) +add_dependencies(SYCLKernelJIT sycl-headers) if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelFusion PRIVATE FUSION_JIT_SUPPORT_PTX) + target_compile_definitions(SYCLKernelJIT PRIVATE JIT_SUPPORT_PTX) endif() if("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelFusion PRIVATE FUSION_JIT_SUPPORT_AMDGCN) + target_compile_definitions(SYCLKernelJIT PRIVATE JIT_SUPPORT_AMDGCN) endif() # Static library for linking with the jit_compiler -add_llvm_library(SYCLKernelFusionPasses +add_llvm_library(SYCLKernelJITPasses SYCLFusionPasses.cpp kernel-fusion/Builtins.cpp kernel-fusion/SYCLKernelFusion.cpp @@ -71,33 +71,33 @@ add_llvm_library(SYCLKernelFusionPasses SYCLLowerIR ) -target_compile_options(SYCLKernelFusionPasses PRIVATE ${SYCL_FUSION_WARNING_FLAGS}) +target_compile_options(SYCLKernelJITPasses 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(SYCLKernelFusionPasses +target_include_directories(SYCLKernelJITPasses SYSTEM PRIVATE ${LLVM_MAIN_INCLUDE_DIR} ) -target_include_directories(SYCLKernelFusionPasses +target_include_directories(SYCLKernelJITPasses PUBLIC $ PRIVATE ${SYCL_JIT_BASE_DIR}/common/include ) -target_link_libraries(SYCLKernelFusionPasses +target_link_libraries(SYCLKernelJITPasses PRIVATE - sycl-fusion-common + sycl-jit-common ) -add_dependencies(SYCLKernelFusionPasses sycl-headers) +add_dependencies(SYCLKernelJITPasses sycl-headers) if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelFusionPasses PRIVATE FUSION_JIT_SUPPORT_PTX) + target_compile_definitions(SYCLKernelJITPasses PRIVATE JIT_SUPPORT_PTX) endif() if("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelFusionPasses PRIVATE FUSION_JIT_SUPPORT_AMDGCN) + target_compile_definitions(SYCLKernelJITPasses PRIVATE JIT_SUPPORT_AMDGCN) endif() diff --git a/sycl-fusion/passes/SYCLFusionPasses.cpp b/sycl-jit/passes/SYCLFusionPasses.cpp similarity index 100% rename from sycl-fusion/passes/SYCLFusionPasses.cpp rename to sycl-jit/passes/SYCLFusionPasses.cpp diff --git a/sycl-fusion/passes/cleanup/Cleanup.cpp b/sycl-jit/passes/cleanup/Cleanup.cpp similarity index 100% rename from sycl-fusion/passes/cleanup/Cleanup.cpp rename to sycl-jit/passes/cleanup/Cleanup.cpp index 599f2ca523e2b..43fc84f9f1b1e 100644 --- a/sycl-fusion/passes/cleanup/Cleanup.cpp +++ b/sycl-jit/passes/cleanup/Cleanup.cpp @@ -15,8 +15,8 @@ #include #include #include -#include #include +#include #include "Kernel.h" #include "kernel-info/SYCLKernelInfo.h" diff --git a/sycl-fusion/passes/cleanup/Cleanup.h b/sycl-jit/passes/cleanup/Cleanup.h similarity index 100% rename from sycl-fusion/passes/cleanup/Cleanup.h rename to sycl-jit/passes/cleanup/Cleanup.h diff --git a/sycl-fusion/passes/debug/PassDebug.cpp b/sycl-jit/passes/debug/PassDebug.cpp similarity index 100% rename from sycl-fusion/passes/debug/PassDebug.cpp rename to sycl-jit/passes/debug/PassDebug.cpp diff --git a/sycl-fusion/passes/debug/PassDebug.h b/sycl-jit/passes/debug/PassDebug.h similarity index 100% rename from sycl-fusion/passes/debug/PassDebug.h rename to sycl-jit/passes/debug/PassDebug.h diff --git a/sycl-fusion/passes/internalization/Internalization.cpp b/sycl-jit/passes/internalization/Internalization.cpp similarity index 100% rename from sycl-fusion/passes/internalization/Internalization.cpp rename to sycl-jit/passes/internalization/Internalization.cpp diff --git a/sycl-fusion/passes/internalization/Internalization.h b/sycl-jit/passes/internalization/Internalization.h similarity index 100% rename from sycl-fusion/passes/internalization/Internalization.h rename to sycl-jit/passes/internalization/Internalization.h diff --git a/sycl-fusion/passes/kernel-fusion/Builtins.cpp b/sycl-jit/passes/kernel-fusion/Builtins.cpp similarity index 100% rename from sycl-fusion/passes/kernel-fusion/Builtins.cpp rename to sycl-jit/passes/kernel-fusion/Builtins.cpp diff --git a/sycl-fusion/passes/kernel-fusion/Builtins.h b/sycl-jit/passes/kernel-fusion/Builtins.h similarity index 100% rename from sycl-fusion/passes/kernel-fusion/Builtins.h rename to sycl-jit/passes/kernel-fusion/Builtins.h diff --git a/sycl-fusion/passes/kernel-fusion/SYCLKernelFusion.cpp b/sycl-jit/passes/kernel-fusion/SYCLKernelFusion.cpp similarity index 100% rename from sycl-fusion/passes/kernel-fusion/SYCLKernelFusion.cpp rename to sycl-jit/passes/kernel-fusion/SYCLKernelFusion.cpp diff --git a/sycl-fusion/passes/kernel-fusion/SYCLKernelFusion.h b/sycl-jit/passes/kernel-fusion/SYCLKernelFusion.h similarity index 100% rename from sycl-fusion/passes/kernel-fusion/SYCLKernelFusion.h rename to sycl-jit/passes/kernel-fusion/SYCLKernelFusion.h diff --git a/sycl-fusion/passes/kernel-fusion/SYCLSpecConstMaterializer.cpp b/sycl-jit/passes/kernel-fusion/SYCLSpecConstMaterializer.cpp similarity index 100% rename from sycl-fusion/passes/kernel-fusion/SYCLSpecConstMaterializer.cpp rename to sycl-jit/passes/kernel-fusion/SYCLSpecConstMaterializer.cpp diff --git a/sycl-fusion/passes/kernel-fusion/SYCLSpecConstMaterializer.h b/sycl-jit/passes/kernel-fusion/SYCLSpecConstMaterializer.h similarity index 100% rename from sycl-fusion/passes/kernel-fusion/SYCLSpecConstMaterializer.h rename to sycl-jit/passes/kernel-fusion/SYCLSpecConstMaterializer.h diff --git a/sycl-fusion/passes/kernel-info/SYCLKernelInfo.cpp b/sycl-jit/passes/kernel-info/SYCLKernelInfo.cpp similarity index 99% rename from sycl-fusion/passes/kernel-info/SYCLKernelInfo.cpp rename to sycl-jit/passes/kernel-info/SYCLKernelInfo.cpp index ce0ac4d825130..36459a07a9aa7 100644 --- a/sycl-fusion/passes/kernel-info/SYCLKernelInfo.cpp +++ b/sycl-jit/passes/kernel-info/SYCLKernelInfo.cpp @@ -19,7 +19,7 @@ using namespace jit_compiler; llvm::AnalysisKey SYCLModuleInfoAnalysis::Key; // Keep this in sync with the enum definition in -// `sycl-fusion/common/include/Kernel.h`. +// `sycl-jit/common/include/Kernel.h`. static constexpr unsigned NumParameterKinds = 6; static constexpr std::array ParameterKindStrings = { diff --git a/sycl-fusion/passes/kernel-info/SYCLKernelInfo.h b/sycl-jit/passes/kernel-info/SYCLKernelInfo.h similarity index 100% rename from sycl-fusion/passes/kernel-info/SYCLKernelInfo.h rename to sycl-jit/passes/kernel-info/SYCLKernelInfo.h diff --git a/sycl-fusion/passes/metadata/MDParsing.h b/sycl-jit/passes/metadata/MDParsing.h similarity index 100% rename from sycl-fusion/passes/metadata/MDParsing.h rename to sycl-jit/passes/metadata/MDParsing.h diff --git a/sycl-fusion/passes/syclcp/SYCLCP.cpp b/sycl-jit/passes/syclcp/SYCLCP.cpp similarity index 100% rename from sycl-fusion/passes/syclcp/SYCLCP.cpp rename to sycl-jit/passes/syclcp/SYCLCP.cpp diff --git a/sycl-fusion/passes/syclcp/SYCLCP.h b/sycl-jit/passes/syclcp/SYCLCP.h similarity index 100% rename from sycl-fusion/passes/syclcp/SYCLCP.h rename to sycl-jit/passes/syclcp/SYCLCP.h diff --git a/sycl-fusion/passes/target/TargetFusionInfo.cpp b/sycl-jit/passes/target/TargetFusionInfo.cpp similarity index 99% rename from sycl-fusion/passes/target/TargetFusionInfo.cpp rename to sycl-jit/passes/target/TargetFusionInfo.cpp index db0aebe392b51..236d8ff7c4718 100644 --- a/sycl-fusion/passes/target/TargetFusionInfo.cpp +++ b/sycl-jit/passes/target/TargetFusionInfo.cpp @@ -711,7 +711,7 @@ class NVPTXAMDGCNTargetFusionInfoBase : public TargetFusionInfoImpl { // // NVPTXTargetFusionInfo // -#ifdef FUSION_JIT_SUPPORT_PTX +#ifdef JIT_SUPPORT_PTX class NVPTXTargetFusionInfo final : public NVPTXAMDGCNTargetFusionInfoBase { public: using NVPTXAMDGCNTargetFusionInfoBase::NVPTXAMDGCNTargetFusionInfoBase; @@ -813,12 +813,12 @@ class NVPTXTargetFusionInfo final : public NVPTXAMDGCNTargetFusionInfoBase { FusedNDRange); } }; -#endif // FUSION_JIT_SUPPORT_PTX +#endif // JIT_SUPPORT_PTX // // AMDGCNTargetFusionInfo // -#ifdef FUSION_JIT_SUPPORT_AMDGCN +#ifdef JIT_SUPPORT_AMDGCN class AMDGCNTargetFusionInfo final : public NVPTXAMDGCNTargetFusionInfoBase { using Base = NVPTXAMDGCNTargetFusionInfoBase; @@ -1099,18 +1099,18 @@ class AMDGCNTargetFusionInfo final : public NVPTXAMDGCNTargetFusionInfoBase { TargetFusionInfo::TargetFusionInfo(llvm::Module *Mod) { llvm::Triple Tri(Mod->getTargetTriple()); -#ifdef FUSION_JIT_SUPPORT_PTX +#ifdef JIT_SUPPORT_PTX if (Tri.isNVPTX()) { Impl = std::make_shared(Mod); return; } -#endif // FUSION_JIT_SUPPORT_PTX -#ifdef FUSION_JIT_SUPPORT_AMDGCN +#endif // JIT_SUPPORT_PTX +#ifdef JIT_SUPPORT_AMDGCN if (Tri.isAMDGCN()) { Impl = std::make_shared(Mod); return; } -#endif // FUSION_JIT_SUPPORT_AMDGCN +#endif // JIT_SUPPORT_AMDGCN if (Tri.isSPIRV() || Tri.isSPIR()) { Impl = std::make_shared(Mod); return; diff --git a/sycl-fusion/passes/target/TargetFusionInfo.h b/sycl-jit/passes/target/TargetFusionInfo.h similarity index 100% rename from sycl-fusion/passes/target/TargetFusionInfo.h rename to sycl-jit/passes/target/TargetFusionInfo.h diff --git a/sycl-fusion/test/CMakeLists.txt b/sycl-jit/test/CMakeLists.txt similarity index 72% rename from sycl-fusion/test/CMakeLists.txt rename to sycl-jit/test/CMakeLists.txt index 10493088c9b02..8f13b78354a25 100644 --- a/sycl-fusion/test/CMakeLists.txt +++ b/sycl-jit/test/CMakeLists.txt @@ -12,14 +12,14 @@ configure_lit_site_cfg( ) # Add a target to invoke tests via Ninja/make. -add_lit_testsuite(check-sycl-fusion-lit-tests "Running SYCL Fusion lit tests" +add_lit_testsuite(check-sycl-jit-lit-tests "Running SYCL Fusion lit tests" "${CMAKE_CURRENT_BINARY_DIR}" DEPENDS - SYCLKernelFusion + SYCLKernelJIT opt FileCheck ) -add_custom_target(check-sycl-fusion) -add_dependencies(check-sycl-fusion check-sycl-fusion-lit-tests) +add_custom_target(check-sycl-jit) +add_dependencies(check-sycl-jit check-sycl-jit-lit-tests) diff --git a/sycl-fusion/test/internalization/abort-promote-call.ll b/sycl-jit/test/internalization/abort-promote-call.ll similarity index 95% rename from sycl-fusion/test/internalization/abort-promote-call.ll rename to sycl-jit/test/internalization/abort-promote-call.ll index 5f96dec37675c..a1cf8a8a83520 100644 --- a/sycl-fusion/test/internalization/abort-promote-call.ll +++ b/sycl-jit/test/internalization/abort-promote-call.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" diff --git a/sycl-fusion/test/internalization/abort-promote-stored-ptr.ll b/sycl-jit/test/internalization/abort-promote-stored-ptr.ll similarity index 95% rename from sycl-fusion/test/internalization/abort-promote-stored-ptr.ll rename to sycl-jit/test/internalization/abort-promote-stored-ptr.ll index 623fcdce8b6ee..8a52e3a4dce97 100644 --- a/sycl-fusion/test/internalization/abort-promote-stored-ptr.ll +++ b/sycl-jit/test/internalization/abort-promote-stored-ptr.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" diff --git a/sycl-fusion/test/internalization/promote-local-nested.ll b/sycl-jit/test/internalization/promote-local-nested.ll similarity index 99% rename from sycl-fusion/test/internalization/promote-local-nested.ll rename to sycl-jit/test/internalization/promote-local-nested.ll index 50cd7d5a1bc2c..fc0405669950f 100644 --- a/sycl-fusion/test/internalization/promote-local-nested.ll +++ b/sycl-jit/test/internalization/promote-local-nested.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" diff --git a/sycl-fusion/test/internalization/promote-local-scalar.ll b/sycl-jit/test/internalization/promote-local-scalar.ll similarity index 99% rename from sycl-fusion/test/internalization/promote-local-scalar.ll rename to sycl-jit/test/internalization/promote-local-scalar.ll index 7bd7ae4903cff..732c844eead6d 100644 --- a/sycl-fusion/test/internalization/promote-local-scalar.ll +++ b/sycl-jit/test/internalization/promote-local-scalar.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" diff --git a/sycl-fusion/test/internalization/promote-local-vec.ll b/sycl-jit/test/internalization/promote-local-vec.ll similarity index 99% rename from sycl-fusion/test/internalization/promote-local-vec.ll rename to sycl-jit/test/internalization/promote-local-vec.ll index 984b60bfe8542..324e1d04113b1 100644 --- a/sycl-fusion/test/internalization/promote-local-vec.ll +++ b/sycl-jit/test/internalization/promote-local-vec.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" diff --git a/sycl-fusion/test/internalization/promote-private-nested.ll b/sycl-jit/test/internalization/promote-private-nested.ll similarity index 99% rename from sycl-fusion/test/internalization/promote-private-nested.ll rename to sycl-jit/test/internalization/promote-private-nested.ll index 5086a3cc1febc..13bc97f5a09f8 100644 --- a/sycl-fusion/test/internalization/promote-private-nested.ll +++ b/sycl-jit/test/internalization/promote-private-nested.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" target triple = "spir64-unknown-unknown" diff --git a/sycl-fusion/test/internalization/promote-private-non-unit-cuda.ll b/sycl-jit/test/internalization/promote-private-non-unit-cuda.ll similarity index 99% rename from sycl-fusion/test/internalization/promote-private-non-unit-cuda.ll rename to sycl-jit/test/internalization/promote-private-non-unit-cuda.ll index d786ab1529667..9246c34da94f4 100644 --- a/sycl-fusion/test/internalization/promote-private-non-unit-cuda.ll +++ b/sycl-jit/test/internalization/promote-private-non-unit-cuda.ll @@ -1,5 +1,5 @@ ; REQUIRES: cuda -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s ; This test is a reduced IR version of diff --git a/sycl-fusion/test/internalization/promote-private-non-unit-hip.ll b/sycl-jit/test/internalization/promote-private-non-unit-hip.ll similarity index 99% rename from sycl-fusion/test/internalization/promote-private-non-unit-hip.ll rename to sycl-jit/test/internalization/promote-private-non-unit-hip.ll index 8bff3b1536888..592987b6ae0d6 100644 --- a/sycl-fusion/test/internalization/promote-private-non-unit-hip.ll +++ b/sycl-jit/test/internalization/promote-private-non-unit-hip.ll @@ -1,5 +1,5 @@ ; REQUIRES: hip_amd -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s ; This test is the IR version of diff --git a/sycl-fusion/test/internalization/promote-private-non-unit.ll b/sycl-jit/test/internalization/promote-private-non-unit.ll similarity index 98% rename from sycl-fusion/test/internalization/promote-private-non-unit.ll rename to sycl-jit/test/internalization/promote-private-non-unit.ll index ac95f0606c628..0c978dd6a107a 100644 --- a/sycl-fusion/test/internalization/promote-private-non-unit.ll +++ b/sycl-jit/test/internalization/promote-private-non-unit.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s ; This test is a reduced IR version of diff --git a/sycl-fusion/test/internalization/promote-private-scalar.ll b/sycl-jit/test/internalization/promote-private-scalar.ll similarity index 99% rename from sycl-fusion/test/internalization/promote-private-scalar.ll rename to sycl-jit/test/internalization/promote-private-scalar.ll index e254520da9938..47d685c04e6c6 100644 --- a/sycl-fusion/test/internalization/promote-private-scalar.ll +++ b/sycl-jit/test/internalization/promote-private-scalar.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" diff --git a/sycl-fusion/test/internalization/promote-private-vec.ll b/sycl-jit/test/internalization/promote-private-vec.ll similarity index 99% rename from sycl-fusion/test/internalization/promote-private-vec.ll rename to sycl-jit/test/internalization/promote-private-vec.ll index 4bdef4a426350..74ad41a94826a 100644 --- a/sycl-fusion/test/internalization/promote-private-vec.ll +++ b/sycl-jit/test/internalization/promote-private-vec.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-internalization -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" diff --git a/sycl-fusion/test/kernel-fusion/check-failed-remapping-amdgpu.ll b/sycl-jit/test/kernel-fusion/check-failed-remapping-amdgpu.ll similarity index 96% rename from sycl-fusion/test/kernel-fusion/check-failed-remapping-amdgpu.ll rename to sycl-jit/test/kernel-fusion/check-failed-remapping-amdgpu.ll index 7015631a5d5a6..69b9ab3b7f293 100644 --- a/sycl-fusion/test/kernel-fusion/check-failed-remapping-amdgpu.ll +++ b/sycl-jit/test/kernel-fusion/check-failed-remapping-amdgpu.ll @@ -1,5 +1,5 @@ ; REQUIRES: hip_amd -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s ; This tests checks that kernel fusion fails when a not-remappable AMDGPU diff --git a/sycl-fusion/test/kernel-fusion/check-failed-remapping-cuda.ll b/sycl-jit/test/kernel-fusion/check-failed-remapping-cuda.ll similarity index 95% rename from sycl-fusion/test/kernel-fusion/check-failed-remapping-cuda.ll rename to sycl-jit/test/kernel-fusion/check-failed-remapping-cuda.ll index 4a6cc2777423b..d862ebce7d3fd 100644 --- a/sycl-fusion/test/kernel-fusion/check-failed-remapping-cuda.ll +++ b/sycl-jit/test/kernel-fusion/check-failed-remapping-cuda.ll @@ -1,5 +1,5 @@ ; REQUIRES: cuda -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s ; This tests checks that kernel fusion fails when a not-remappable PTX intrinsic diff --git a/sycl-fusion/test/kernel-fusion/check-failed-remapping.ll b/sycl-jit/test/kernel-fusion/check-failed-remapping.ll similarity index 96% rename from sycl-fusion/test/kernel-fusion/check-failed-remapping.ll rename to sycl-jit/test/kernel-fusion/check-failed-remapping.ll index 191f9aa5eb0d7..1571c890b85af 100644 --- a/sycl-fusion/test/kernel-fusion/check-failed-remapping.ll +++ b/sycl-jit/test/kernel-fusion/check-failed-remapping.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s ; This tests checks that kernel fusion fails when an unknown builtin diff --git a/sycl-fusion/test/kernel-fusion/check-remapping-amdgpu.ll b/sycl-jit/test/kernel-fusion/check-remapping-amdgpu.ll similarity index 99% rename from sycl-fusion/test/kernel-fusion/check-remapping-amdgpu.ll rename to sycl-jit/test/kernel-fusion/check-remapping-amdgpu.ll index 91428cae55773..52e4710fd0c2d 100644 --- a/sycl-fusion/test/kernel-fusion/check-remapping-amdgpu.ll +++ b/sycl-jit/test/kernel-fusion/check-remapping-amdgpu.ll @@ -1,5 +1,5 @@ ; REQUIRES: hip_amd -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s ; This tests checks that AMDGPU intrinsics are correctly remapped when fusing diff --git a/sycl-fusion/test/kernel-fusion/check-remapping-cuda.ll b/sycl-jit/test/kernel-fusion/check-remapping-cuda.ll similarity index 99% rename from sycl-fusion/test/kernel-fusion/check-remapping-cuda.ll rename to sycl-jit/test/kernel-fusion/check-remapping-cuda.ll index e1ec7127498ca..fb75d09307e4a 100644 --- a/sycl-fusion/test/kernel-fusion/check-remapping-cuda.ll +++ b/sycl-jit/test/kernel-fusion/check-remapping-cuda.ll @@ -1,5 +1,5 @@ ; REQUIRES: cuda -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s ; This tests checks that PTX intrinsics are correctly remapped when fusing diff --git a/sycl-fusion/test/kernel-fusion/check-remapping-interproc-amdgpu.ll b/sycl-jit/test/kernel-fusion/check-remapping-interproc-amdgpu.ll similarity index 99% rename from sycl-fusion/test/kernel-fusion/check-remapping-interproc-amdgpu.ll rename to sycl-jit/test/kernel-fusion/check-remapping-interproc-amdgpu.ll index 436ce4ad98b64..7a214eef3b2dd 100644 --- a/sycl-fusion/test/kernel-fusion/check-remapping-interproc-amdgpu.ll +++ b/sycl-jit/test/kernel-fusion/check-remapping-interproc-amdgpu.ll @@ -1,5 +1,5 @@ ; REQUIRES: hip_amd -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s ; This tests checks that AMDGPU intrinsics are correctly remapped when fusing diff --git a/sycl-fusion/test/kernel-fusion/check-remapping-interproc-cuda.ll b/sycl-jit/test/kernel-fusion/check-remapping-interproc-cuda.ll similarity index 99% rename from sycl-fusion/test/kernel-fusion/check-remapping-interproc-cuda.ll rename to sycl-jit/test/kernel-fusion/check-remapping-interproc-cuda.ll index 87f010e9eb62d..ecc5706d14074 100644 --- a/sycl-fusion/test/kernel-fusion/check-remapping-interproc-cuda.ll +++ b/sycl-jit/test/kernel-fusion/check-remapping-interproc-cuda.ll @@ -1,5 +1,5 @@ ; REQUIRES: cuda -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext \ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext \ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s ; This tests checks that PTX intrinsics are correctly remapped when fusing diff --git a/sycl-fusion/test/kernel-fusion/check-remapping-interproc.ll b/sycl-jit/test/kernel-fusion/check-remapping-interproc.ll similarity index 99% rename from sycl-fusion/test/kernel-fusion/check-remapping-interproc.ll rename to sycl-jit/test/kernel-fusion/check-remapping-interproc.ll index 539587f53291a..37e74e5dee5cf 100644 --- a/sycl-fusion/test/kernel-fusion/check-remapping-interproc.ll +++ b/sycl-jit/test/kernel-fusion/check-remapping-interproc.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s ; This tests checks that SPIR-V builtins are correctly remapped when fusing diff --git a/sycl-fusion/test/kernel-fusion/check-remapping.ll b/sycl-jit/test/kernel-fusion/check-remapping.ll similarity index 99% rename from sycl-fusion/test/kernel-fusion/check-remapping.ll rename to sycl-jit/test/kernel-fusion/check-remapping.ll index 251d38eee3ea8..d687c201a1d71 100644 --- a/sycl-fusion/test/kernel-fusion/check-remapping.ll +++ b/sycl-jit/test/kernel-fusion/check-remapping.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion -S %s\ ; RUN: | FileCheck %s diff --git a/sycl-fusion/test/kernel-fusion/required_work_group_size.ll b/sycl-jit/test/kernel-fusion/required_work_group_size.ll similarity index 97% rename from sycl-fusion/test/kernel-fusion/required_work_group_size.ll rename to sycl-jit/test/kernel-fusion/required_work_group_size.ll index f451623193420..75de88084f624 100644 --- a/sycl-fusion/test/kernel-fusion/required_work_group_size.ll +++ b/sycl-jit/test/kernel-fusion/required_work_group_size.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" diff --git a/sycl-fusion/test/kernel-fusion/two-kernels-no-identities.ll b/sycl-jit/test/kernel-fusion/two-kernels-no-identities.ll similarity index 98% rename from sycl-fusion/test/kernel-fusion/two-kernels-no-identities.ll rename to sycl-jit/test/kernel-fusion/two-kernels-no-identities.ll index 8e4ad8ce65def..5aaf275bc8ba5 100644 --- a/sycl-fusion/test/kernel-fusion/two-kernels-no-identities.ll +++ b/sycl-jit/test/kernel-fusion/two-kernels-no-identities.ll @@ -1,15 +1,15 @@ ; Check IR produced by fusion pass: -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes="sycl-kernel-fusion" -S %s\ ; RUN: | FileCheck %s --implicit-check-not fused_kernel --check-prefix FUSION ; Check metadata attached to kernel by fusion pass: -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion -S %s\ ; RUN: | FileCheck %s --check-prefix MD ; Check kernel information produced by fusion pass: -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion,print-sycl-module-info -disable-output -S %s\ ; RUN: | FileCheck %s --check-prefix INFO diff --git a/sycl-fusion/test/kernel-fusion/two-kernels-out-is-in.ll b/sycl-jit/test/kernel-fusion/two-kernels-out-is-in.ll similarity index 98% rename from sycl-fusion/test/kernel-fusion/two-kernels-out-is-in.ll rename to sycl-jit/test/kernel-fusion/two-kernels-out-is-in.ll index 2e617bc0e0483..ab2b1dcb0ed46 100644 --- a/sycl-fusion/test/kernel-fusion/two-kernels-out-is-in.ll +++ b/sycl-jit/test/kernel-fusion/two-kernels-out-is-in.ll @@ -1,15 +1,15 @@ ; Check IR produced by fusion pass: -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes="sycl-kernel-fusion" -S %s\ ; RUN: | FileCheck %s --implicit-check-not fused_kernel --check-prefix FUSION ; Check metadata attached to kernel by fusion pass: -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion -S %s\ ; RUN: | FileCheck %s --check-prefix MD ; Check kernel information produced by fusion pass: -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion,print-sycl-module-info -disable-output -S %s\ ; RUN: | FileCheck %s --check-prefix INFO diff --git a/sycl-fusion/test/kernel-fusion/work_group_size_hint.ll b/sycl-jit/test/kernel-fusion/work_group_size_hint.ll similarity index 97% rename from sycl-fusion/test/kernel-fusion/work_group_size_hint.ll rename to sycl-jit/test/kernel-fusion/work_group_size_hint.ll index 76588ed622d5d..b922071407fed 100644 --- a/sycl-fusion/test/kernel-fusion/work_group_size_hint.ll +++ b/sycl-jit/test/kernel-fusion/work_group_size_hint.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s diff --git a/sycl-fusion/test/kernel-info/kernel-info.ll b/sycl-jit/test/kernel-info/kernel-info.ll similarity index 96% rename from sycl-fusion/test/kernel-info/kernel-info.ll rename to sycl-jit/test/kernel-info/kernel-info.ll index 8efd594e0c6c7..87453e8066a99 100644 --- a/sycl-fusion/test/kernel-info/kernel-info.ll +++ b/sycl-jit/test/kernel-info/kernel-info.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=print-sycl-module-info -disable-output %s | FileCheck %s target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" diff --git a/sycl-fusion/test/lit.cfg.py b/sycl-jit/test/lit.cfg.py similarity index 100% rename from sycl-fusion/test/lit.cfg.py rename to sycl-jit/test/lit.cfg.py diff --git a/sycl-fusion/test/lit.site.cfg.py.in b/sycl-jit/test/lit.site.cfg.py.in similarity index 100% rename from sycl-fusion/test/lit.site.cfg.py.in rename to sycl-jit/test/lit.site.cfg.py.in diff --git a/sycl-fusion/test/materializer/basic.ll b/sycl-jit/test/materializer/basic.ll similarity index 94% rename from sycl-fusion/test/materializer/basic.ll rename to sycl-jit/test/materializer/basic.ll index 19bab11479671..524322116a384 100644 --- a/sycl-fusion/test/materializer/basic.ll +++ b/sycl-jit/test/materializer/basic.ll @@ -1,16 +1,16 @@ -; RUN: %if hip_amd %{ opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: %if hip_amd %{ opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: --mtriple amdgcn-amd-amdhsa -passes=sycl-spec-const-materializer -S %s |\ ; RUN: FileCheck --check-prefix=CHECK-MATERIALIZER %s %} -; RUN: %if cuda %{ opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: %if cuda %{ opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: --mtriple nvptx64-nvidia-cuda -passes=sycl-spec-const-materializer -S %s |\ ; RUN: FileCheck --check-prefix=CHECK-MATERIALIZER %s %} -; RUN: %if hip_amd %{ opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: %if hip_amd %{ opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: --mtriple amdgcn-amd-amdhsa -passes=sycl-spec-const-materializer,early-cse,adce -S %s |\ ; RUN: FileCheck --check-prefix=CHECK-MATERIALIZER-CSE %s %} -; RUN: %if cuda %{ opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: %if cuda %{ opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: --mtriple nvptx64-nvidia-cuda -passes=sycl-spec-const-materializer,early-cse,adce -S %s |\ ; RUN: FileCheck --check-prefix=CHECK-MATERIALIZER-CSE %s %} diff --git a/sycl-fusion/test/materializer/debug_output.ll b/sycl-jit/test/materializer/debug_output.ll similarity index 94% rename from sycl-fusion/test/materializer/debug_output.ll rename to sycl-jit/test/materializer/debug_output.ll index b9d0250118f1c..c1035d92c93dd 100644 --- a/sycl-fusion/test/materializer/debug_output.ll +++ b/sycl-jit/test/materializer/debug_output.ll @@ -1,13 +1,13 @@ ; https://github.com/intel/llvm/issues/14783 ; REQUIRES: TEMPORARY_DISABLED ; RUN: %if hip_amd %{ env SYCL_JIT_COMPILER_DEBUG="sycl-spec-const-materializer" opt\ -; RUN: -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext --mtriple amdgcn-amd-amdhsa\ +; RUN: -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext --mtriple amdgcn-amd-amdhsa\ ; RUN: -passes=sycl-spec-const-materializer,sccp -S %s 2> %t.stderr\ ; RUN: | FileCheck %s %} ; RUN: %if hip_amd %{ FileCheck --input-file=%t.stderr --check-prefix=CHECK-DEBUG %s %} ; RUN: %if cuda %{ env SYCL_JIT_COMPILER_DEBUG="sycl-spec-const-materializer" opt\ -; RUN: -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: --mtriple nvptx64-nvidia-cuda -passes=sycl-spec-const-materializer,sccp -S %s 2> %t.stderr\ ; RUN: | FileCheck %s %} ; RUN: %if hip_amd %{ FileCheck --input-file=%t.stderr --check-prefix=CHECK-DEBUG %s %} diff --git a/sycl-fusion/test/materializer/multi_type.ll b/sycl-jit/test/materializer/multi_type.ll similarity index 94% rename from sycl-fusion/test/materializer/multi_type.ll rename to sycl-jit/test/materializer/multi_type.ll index 112b685b959b5..f69bd057748e0 100644 --- a/sycl-fusion/test/materializer/multi_type.ll +++ b/sycl-jit/test/materializer/multi_type.ll @@ -1,16 +1,16 @@ -; RUN: %if hip_amd %{ opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: %if hip_amd %{ opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: --mtriple amdgcn-amd-amdhsa -passes=sycl-spec-const-materializer -S %s |\ ; RUN: FileCheck --check-prefix=CHECK-MATERIALIZER %s %} -; RUN: %if cuda %{ opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: %if cuda %{ opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: --mtriple nvptx64-nvidia-cuda -passes=sycl-spec-const-materializer -S %s |\ ; RUN: FileCheck --check-prefix=CHECK-MATERIALIZER %s %} -; RUN: %if hip_amd %{ opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: %if hip_amd %{ opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: --mtriple amdgcn-amd-amdhsa -passes=sycl-spec-const-materializer,early-cse -S %s |\ ; RUN: FileCheck --check-prefix=CHECK-MATERIALIZER-CSE %s %} -; RUN: %if cuda %{ opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: %if cuda %{ opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: --mtriple nvptx64-nvidia-cuda -passes=sycl-spec-const-materializer,early-cse -S %s |\ ; RUN: FileCheck --check-prefix=CHECK-MATERIALIZER-CSE %s %} diff --git a/sycl-fusion/test/syclcp/syclcp.ll b/sycl-jit/test/syclcp/syclcp.ll similarity index 99% rename from sycl-fusion/test/syclcp/syclcp.ll rename to sycl-jit/test/syclcp/syclcp.ll index d56a2858c13ec..c3c568cb91177 100644 --- a/sycl-fusion/test/syclcp/syclcp.ll +++ b/sycl-jit/test/syclcp/syclcp.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelFusion%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-cp -S %s | FileCheck %s target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024" target triple = "spir64-unknown-unknown" diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index 1d5859cca5016..fbf9e8c681f4e 100644 --- a/sycl/CMakeLists.txt +++ b/sycl/CMakeLists.txt @@ -22,12 +22,13 @@ if (NOT DEFINED SYCL_ENABLE_PLUGINS) set(SYCL_ENABLE_PLUGINS "opencl;level_zero") endif() -# Option to enable online kernel fusion via a JIT compiler -option(SYCL_ENABLE_KERNEL_FUSION "Enable kernel fusion via JIT compiler" ON) -if(SYCL_ENABLE_KERNEL_FUSION AND WIN32) - message(WARNING "Kernel fusion not yet supported on Windows") - set(SYCL_ENABLE_KERNEL_FUSION OFF CACHE - BOOL "Kernel fusion not yet supported on Windows" FORCE) +# 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) diff --git a/sycl/cmake/modules/AddSYCLUnitTest.cmake b/sycl/cmake/modules/AddSYCLUnitTest.cmake index 7843e4b54da1a..2464cb13013cf 100644 --- a/sycl/cmake/modules/AddSYCLUnitTest.cmake +++ b/sycl/cmake/modules/AddSYCLUnitTest.cmake @@ -78,9 +78,9 @@ macro(add_sycl_unittest test_dirname link_variant) add_dependencies(${test_dirname} ur_adapter_mock) - if(SYCL_ENABLE_KERNEL_FUSION) - target_link_libraries(${test_dirname} PRIVATE sycl-fusion) - endif(SYCL_ENABLE_KERNEL_FUSION) + if(SYCL_ENABLE_EXTENSION_JIT) + target_link_libraries(${test_dirname} PRIVATE sycl-jit) + endif(SYCL_ENABLE_EXTENSION_JIT) target_include_directories(${test_dirname} PRIVATE SYSTEM diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index ca412011b33b1..7d9a52bdc73a1 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -40,7 +40,7 @@ if("cuda" IN_LIST SYCL_ENABLE_PLUGINS) endif() if("hip" IN_LIST SYCL_ENABLE_PLUGINS) set(UR_BUILD_ADAPTER_HIP ON) - if (SYCL_ENABLE_KERNEL_FUSION) + if (SYCL_ENABLE_EXTENSION_JIT) set(UR_ENABLE_COMGR ON) endif() endif() diff --git a/sycl/doc/GetStartedGuide.md b/sycl/doc/GetStartedGuide.md index 50326f35aa45b..7933d33cda344 100644 --- a/sycl/doc/GetStartedGuide.md +++ b/sycl/doc/GetStartedGuide.md @@ -12,7 +12,7 @@ and a wide range of compute accelerators such as GPU and FPGA. * [Build DPC++ toolchain with support for NVIDIA CUDA](#build-dpc-toolchain-with-support-for-nvidia-cuda) * [Build DPC++ toolchain with support for HIP AMD](#build-dpc-toolchain-with-support-for-hip-amd) * [Build DPC++ toolchain with support for HIP NVIDIA](#build-dpc-toolchain-with-support-for-hip-nvidia) - * [Build DPC++ toolchain with support for runtime kernel fusion](#build-dpc-toolchain-with-support-for-runtime-kernel-fusion) + * [Build DPC++ toolchain with support for runtime kernel fusion and JIT compilation](#build-dpc-toolchain-with-support-for-runtime-kernel-fusion-and-jit-compilation) * [Build DPC++ toolchain with a custom Unified Runtime](#build-dpc-toolchain-with-a-custom-unified-runtime) * [Build Doxygen documentation](#build-doxygen-documentation) * [Deployment](#deployment) @@ -303,15 +303,17 @@ as well as the CUDA Runtime API to be installed, see Currently, this has only been tried on Linux, with ROCm 4.2.0 or 4.3.0, with CUDA 11, and using a GeForce 1060 device. -### Build DPC++ toolchain with support for runtime kernel fusion +### Build DPC++ toolchain with support for runtime kernel fusion and JIT compilation Support for the experimental SYCL extension for user-driven kernel fusion at -runtime is enabled by default. +runtime is enabled by default. The same mechanism is used to allow JIT +compilation of AMD and Nvidia kernels. -To disable support for this feature, follow the instructions for the Linux DPC++ -toolchain, but add the `--disable-fusion` flag. +To disable support for these features, follow the instructions for the Linux +DPC++ toolchain, but add the `--disable-jit` flag. -Kernel fusion is currently not yet supported on the Windows platform. +Both kernel fusion and JIT compilation of AMD and Nvidia kernels are currently +not yet supported on the Windows platform. ### Build Doxygen documentation diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index d02dbb725637a..6b483615ce0b2 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -89,7 +89,7 @@ function(add_sycl_rt_library LIB_NAME LIB_OBJ_NAME) target_compile_definitions(${LIB_OBJ_NAME} PRIVATE __SYCL_INTERNAL_API ) - if (WIN32) + if (WIN32) target_compile_definitions(${LIB_OBJ_NAME} PRIVATE __SYCL_BUILD_SYCL_DLL ) target_link_libraries(${LIB_NAME} PRIVATE shlwapi) if (ARG_IMPLIB_NAME) @@ -110,7 +110,7 @@ function(add_sycl_rt_library LIB_NAME LIB_OBJ_NAME) else() target_compile_options(${LIB_OBJ_NAME} PUBLIC -fvisibility=hidden -fvisibility-inlines-hidden) - + # Sycl math built-in macros cause a GCC 4.6 'note' to be output repeatedly. # => note: the ABI for passing parameters with 32-byte alignment has changed in GCC 4.6 # Seems to be no way to suppress it except use -Wno-psabi @@ -140,20 +140,20 @@ function(add_sycl_rt_library LIB_NAME LIB_OBJ_NAME) PRIVATE OpenCL-Headers ) - if(SYCL_ENABLE_KERNEL_FUSION) - if(NOT DEFINED LLVM_EXTERNAL_SYCL_FUSION_SOURCE_DIR) - message(FATAL_ERROR "Undefined LLVM_EXTERNAL_SYCL_FUSION_SOURCE_DIR variable: Must be set when SYCL fusion is enabled") + if(SYCL_ENABLE_EXTENSION_JIT) + if(NOT DEFINED LLVM_EXTERNAL_SYCL_JIT_SOURCE_DIR) + message(FATAL_ERROR "Undefined LLVM_EXTERNAL_SYCL_JIT_SOURCE_DIR variable: Must be set when extension to JIT SYCL kernels is enabled") endif() - set(SYCL_FUSION_INCLUDE_DIRS - ${LLVM_EXTERNAL_SYCL_FUSION_SOURCE_DIR}/common/include - ${LLVM_EXTERNAL_SYCL_FUSION_SOURCE_DIR}/jit-compiler/include) - add_dependencies(${LIB_NAME} sycl-fusion) - add_dependencies(${LIB_OBJ_NAME} sycl-fusion) - target_include_directories(${LIB_NAME} PRIVATE ${SYCL_FUSION_INCLUDE_DIRS}) - target_include_directories(${LIB_OBJ_NAME} PRIVATE ${SYCL_FUSION_INCLUDE_DIRS}) + set(SYCL_JIT_INCLUDE_DIRS + ${LLVM_EXTERNAL_SYCL_JIT_SOURCE_DIR}/common/include + ${LLVM_EXTERNAL_SYCL_JIT_SOURCE_DIR}/jit-compiler/include) + add_dependencies(${LIB_NAME} sycl-jit) + add_dependencies(${LIB_OBJ_NAME} sycl-jit) + target_include_directories(${LIB_NAME} PRIVATE ${SYCL_JIT_INCLUDE_DIRS}) + target_include_directories(${LIB_OBJ_NAME} PRIVATE ${SYCL_JIT_INCLUDE_DIRS}) set_property(GLOBAL APPEND PROPERTY SYCL_TOOLCHAIN_INSTALL_COMPONENTS - sycl-fusion) - endif(SYCL_ENABLE_KERNEL_FUSION) + sycl-jit) + endif(SYCL_ENABLE_EXTENSION_JIT) find_package(Threads REQUIRED) diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 909fc751772dc..cf78877d254c9 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -30,7 +30,7 @@ static inline void printPerformanceWarning(const std::string &Message) { jit_compiler::jit_compiler() { auto checkJITLibrary = [this]() -> bool { - static const std::string JITLibraryName = "libsycl-fusion.so"; + static const std::string JITLibraryName = "libsycl-jit.so"; void *LibraryPtr = sycl::detail::ur::loadOsLibrary(JITLibraryName); if (LibraryPtr == nullptr) { diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index d3724c49bad7e..ac0ba886240e3 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -88,8 +88,8 @@ inline namespace _V1 { #define SYCL_EXT_ONEAPI_COPY_OPTIMIZE 1 #define SYCL_EXT_ONEAPI_VIRTUAL_MEM 1 #define SYCL_EXT_ONEAPI_USM_MALLOC_PROPERTIES 1 -#cmakedefine01 SYCL_ENABLE_KERNEL_FUSION -#if SYCL_ENABLE_KERNEL_FUSION +#cmakedefine01 SYCL_ENABLE_EXTENSION_JIT +#if SYCL_ENABLE_EXTENSION_JIT #define SYCL_EXT_CODEPLAY_KERNEL_FUSION 1 #endif #define SYCL_EXT_INTEL_CACHE_CONFIG 1 diff --git a/sycl/test/CMakeLists.txt b/sycl/test/CMakeLists.txt index e41ef32d5f1ff..412a7bbe64f01 100644 --- a/sycl/test/CMakeLists.txt +++ b/sycl/test/CMakeLists.txt @@ -137,6 +137,6 @@ if(SYCL_BUILD_BACKEND_HIP) endif() endif() -if(SYCL_ENABLE_KERNEL_FUSION) - add_dependencies(check-sycl check-sycl-fusion) -endif(SYCL_ENABLE_KERNEL_FUSION) +if(SYCL_ENABLE_EXTENSION_JIT) + add_dependencies(check-sycl check-sycl-jit) +endif(SYCL_ENABLE_EXTENSION_JIT)