From 544f1758f272a5429b8755e729dc2aaf335ccc10 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Thu, 25 Jul 2024 06:47:04 +0000 Subject: [PATCH 1/9] [SYCL] Rename sycl-fusion to sycl-jit This patch renames the directory (`sycl-fusion` to `sycl-jit`) as well as the involved library names (following the same principle: `fusion` to `jit`). To keep the user facing names consistent buildbot switches were updated. And finally the codeowners file had to reflect the directory change. --- .github/CODEOWNERS | 2 +- .github/workflows/sycl-detect-changes.yml | 8 ++--- buildbot/configure.py | 16 ++++----- {sycl-fusion => sycl-jit}/CMakeLists.txt | 0 {sycl-fusion => sycl-jit}/README.md | 0 .../common/CMakeLists.txt | 14 ++++---- .../common/include/DynArray.h | 0 .../common/include/Kernel.h | 0 .../common/include/View.h | 0 .../common/ld-version-script.txt | 0 .../common/lib/ModuleInfo.h | 0 .../common/lib/NDRangesHelper.cpp | 0 .../common/lib/NDRangesHelper.h | 0 .../jit-compiler/CMakeLists.txt | 24 ++++++------- .../jit-compiler/include/KernelFusion.h | 0 .../jit-compiler/include/Options.h | 0 .../jit-compiler/include/Parameter.h | 0 .../jit-compiler/ld-version-script.txt | 0 .../jit-compiler/lib/KernelFusion.cpp | 0 .../jit-compiler/lib/fusion/FusionHelper.cpp | 0 .../jit-compiler/lib/fusion/FusionHelper.h | 0 .../lib/fusion/FusionPipeline.cpp | 0 .../jit-compiler/lib/fusion/FusionPipeline.h | 0 .../jit-compiler/lib/fusion/Hashing.h | 0 .../jit-compiler/lib/fusion/JITContext.cpp | 0 .../jit-compiler/lib/fusion/JITContext.h | 0 .../jit-compiler/lib/fusion/ModuleHelper.cpp | 0 .../jit-compiler/lib/fusion/ModuleHelper.h | 0 .../jit-compiler/lib/helper/ConfigHelper.cpp | 0 .../jit-compiler/lib/helper/ConfigHelper.h | 0 .../jit-compiler/lib/helper/ErrorHandling.h | 0 .../lib/translation/KernelTranslation.cpp | 0 .../lib/translation/KernelTranslation.h | 0 .../lib/translation/SPIRVLLVMTranslation.cpp | 0 .../lib/translation/SPIRVLLVMTranslation.h | 0 .../passes/CMakeLists.txt | 36 +++++++++---------- .../passes/SYCLFusionPasses.cpp | 0 .../passes/cleanup/Cleanup.cpp | 2 +- .../passes/cleanup/Cleanup.h | 0 .../passes/debug/PassDebug.cpp | 0 .../passes/debug/PassDebug.h | 0 .../internalization/Internalization.cpp | 0 .../passes/internalization/Internalization.h | 0 .../passes/kernel-fusion/Builtins.cpp | 0 .../passes/kernel-fusion/Builtins.h | 0 .../passes/kernel-fusion/SYCLKernelFusion.cpp | 0 .../passes/kernel-fusion/SYCLKernelFusion.h | 0 .../SYCLSpecConstMaterializer.cpp | 0 .../kernel-fusion/SYCLSpecConstMaterializer.h | 0 .../passes/kernel-info/SYCLKernelInfo.cpp | 2 +- .../passes/kernel-info/SYCLKernelInfo.h | 0 .../passes/metadata/MDParsing.h | 0 .../passes/syclcp/SYCLCP.cpp | 0 .../passes/syclcp/SYCLCP.h | 0 .../passes/target/TargetFusionInfo.cpp | 0 .../passes/target/TargetFusionInfo.h | 0 {sycl-fusion => sycl-jit}/test/CMakeLists.txt | 8 ++--- .../internalization/abort-promote-call.ll | 2 +- .../abort-promote-stored-ptr.ll | 2 +- .../internalization/promote-local-nested.ll | 2 +- .../internalization/promote-local-scalar.ll | 2 +- .../test/internalization/promote-local-vec.ll | 2 +- .../internalization/promote-private-nested.ll | 2 +- .../promote-private-non-unit-cuda.ll | 2 +- .../promote-private-non-unit-hip.ll | 2 +- .../promote-private-non-unit.ll | 2 +- .../internalization/promote-private-scalar.ll | 2 +- .../internalization/promote-private-vec.ll | 2 +- .../check-failed-remapping-amdgpu.ll | 2 +- .../check-failed-remapping-cuda.ll | 2 +- .../kernel-fusion/check-failed-remapping.ll | 2 +- .../kernel-fusion/check-remapping-amdgpu.ll | 2 +- .../kernel-fusion/check-remapping-cuda.ll | 2 +- .../check-remapping-interproc-amdgpu.ll | 2 +- .../check-remapping-interproc-cuda.ll | 2 +- .../check-remapping-interproc.ll | 2 +- .../test/kernel-fusion/check-remapping.ll | 2 +- .../kernel-fusion/required_work_group_size.ll | 2 +- .../two-kernels-no-identities.ll | 6 ++-- .../kernel-fusion/two-kernels-out-is-in.ll | 6 ++-- .../kernel-fusion/work_group_size_hint.ll | 2 +- .../test/kernel-info/kernel-info.ll | 2 +- {sycl-fusion => sycl-jit}/test/lit.cfg.py | 0 .../test/lit.site.cfg.py.in | 0 .../test/materializer/basic.ll | 8 ++--- .../test/materializer/debug_output.ll | 4 +-- .../test/materializer/multi_type.ll | 8 ++--- .../test/syclcp/syclcp.ll | 2 +- sycl/cmake/modules/AddSYCLUnitTest.cmake | 2 +- sycl/doc/GetStartedGuide.md | 2 +- sycl/source/CMakeLists.txt | 6 ++-- sycl/source/detail/jit_compiler.cpp | 2 +- sycl/test/CMakeLists.txt | 2 +- 93 files changed, 103 insertions(+), 103 deletions(-) rename {sycl-fusion => sycl-jit}/CMakeLists.txt (100%) rename {sycl-fusion => sycl-jit}/README.md (100%) rename {sycl-fusion => sycl-jit}/common/CMakeLists.txt (65%) rename {sycl-fusion => sycl-jit}/common/include/DynArray.h (100%) rename {sycl-fusion => sycl-jit}/common/include/Kernel.h (100%) rename {sycl-fusion => sycl-jit}/common/include/View.h (100%) rename {sycl-fusion => sycl-jit}/common/ld-version-script.txt (100%) rename {sycl-fusion => sycl-jit}/common/lib/ModuleInfo.h (100%) rename {sycl-fusion => sycl-jit}/common/lib/NDRangesHelper.cpp (100%) rename {sycl-fusion => sycl-jit}/common/lib/NDRangesHelper.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/CMakeLists.txt (70%) rename {sycl-fusion => sycl-jit}/jit-compiler/include/KernelFusion.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/include/Options.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/include/Parameter.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/ld-version-script.txt (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/KernelFusion.cpp (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/fusion/FusionHelper.cpp (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/fusion/FusionHelper.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/fusion/FusionPipeline.cpp (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/fusion/FusionPipeline.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/fusion/Hashing.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/fusion/JITContext.cpp (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/fusion/JITContext.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/fusion/ModuleHelper.cpp (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/fusion/ModuleHelper.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/helper/ConfigHelper.cpp (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/helper/ConfigHelper.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/helper/ErrorHandling.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/translation/KernelTranslation.cpp (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/translation/KernelTranslation.h (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp (100%) rename {sycl-fusion => sycl-jit}/jit-compiler/lib/translation/SPIRVLLVMTranslation.h (100%) rename {sycl-fusion => sycl-jit}/passes/CMakeLists.txt (63%) rename {sycl-fusion => sycl-jit}/passes/SYCLFusionPasses.cpp (100%) rename {sycl-fusion => sycl-jit}/passes/cleanup/Cleanup.cpp (100%) rename {sycl-fusion => sycl-jit}/passes/cleanup/Cleanup.h (100%) rename {sycl-fusion => sycl-jit}/passes/debug/PassDebug.cpp (100%) rename {sycl-fusion => sycl-jit}/passes/debug/PassDebug.h (100%) rename {sycl-fusion => sycl-jit}/passes/internalization/Internalization.cpp (100%) rename {sycl-fusion => sycl-jit}/passes/internalization/Internalization.h (100%) rename {sycl-fusion => sycl-jit}/passes/kernel-fusion/Builtins.cpp (100%) rename {sycl-fusion => sycl-jit}/passes/kernel-fusion/Builtins.h (100%) rename {sycl-fusion => sycl-jit}/passes/kernel-fusion/SYCLKernelFusion.cpp (100%) rename {sycl-fusion => sycl-jit}/passes/kernel-fusion/SYCLKernelFusion.h (100%) rename {sycl-fusion => sycl-jit}/passes/kernel-fusion/SYCLSpecConstMaterializer.cpp (100%) rename {sycl-fusion => sycl-jit}/passes/kernel-fusion/SYCLSpecConstMaterializer.h (100%) rename {sycl-fusion => sycl-jit}/passes/kernel-info/SYCLKernelInfo.cpp (99%) rename {sycl-fusion => sycl-jit}/passes/kernel-info/SYCLKernelInfo.h (100%) rename {sycl-fusion => sycl-jit}/passes/metadata/MDParsing.h (100%) rename {sycl-fusion => sycl-jit}/passes/syclcp/SYCLCP.cpp (100%) rename {sycl-fusion => sycl-jit}/passes/syclcp/SYCLCP.h (100%) rename {sycl-fusion => sycl-jit}/passes/target/TargetFusionInfo.cpp (100%) rename {sycl-fusion => sycl-jit}/passes/target/TargetFusionInfo.h (100%) rename {sycl-fusion => sycl-jit}/test/CMakeLists.txt (72%) rename {sycl-fusion => sycl-jit}/test/internalization/abort-promote-call.ll (95%) rename {sycl-fusion => sycl-jit}/test/internalization/abort-promote-stored-ptr.ll (95%) rename {sycl-fusion => sycl-jit}/test/internalization/promote-local-nested.ll (99%) rename {sycl-fusion => sycl-jit}/test/internalization/promote-local-scalar.ll (99%) rename {sycl-fusion => sycl-jit}/test/internalization/promote-local-vec.ll (99%) rename {sycl-fusion => sycl-jit}/test/internalization/promote-private-nested.ll (99%) rename {sycl-fusion => sycl-jit}/test/internalization/promote-private-non-unit-cuda.ll (99%) rename {sycl-fusion => sycl-jit}/test/internalization/promote-private-non-unit-hip.ll (99%) rename {sycl-fusion => sycl-jit}/test/internalization/promote-private-non-unit.ll (98%) rename {sycl-fusion => sycl-jit}/test/internalization/promote-private-scalar.ll (99%) rename {sycl-fusion => sycl-jit}/test/internalization/promote-private-vec.ll (99%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/check-failed-remapping-amdgpu.ll (96%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/check-failed-remapping-cuda.ll (95%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/check-failed-remapping.ll (96%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/check-remapping-amdgpu.ll (99%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/check-remapping-cuda.ll (99%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/check-remapping-interproc-amdgpu.ll (99%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/check-remapping-interproc-cuda.ll (99%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/check-remapping-interproc.ll (99%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/check-remapping.ll (99%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/required_work_group_size.ll (97%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/two-kernels-no-identities.ll (98%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/two-kernels-out-is-in.ll (98%) rename {sycl-fusion => sycl-jit}/test/kernel-fusion/work_group_size_hint.ll (97%) rename {sycl-fusion => sycl-jit}/test/kernel-info/kernel-info.ll (96%) rename {sycl-fusion => sycl-jit}/test/lit.cfg.py (100%) rename {sycl-fusion => sycl-jit}/test/lit.site.cfg.py.in (100%) rename {sycl-fusion => sycl-jit}/test/materializer/basic.ll (94%) rename {sycl-fusion => sycl-jit}/test/materializer/debug_output.ll (94%) rename {sycl-fusion => sycl-jit}/test/materializer/multi_type.ll (94%) rename {sycl-fusion => sycl-jit}/test/syclcp/syclcp.ll (99%) diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS index 772af752ea60f..51697252ad387 100644 --- a/.github/CODEOWNERS +++ b/.github/CODEOWNERS @@ -104,7 +104,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 f172be352ba7d..bb7d97eb79a5e 100644 --- a/buildbot/configure.py +++ b/buildbot/configure.py @@ -24,10 +24,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(",", ";") @@ -38,7 +38,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' @@ -167,7 +167,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_FUSION_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", @@ -182,7 +182,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_KERNEL_FUSION={}".format(sycl_enable_jit), "-DSYCL_ENABLE_MAJOR_RELEASE_PREVIEW_LIB={}".format(sycl_preview_lib), "-DBUG_REPORT_URL=https://github.com/intel/llvm/issues", ] @@ -278,7 +278,7 @@ def main(): parser.add_argument("--ci-defaults", action="store_true", help="Enable default CI parameters") parser.add_argument("--enable-plugin", action='append', help="Enable SYCL plugin") parser.add_argument("--disable-preview-lib", action='store_true', help="Disable building of the SYCL runtime major release preview library") - parser.add_argument("--disable-fusion", action="store_true", help="Disable the kernel fusion JIT compiler") + parser.add_argument("--disable-jit", action="store_true", help="Disable the kernel JIT compiler for AMD and Nvidia") parser.add_argument("--add_security_flags", type=str, choices=['none', 'default', 'sanitize'], default=None, help="Enables security flags for compile & link. Two values are supported: 'default' and 'sanitize'. 'Sanitize' option is an extension of 'default' set.") parser.add_argument('--native-cpu-libclc-targets', help='Target triples for libclc, used by the Native CPU backend') args = parser.parse_args() diff --git a/sycl-fusion/CMakeLists.txt b/sycl-jit/CMakeLists.txt similarity index 100% rename from sycl-fusion/CMakeLists.txt rename to sycl-jit/CMakeLists.txt 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..f6a67a788785f 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_FUSION_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..9035e64d8dbf6 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_FUSION_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 FUSION_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 FUSION_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 100% rename from sycl-fusion/jit-compiler/lib/KernelFusion.cpp rename to sycl-jit/jit-compiler/lib/KernelFusion.cpp 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 100% rename from sycl-fusion/jit-compiler/lib/translation/KernelTranslation.cpp rename to sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp 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 63% rename from sycl-fusion/passes/CMakeLists.txt rename to sycl-jit/passes/CMakeLists.txt index e3e8af1518560..01e6645939a36 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_FUSION_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 FUSION_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 FUSION_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 @@ -70,33 +70,33 @@ add_llvm_library(SYCLKernelFusionPasses TargetParser ) -target_compile_options(SYCLKernelFusionPasses PRIVATE ${SYCL_FUSION_WARNING_FLAGS}) +target_compile_options(SYCLKernelJitPasses PRIVATE ${SYCL_FUSION_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 FUSION_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 FUSION_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 100% rename from sycl-fusion/passes/target/TargetFusionInfo.cpp rename to sycl-jit/passes/target/TargetFusionInfo.cpp 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..b7daf21d4b280 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..8504c7daf32d0 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..b4fc0e9f5d419 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..4284947840cb3 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..768f46a7a0ed7 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..fd78fd26395b8 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..27b3971ea6c48 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..ef984571bcd67 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 b08d7ba472e57..80dc1d0b489ff 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..86e9ad09d87da 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..7ab8b2d2d4ddc 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..ff294a5c092a0 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..79d8b2f7efeb4 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..126c3dedff640 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..bf4da304adbbd 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 2c88036564ccd..ff8084ff89a58 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..1a80f79c85fb9 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 8b7a44a147099..16f12aef59324 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..48093162d679a 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..11dfb72b8b795 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..077f5cf3dc029 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..a276d4844dc68 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..88f38bc0ae71e 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..0768c1fc089f8 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..9612926430687 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..27da7ddde3097 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..f195e34688530 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 b7949e7d3f401..6203b74194659 100644 --- a/sycl-fusion/test/materializer/debug_output.ll +++ b/sycl-jit/test/materializer/debug_output.ll @@ -1,11 +1,11 @@ ; 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..6821bd55244dc 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..db9726d008116 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/cmake/modules/AddSYCLUnitTest.cmake b/sycl/cmake/modules/AddSYCLUnitTest.cmake index 9571d43cc07c9..712a466ff0402 100644 --- a/sycl/cmake/modules/AddSYCLUnitTest.cmake +++ b/sycl/cmake/modules/AddSYCLUnitTest.cmake @@ -76,7 +76,7 @@ macro(add_sycl_unittest test_dirname link_variant) ) if(SYCL_ENABLE_KERNEL_FUSION) - target_link_libraries(${test_dirname} PRIVATE sycl-fusion) + target_link_libraries(${test_dirname} PRIVATE sycl-jit) endif(SYCL_ENABLE_KERNEL_FUSION) target_include_directories(${test_dirname} diff --git a/sycl/doc/GetStartedGuide.md b/sycl/doc/GetStartedGuide.md index 70d180e686c02..f60d2ba231d85 100644 --- a/sycl/doc/GetStartedGuide.md +++ b/sycl/doc/GetStartedGuide.md @@ -307,7 +307,7 @@ Support for the experimental SYCL extension for user-driven kernel fusion at runtime is enabled by default. To disable support for this feature, follow the instructions for the Linux DPC++ -toolchain, but add the `--disable-fusion` flag. +toolchain, but add the `--disable-jit` flag. Kernel fusion is currently not yet supported on the Windows platform. diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index 70b2e80c35dd0..0b64a66a3a4d4 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -147,12 +147,12 @@ function(add_sycl_rt_library LIB_NAME LIB_OBJ_NAME) 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) + add_dependencies(${LIB_NAME} sycl-jit) + add_dependencies(${LIB_OBJ_NAME} sycl-jit) target_include_directories(${LIB_NAME} PRIVATE ${SYCL_FUSION_INCLUDE_DIRS}) target_include_directories(${LIB_OBJ_NAME} PRIVATE ${SYCL_FUSION_INCLUDE_DIRS}) set_property(GLOBAL APPEND PROPERTY SYCL_TOOLCHAIN_INSTALL_COMPONENTS - sycl-fusion) + sycl-jit) endif(SYCL_ENABLE_KERNEL_FUSION) find_package(Threads REQUIRED) diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 48d1d0c3e6ba7..2715931801b9e 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::pi::loadOsLibrary(JITLibraryName); if (LibraryPtr == nullptr) { diff --git a/sycl/test/CMakeLists.txt b/sycl/test/CMakeLists.txt index e41ef32d5f1ff..f9b85395bb21d 100644 --- a/sycl/test/CMakeLists.txt +++ b/sycl/test/CMakeLists.txt @@ -138,5 +138,5 @@ if(SYCL_BUILD_BACKEND_HIP) endif() if(SYCL_ENABLE_KERNEL_FUSION) - add_dependencies(check-sycl check-sycl-fusion) + add_dependencies(check-sycl check-sycl-jit) endif(SYCL_ENABLE_KERNEL_FUSION) From 7ae1d184ad17c5d2ce7336acaeb628cfe66272db Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Thu, 25 Jul 2024 12:05:33 +0000 Subject: [PATCH 2/9] SYCLKernelJit -> SYCLKernelJIT --- sycl-jit/jit-compiler/CMakeLists.txt | 2 +- sycl-jit/passes/CMakeLists.txt | 32 +++++++++---------- sycl-jit/test/CMakeLists.txt | 2 +- .../internalization/abort-promote-call.ll | 2 +- .../abort-promote-stored-ptr.ll | 2 +- .../internalization/promote-local-nested.ll | 2 +- .../internalization/promote-local-scalar.ll | 2 +- .../test/internalization/promote-local-vec.ll | 2 +- .../internalization/promote-private-nested.ll | 2 +- .../promote-private-non-unit-cuda.ll | 2 +- .../promote-private-non-unit-hip.ll | 2 +- .../promote-private-non-unit.ll | 2 +- .../internalization/promote-private-scalar.ll | 2 +- .../internalization/promote-private-vec.ll | 2 +- .../check-failed-remapping-amdgpu.ll | 2 +- .../check-failed-remapping-cuda.ll | 2 +- .../kernel-fusion/check-failed-remapping.ll | 2 +- .../kernel-fusion/check-remapping-amdgpu.ll | 2 +- .../kernel-fusion/check-remapping-cuda.ll | 2 +- .../check-remapping-interproc-amdgpu.ll | 2 +- .../check-remapping-interproc-cuda.ll | 2 +- .../check-remapping-interproc.ll | 2 +- .../test/kernel-fusion/check-remapping.ll | 2 +- .../kernel-fusion/required_work_group_size.ll | 2 +- .../two-kernels-no-identities.ll | 6 ++-- .../kernel-fusion/two-kernels-out-is-in.ll | 6 ++-- .../kernel-fusion/work_group_size_hint.ll | 2 +- sycl-jit/test/kernel-info/kernel-info.ll | 2 +- sycl-jit/test/materializer/basic.ll | 8 ++--- sycl-jit/test/materializer/debug_output.ll | 4 +-- sycl-jit/test/materializer/multi_type.ll | 8 ++--- sycl-jit/test/syclcp/syclcp.ll | 2 +- 32 files changed, 58 insertions(+), 58 deletions(-) diff --git a/sycl-jit/jit-compiler/CMakeLists.txt b/sycl-jit/jit-compiler/CMakeLists.txt index 9035e64d8dbf6..7fdf73166c40e 100644 --- a/sycl-jit/jit-compiler/CMakeLists.txt +++ b/sycl-jit/jit-compiler/CMakeLists.txt @@ -56,7 +56,7 @@ target_link_libraries(sycl-jit PRIVATE sycl-jit-common LLVMSPIRVLib - SYCLKernelJitPasses + SYCLKernelJITPasses ${CMAKE_THREAD_LIBS_INIT} ) diff --git a/sycl-jit/passes/CMakeLists.txt b/sycl-jit/passes/CMakeLists.txt index 01e6645939a36..28c75f856d783 100644 --- a/sycl-jit/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(SYCLKernelJit MODULE +add_llvm_library(SYCLKernelJIT MODULE SYCLFusionPasses.cpp kernel-fusion/Builtins.cpp kernel-fusion/SYCLKernelFusion.cpp @@ -15,39 +15,39 @@ add_llvm_library(SYCLKernelJit MODULE intrinsics_gen ) -target_compile_options(SYCLKernelJit PRIVATE ${SYCL_FUSION_WARNING_FLAGS}) +target_compile_options(SYCLKernelJIT PRIVATE ${SYCL_FUSION_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 +target_include_directories(SYCLKernelJIT SYSTEM PRIVATE ${LLVM_MAIN_INCLUDE_DIR} ) -target_include_directories(SYCLKernelJit +target_include_directories(SYCLKernelJIT PUBLIC ${CMAKE_CURRENT_SOURCE_DIR} PRIVATE ${SYCL_JIT_BASE_DIR}/common/include ) -target_link_libraries(SYCLKernelJit +target_link_libraries(SYCLKernelJIT PRIVATE sycl-jit-common ) -add_dependencies(SYCLKernelJit sycl-headers) +add_dependencies(SYCLKernelJIT sycl-headers) if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelJit PRIVATE FUSION_JIT_SUPPORT_PTX) + target_compile_definitions(SYCLKernelJIT PRIVATE FUSION_JIT_SUPPORT_PTX) endif() if("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelJit PRIVATE FUSION_JIT_SUPPORT_AMDGCN) + target_compile_definitions(SYCLKernelJIT PRIVATE FUSION_JIT_SUPPORT_AMDGCN) endif() # Static library for linking with the jit_compiler -add_llvm_library(SYCLKernelJitPasses +add_llvm_library(SYCLKernelJITPasses SYCLFusionPasses.cpp kernel-fusion/Builtins.cpp kernel-fusion/SYCLKernelFusion.cpp @@ -70,33 +70,33 @@ add_llvm_library(SYCLKernelJitPasses TargetParser ) -target_compile_options(SYCLKernelJitPasses PRIVATE ${SYCL_FUSION_WARNING_FLAGS}) +target_compile_options(SYCLKernelJITPasses PRIVATE ${SYCL_FUSION_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(SYCLKernelJitPasses +target_include_directories(SYCLKernelJITPasses SYSTEM PRIVATE ${LLVM_MAIN_INCLUDE_DIR} ) -target_include_directories(SYCLKernelJitPasses +target_include_directories(SYCLKernelJITPasses PUBLIC $ PRIVATE ${SYCL_JIT_BASE_DIR}/common/include ) -target_link_libraries(SYCLKernelJitPasses +target_link_libraries(SYCLKernelJITPasses PRIVATE sycl-jit-common ) -add_dependencies(SYCLKernelJitPasses sycl-headers) +add_dependencies(SYCLKernelJITPasses sycl-headers) if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelJitPasses PRIVATE FUSION_JIT_SUPPORT_PTX) + target_compile_definitions(SYCLKernelJITPasses PRIVATE FUSION_JIT_SUPPORT_PTX) endif() if("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelJitPasses PRIVATE FUSION_JIT_SUPPORT_AMDGCN) + target_compile_definitions(SYCLKernelJITPasses PRIVATE FUSION_JIT_SUPPORT_AMDGCN) endif() diff --git a/sycl-jit/test/CMakeLists.txt b/sycl-jit/test/CMakeLists.txt index b7daf21d4b280..8f13b78354a25 100644 --- a/sycl-jit/test/CMakeLists.txt +++ b/sycl-jit/test/CMakeLists.txt @@ -16,7 +16,7 @@ add_lit_testsuite(check-sycl-jit-lit-tests "Running SYCL Fusion lit tests" "${CMAKE_CURRENT_BINARY_DIR}" DEPENDS - SYCLKernelJit + SYCLKernelJIT opt FileCheck ) diff --git a/sycl-jit/test/internalization/abort-promote-call.ll b/sycl-jit/test/internalization/abort-promote-call.ll index 8504c7daf32d0..a1cf8a8a83520 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/abort-promote-stored-ptr.ll b/sycl-jit/test/internalization/abort-promote-stored-ptr.ll index b4fc0e9f5d419..8a52e3a4dce97 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/promote-local-nested.ll b/sycl-jit/test/internalization/promote-local-nested.ll index 4284947840cb3..fc0405669950f 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/promote-local-scalar.ll b/sycl-jit/test/internalization/promote-local-scalar.ll index 768f46a7a0ed7..732c844eead6d 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/promote-local-vec.ll b/sycl-jit/test/internalization/promote-local-vec.ll index fd78fd26395b8..324e1d04113b1 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/promote-private-nested.ll b/sycl-jit/test/internalization/promote-private-nested.ll index 27b3971ea6c48..13bc97f5a09f8 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/promote-private-non-unit-cuda.ll b/sycl-jit/test/internalization/promote-private-non-unit-cuda.ll index ef984571bcd67..9246c34da94f4 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/promote-private-non-unit-hip.ll b/sycl-jit/test/internalization/promote-private-non-unit-hip.ll index 80dc1d0b489ff..3e0ae0e97eac1 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/promote-private-non-unit.ll b/sycl-jit/test/internalization/promote-private-non-unit.ll index 86e9ad09d87da..0c978dd6a107a 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/promote-private-scalar.ll b/sycl-jit/test/internalization/promote-private-scalar.ll index 7ab8b2d2d4ddc..47d685c04e6c6 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/internalization/promote-private-vec.ll b/sycl-jit/test/internalization/promote-private-vec.ll index ff294a5c092a0..74ad41a94826a 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/check-failed-remapping-amdgpu.ll b/sycl-jit/test/kernel-fusion/check-failed-remapping-amdgpu.ll index 79d8b2f7efeb4..69b9ab3b7f293 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/check-failed-remapping-cuda.ll b/sycl-jit/test/kernel-fusion/check-failed-remapping-cuda.ll index 126c3dedff640..d862ebce7d3fd 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/check-failed-remapping.ll b/sycl-jit/test/kernel-fusion/check-failed-remapping.ll index bf4da304adbbd..1571c890b85af 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/check-remapping-amdgpu.ll b/sycl-jit/test/kernel-fusion/check-remapping-amdgpu.ll index ff8084ff89a58..9bc3205f39441 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/check-remapping-cuda.ll b/sycl-jit/test/kernel-fusion/check-remapping-cuda.ll index 1a80f79c85fb9..fb75d09307e4a 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/check-remapping-interproc-amdgpu.ll b/sycl-jit/test/kernel-fusion/check-remapping-interproc-amdgpu.ll index 16f12aef59324..df4b25ee3a400 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/check-remapping-interproc-cuda.ll b/sycl-jit/test/kernel-fusion/check-remapping-interproc-cuda.ll index 48093162d679a..ecc5706d14074 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/check-remapping-interproc.ll b/sycl-jit/test/kernel-fusion/check-remapping-interproc.ll index 11dfb72b8b795..37e74e5dee5cf 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/check-remapping.ll b/sycl-jit/test/kernel-fusion/check-remapping.ll index 077f5cf3dc029..d687c201a1d71 100644 --- a/sycl-jit/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/SYCLKernelJit%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion -S %s\ ; RUN: | FileCheck %s diff --git a/sycl-jit/test/kernel-fusion/required_work_group_size.ll b/sycl-jit/test/kernel-fusion/required_work_group_size.ll index a276d4844dc68..75de88084f624 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/kernel-fusion/two-kernels-no-identities.ll b/sycl-jit/test/kernel-fusion/two-kernels-no-identities.ll index 88f38bc0ae71e..5aaf275bc8ba5 100644 --- a/sycl-jit/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/SYCLKernelJit%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/SYCLKernelJit%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/SYCLKernelJit%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-jit/test/kernel-fusion/two-kernels-out-is-in.ll b/sycl-jit/test/kernel-fusion/two-kernels-out-is-in.ll index 0768c1fc089f8..ab2b1dcb0ed46 100644 --- a/sycl-jit/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/SYCLKernelJit%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/SYCLKernelJit%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/SYCLKernelJit%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-jit/test/kernel-fusion/work_group_size_hint.ll b/sycl-jit/test/kernel-fusion/work_group_size_hint.ll index 9612926430687..b922071407fed 100644 --- a/sycl-jit/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/SYCLKernelJit%shlibext\ +; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJIT%shlibext\ ; RUN: -passes=sycl-kernel-fusion -S %s | FileCheck %s diff --git a/sycl-jit/test/kernel-info/kernel-info.ll b/sycl-jit/test/kernel-info/kernel-info.ll index 27da7ddde3097..87453e8066a99 100644 --- a/sycl-jit/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/SYCLKernelJit%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-jit/test/materializer/basic.ll b/sycl-jit/test/materializer/basic.ll index f195e34688530..524322116a384 100644 --- a/sycl-jit/test/materializer/basic.ll +++ b/sycl-jit/test/materializer/basic.ll @@ -1,16 +1,16 @@ -; RUN: %if hip_amd %{ opt -load-pass-plugin %shlibdir/SYCLKernelJit%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/SYCLKernelJit%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/SYCLKernelJit%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/SYCLKernelJit%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-jit/test/materializer/debug_output.ll b/sycl-jit/test/materializer/debug_output.ll index 6203b74194659..9f389a2e8e321 100644 --- a/sycl-jit/test/materializer/debug_output.ll +++ b/sycl-jit/test/materializer/debug_output.ll @@ -1,11 +1,11 @@ ; RUN: %if hip_amd %{ env SYCL_JIT_COMPILER_DEBUG="sycl-spec-const-materializer" opt\ -; RUN: -load-pass-plugin %shlibdir/SYCLKernelJit%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/SYCLKernelJit%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-jit/test/materializer/multi_type.ll b/sycl-jit/test/materializer/multi_type.ll index 6821bd55244dc..f69bd057748e0 100644 --- a/sycl-jit/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/SYCLKernelJit%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/SYCLKernelJit%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/SYCLKernelJit%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/SYCLKernelJit%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-jit/test/syclcp/syclcp.ll b/sycl-jit/test/syclcp/syclcp.ll index db9726d008116..c3c568cb91177 100644 --- a/sycl-jit/test/syclcp/syclcp.ll +++ b/sycl-jit/test/syclcp/syclcp.ll @@ -1,4 +1,4 @@ -; RUN: opt -load-pass-plugin %shlibdir/SYCLKernelJit%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" From 14a9bc82624ea4ae4eae27e1a80edf94fcf090db Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Thu, 25 Jul 2024 12:42:59 +0000 Subject: [PATCH 3/9] LLVM_EXTERNAL_SYCL_FUSION_SOURCE_DIR -> LLVM_EXTERNAL_SYCL_JIT_SOURCE_DIR SYCL_ENABLE_KERNEL_FUSION -> SYCL_ENABLE_EXTENSION_JIT SYCL_FUSION_INCLUDE_DIRS -> SYCL_JIT_INCLUDE_DIRS --- buildbot/configure.py | 4 ++-- sycl/CMakeLists.txt | 13 +++++++------ sycl/cmake/modules/AddSYCLUnitTest.cmake | 4 ++-- sycl/plugins/hip/CMakeLists.txt | 8 ++++---- sycl/source/CMakeLists.txt | 22 +++++++++++----------- sycl/source/feature_test.hpp.in | 4 ++-- sycl/test/CMakeLists.txt | 4 ++-- 7 files changed, 30 insertions(+), 29 deletions(-) diff --git a/buildbot/configure.py b/buildbot/configure.py index bb7d97eb79a5e..1d380d95ded02 100644 --- a/buildbot/configure.py +++ b/buildbot/configure.py @@ -167,7 +167,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(jit_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", @@ -182,7 +182,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_jit), + "-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", ] diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index ef13946a5b802..ff86be3c8e148 100644 --- a/sycl/CMakeLists.txt +++ b/sycl/CMakeLists.txt @@ -144,12 +144,13 @@ install(DIRECTORY ${OpenCL_INCLUDE_DIR}/CL DESTINATION ${SYCL_INCLUDE_DIR}/sycl COMPONENT OpenCL-Headers) -# 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() # Option for enabling building the SYCL major release preview library. diff --git a/sycl/cmake/modules/AddSYCLUnitTest.cmake b/sycl/cmake/modules/AddSYCLUnitTest.cmake index 712a466ff0402..b12bd6674919e 100644 --- a/sycl/cmake/modules/AddSYCLUnitTest.cmake +++ b/sycl/cmake/modules/AddSYCLUnitTest.cmake @@ -75,9 +75,9 @@ macro(add_sycl_unittest test_dirname link_variant) ${SYCL_LINK_LIBS} ) - if(SYCL_ENABLE_KERNEL_FUSION) + if(SYCL_ENABLE_EXTENSION_JIT) target_link_libraries(${test_dirname} PRIVATE sycl-jit) - endif(SYCL_ENABLE_KERNEL_FUSION) + endif(SYCL_ENABLE_EXTENSION_JIT) target_include_directories(${test_dirname} PRIVATE SYSTEM diff --git a/sycl/plugins/hip/CMakeLists.txt b/sycl/plugins/hip/CMakeLists.txt index b84b7091ff213..ba713801a6487 100644 --- a/sycl/plugins/hip/CMakeLists.txt +++ b/sycl/plugins/hip/CMakeLists.txt @@ -11,11 +11,11 @@ set(SYCL_BUILD_PI_HIP_ROCM_DIR "/opt/rocm" CACHE STRING "ROCm installation dir") set(SYCL_BUILD_PI_HIP_INCLUDE_DIR "" CACHE STRING "Override HIP include dir path (set to \"\" for default behavior)") set(SYCL_BUILD_PI_HIP_HSA_INCLUDE_DIR "" CACHE STRING "Override HSA include dir path (set to \"\" for default behavior)") -if(SYCL_ENABLE_KERNEL_FUSION) +if(SYCL_ENABLE_EXTENSION_JIT) set(SYCL_ENABLE_COMGR ON) -else(SYCL_ENABLE_KERNEL_FUSION) +else(SYCL_ENABLE_EXTENSION_JIT) set(SYCL_ENABLE_COMGR OFF) -endif(SYCL_ENABLE_KERNEL_FUSION) +endif(SYCL_ENABLE_EXTENSION_JIT) if("${SYCL_BUILD_PI_HIP_INCLUDE_DIR}" STREQUAL "") set(PI_HIP_INCLUDE_DIR "${SYCL_BUILD_PI_HIP_ROCM_DIR}/include") @@ -165,7 +165,7 @@ if("${SYCL_BUILD_PI_HIP_PLATFORM}" STREQUAL "AMD") INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "${HIP_HEADERS}" ) target_link_libraries(pi_hip PUBLIC amd_comgr) - target_compile_definitions(pi_hip PRIVATE SYCL_ENABLE_KERNEL_FUSION) + target_compile_definitions(pi_hip PRIVATE SYCL_ENABLE_EXTENSION_JIT) endif(SYCL_ENABLE_COMGR) # Set HIP define to select AMD platform diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index 0b64a66a3a4d4..48f2fb8fb5b2c 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) + 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_FUSION_INCLUDE_DIRS}) - target_include_directories(${LIB_OBJ_NAME} PRIVATE ${SYCL_FUSION_INCLUDE_DIRS}) + 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-jit) - endif(SYCL_ENABLE_KERNEL_FUSION) + endif(SYCL_ENABLE_EXTENSION_JIT) find_package(Threads REQUIRED) diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 660efa9fa98a7..63aa7c54927bd 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 f9b85395bb21d..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) +if(SYCL_ENABLE_EXTENSION_JIT) add_dependencies(check-sycl check-sycl-jit) -endif(SYCL_ENABLE_KERNEL_FUSION) +endif(SYCL_ENABLE_EXTENSION_JIT) From b0ac4af8b8e4ed4ab069b5f3ffe1c2835aad3291 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Thu, 25 Jul 2024 13:09:50 +0000 Subject: [PATCH 4/9] SYCL_FUSION_WARNING_FLAGS -> SYCL_JIT_WARNING_FLAGS SYCL_FUSION_ENABLE_WERROR -> SYCL_JIT_ENABLE_WERROR FUSION_JIT_SUPPORT_PTX -> JIT_SUPPORT_PTX FUSION_JIT_SUPPORT_AMDGCN -> JIT_SUPPORT_AMDGCN --- sycl-jit/CMakeLists.txt | 12 ++++++------ sycl-jit/common/CMakeLists.txt | 2 +- sycl-jit/jit-compiler/CMakeLists.txt | 6 +++--- sycl-jit/jit-compiler/lib/KernelFusion.cpp | 12 ++++++------ .../lib/translation/KernelTranslation.cpp | 12 ++++++------ sycl-jit/passes/CMakeLists.txt | 12 ++++++------ sycl-jit/passes/target/TargetFusionInfo.cpp | 14 +++++++------- 7 files changed, 35 insertions(+), 35 deletions(-) diff --git a/sycl-jit/CMakeLists.txt b/sycl-jit/CMakeLists.txt index 313635862acb3..874856a63d363 100644 --- a/sycl-jit/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-jit/common/CMakeLists.txt b/sycl-jit/common/CMakeLists.txt index f6a67a788785f..1e51b3e336d8b 100644 --- a/sycl-jit/common/CMakeLists.txt +++ b/sycl-jit/common/CMakeLists.txt @@ -5,7 +5,7 @@ add_llvm_library(sycl-jit-common Support ) -target_compile_options(sycl-jit-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 diff --git a/sycl-jit/jit-compiler/CMakeLists.txt b/sycl-jit/jit-compiler/CMakeLists.txt index 7fdf73166c40e..f49833b5ac7f8 100644 --- a/sycl-jit/jit-compiler/CMakeLists.txt +++ b/sycl-jit/jit-compiler/CMakeLists.txt @@ -31,7 +31,7 @@ add_llvm_library(sycl-jit ${LLVM_TARGETS_TO_BUILD} ) -target_compile_options(sycl-jit 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 @@ -63,11 +63,11 @@ target_link_libraries(sycl-jit add_dependencies(sycl-jit sycl-headers) if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(sycl-jit 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-jit PRIVATE FUSION_JIT_SUPPORT_AMDGCN) + target_compile_definitions(sycl-jit PRIVATE JIT_SUPPORT_AMDGCN) endif() if(NOT MSVC AND NOT APPLE) diff --git a/sycl-jit/jit-compiler/lib/KernelFusion.cpp b/sycl-jit/jit-compiler/lib/KernelFusion.cpp index 2277c99374919..27886566b5199 100644 --- a/sycl-jit/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-jit/jit-compiler/lib/translation/KernelTranslation.cpp b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp index 0d45a22c9e5a1..f410d41d3f4bf 100644 --- a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp +++ b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp @@ -230,13 +230,13 @@ KernelTranslator::translateToSPIRV(llvm::Module &Mod, JITContext &JITCtx) { llvm::Expected KernelTranslator::translateToPTX( SYCLKernelInfo &KernelInfo, llvm::Module &Mod, JITContext &JITCtx, const std::string &TargetCPU, 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(); @@ -305,19 +305,19 @@ 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, const std::string &TargetCPU, 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(); @@ -381,5 +381,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-jit/passes/CMakeLists.txt b/sycl-jit/passes/CMakeLists.txt index 28c75f856d783..0841465cbf0dd 100644 --- a/sycl-jit/passes/CMakeLists.txt +++ b/sycl-jit/passes/CMakeLists.txt @@ -15,7 +15,7 @@ add_llvm_library(SYCLKernelJIT MODULE intrinsics_gen ) -target_compile_options(SYCLKernelJIT 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 @@ -39,11 +39,11 @@ target_link_libraries(SYCLKernelJIT add_dependencies(SYCLKernelJIT sycl-headers) if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelJIT 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(SYCLKernelJIT PRIVATE FUSION_JIT_SUPPORT_AMDGCN) + target_compile_definitions(SYCLKernelJIT PRIVATE JIT_SUPPORT_AMDGCN) endif() # Static library for linking with the jit_compiler @@ -70,7 +70,7 @@ add_llvm_library(SYCLKernelJITPasses TargetParser ) -target_compile_options(SYCLKernelJITPasses 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 @@ -94,9 +94,9 @@ target_link_libraries(SYCLKernelJITPasses add_dependencies(SYCLKernelJITPasses sycl-headers) if("NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - target_compile_definitions(SYCLKernelJITPasses 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(SYCLKernelJITPasses PRIVATE FUSION_JIT_SUPPORT_AMDGCN) + target_compile_definitions(SYCLKernelJITPasses PRIVATE JIT_SUPPORT_AMDGCN) endif() diff --git a/sycl-jit/passes/target/TargetFusionInfo.cpp b/sycl-jit/passes/target/TargetFusionInfo.cpp index ff1e2e0a352c5..74a96ccf80581 100644 --- a/sycl-jit/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; @@ -1104,18 +1104,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; From 5c8687300aaba14028d2008369a7d18cfb36687f Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Mon, 29 Jul 2024 09:21:19 +0000 Subject: [PATCH 5/9] Mention JIT in getting started (alongside fusion) --- sycl/doc/GetStartedGuide.md | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/sycl/doc/GetStartedGuide.md b/sycl/doc/GetStartedGuide.md index 004ef00c99694..011f41df5026a 100644 --- a/sycl/doc/GetStartedGuide.md +++ b/sycl/doc/GetStartedGuide.md @@ -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-jit` flag. +To disable support for this 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 From e3122f83c6988803c99434513945faf50a8d2c71 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Wed, 31 Jul 2024 09:04:45 +0100 Subject: [PATCH 6/9] Doc build fix --- sycl/doc/GetStartedGuide.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/GetStartedGuide.md b/sycl/doc/GetStartedGuide.md index 011f41df5026a..f9b6e585dbc35 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) From 665b0289c215c6a167dedfe6c05ef74ece8f8613 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Wed, 31 Jul 2024 11:58:48 +0100 Subject: [PATCH 7/9] Use correct ifndef identifier in jit --- sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp index 55747990673c5..2980d19e9f4f7 100644 --- a/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp +++ b/sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp @@ -231,7 +231,7 @@ 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; @@ -313,7 +313,7 @@ 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; From 1608b4b226c8809b7a031d4b7a467deb438a2acc Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Thu, 1 Aug 2024 10:01:20 +0200 Subject: [PATCH 8/9] Update sycl/doc/GetStartedGuide.md Co-authored-by: Alexey Bader --- sycl/doc/GetStartedGuide.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/GetStartedGuide.md b/sycl/doc/GetStartedGuide.md index f9b6e585dbc35..7933d33cda344 100644 --- a/sycl/doc/GetStartedGuide.md +++ b/sycl/doc/GetStartedGuide.md @@ -309,7 +309,7 @@ Support for the experimental SYCL extension for user-driven kernel fusion at runtime is enabled by default. The same mechanism is used to allow JIT compilation of AMD and Nvidia kernels. -To disable support for this features, follow the instructions for the Linux +To disable support for these features, follow the instructions for the Linux DPC++ toolchain, but add the `--disable-jit` flag. Both kernel fusion and JIT compilation of AMD and Nvidia kernels are currently From 3c0d61d300a35b9f27757ca325240904405a5653 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Fri, 2 Aug 2024 07:53:09 +0000 Subject: [PATCH 9/9] Use SYCL_ENABLE_EXTENSION_JIT in fetching UR cmake module --- sycl/cmake/modules/FetchUnifiedRuntime.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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()