diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index ba052e5fe3a87..903e85aa42a59 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -10677,6 +10677,11 @@ static void getNonTripleBasedSYCLPostLinkOpts(const ToolChain &TC, if (TCArgs.hasFlag(options::OPT_fno_sycl_esimd_force_stateless_mem, options::OPT_fsycl_esimd_force_stateless_mem, false)) addArgs(PostLinkArgs, TCArgs, {"-lower-esimd-force-stateless-mem=false"}); + + bool IsUsingLTO = TC.getDriver().isUsingLTO(/*IsDeviceOffloadAction=*/true); + auto LTOMode = TC.getDriver().getLTOMode(/*IsDeviceOffloadAction=*/true); + if (!IsUsingLTO || LTOMode != LTOK_Thin) + addArgs(PostLinkArgs, TCArgs, {"-properties"}); } // Add any sycl-post-link options that rely on a specific Triple in addition @@ -10731,9 +10736,12 @@ static void getTripleBasedSYCLPostLinkOpts(const ToolChain &TC, bool SplitEsimd = TCArgs.hasFlag( options::OPT_fsycl_device_code_split_esimd, options::OPT_fno_sycl_device_code_split_esimd, SplitEsimdByDefault); - // Symbol file and specialization constant info generation is mandatory - + bool IsUsingLTO = TC.getDriver().isUsingLTO(/*IsDeviceOffloadAction=*/true); + auto LTOMode = TC.getDriver().getLTOMode(/*IsDeviceOffloadAction=*/true); + if (!IsUsingLTO || LTOMode != LTOK_Thin) + addArgs(PostLinkArgs, TCArgs, {"-symbols"}); + // Specialization constant info generation is mandatory - // add options unconditionally - addArgs(PostLinkArgs, TCArgs, {"-symbols"}); addArgs(PostLinkArgs, TCArgs, {"-emit-exported-symbols"}); addArgs(PostLinkArgs, TCArgs, {"-emit-imported-symbols"}); if (SplitEsimd) diff --git a/clang/test/Driver/sycl-linker-wrapper-image.cpp b/clang/test/Driver/sycl-linker-wrapper-image.cpp index cfbebd7c01a6a..a8be834a19690 100644 --- a/clang/test/Driver/sycl-linker-wrapper-image.cpp +++ b/clang/test/Driver/sycl-linker-wrapper-image.cpp @@ -5,7 +5,7 @@ // RUN: clang-offload-packager -o %t.fat --image=file=%t.device.bc,kind=sycl,triple=spir64-unknown-unknown // RUN: %clang -cc1 %s -triple=x86_64-unknown-linux-gnu -emit-obj -o %t.o -fembed-offload-object=%t.fat // RUN: clang-linker-wrapper --print-wrapped-module --host-triple=x86_64-unknown-linux-gnu \ -// RUN: -sycl-device-library-location=%S/Inputs -sycl-post-link-options="-split=auto -symbols" \ +// RUN: -sycl-device-library-location=%S/Inputs -sycl-post-link-options="-split=auto -symbols -properties" \ // RUN: %t.o -o %t.out 2>&1 --linker-path="/usr/bin/ld" | FileCheck %s template diff --git a/clang/test/Driver/sycl-lto.cpp b/clang/test/Driver/sycl-lto.cpp index 1c14219daa342..77c75b348f7fd 100644 --- a/clang/test/Driver/sycl-lto.cpp +++ b/clang/test/Driver/sycl-lto.cpp @@ -12,6 +12,7 @@ // RUN: %clangxx -fsycl --offload-new-driver -foffload-lto=thin %s -### 2>&1 | FileCheck -check-prefix=CHECK_SUPPORTED %s // CHECK_SUPPORTED: clang{{.*}} "-cc1" "-triple" "spir64-unknown-unknown" {{.*}} "-flto=thin" "-flto-unit" // CHECK_SUPPORTED: sycl-post-link{{.*}} +// CHECK_SUPPORTED-NOT: -properties // CHECK_SUPPORTED-NEXT: file-table-tform{{.*}} // CHECK_SUPPORTED-NEXT: llvm-foreach{{.*}} "--" {{.*}}clang{{.*}} "-fsycl-is-device"{{.*}} "-flto=thin" "-flto-unit" // CHECK_SUPPORTED-NEXT: file-table-tform{{.*}} diff --git a/clang/test/Driver/sycl-offload-intelfpga-emu.cpp b/clang/test/Driver/sycl-offload-intelfpga-emu.cpp index 2a01cf719801e..cabbdd17b2d0b 100644 --- a/clang/test/Driver/sycl-offload-intelfpga-emu.cpp +++ b/clang/test/Driver/sycl-offload-intelfpga-emu.cpp @@ -178,7 +178,7 @@ // CHK-FPGA-AOCX-SRC: clang-offload-wrapper{{.*}} "-o=[[WRAPOUT:.+\.bc]]" {{.*}} "-target=spir64_fpga" "-kind=sycl" "--sym-prop-bc-files=[[SYM_AND_PROP]]" "-batch" "[[TABLEOUT]]" // CHK-FPGA-AOCX-SRC: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT:.+\.(o|obj)]]" "[[WRAPOUT]]" // CHK-FPGA-AOCX-SRC: llvm-link{{.*}} "[[DEVICEBC]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings" -// CHK-FPGA-AOCX-SRC: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=emulation"{{.*}} "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]] +// CHK-FPGA-AOCX-SRC: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=emulation"{{.*}} "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]] // CHK-FPGA-AOCX-SRC: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]" // CHK-FPGA-AOCX-SRC: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]" // CHK-FPGA-AOCX-SRC: opencl-aot{{.*}} "-device=fpga_fast_emu" "-spv=[[LLVMSPVOUT]]" "-ir=[[OUTPUT4:.+\.aocx]]" "--bo=-g" diff --git a/clang/test/Driver/sycl-offload-intelfpga-link.cpp b/clang/test/Driver/sycl-offload-intelfpga-link.cpp index d816f01b4a6e7..87b4e60de1339 100644 --- a/clang/test/Driver/sycl-offload-intelfpga-link.cpp +++ b/clang/test/Driver/sycl-offload-intelfpga-link.cpp @@ -13,7 +13,7 @@ // CHK-FPGA-LINK-NOT: clang-offload-bundler{{.*}} // CHK-FPGA-LINK: spirv-to-ir-wrapper{{.*}} "[[OUTPUT1]]" "-o" "[[IROUTPUT1:.+\.bc]]" // CHK-FPGA-LINK: llvm-link{{.*}} "[[IROUTPUT1]]" "-o" "[[OUTPUT2_1:.+\.bc]]" -// CHK-FPGA-LINK: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=emulation"{{.*}} "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]" +// CHK-FPGA-LINK: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=emulation"{{.*}} "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]" // CHK-FPGA-LINK: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[OUTPUT2]]" // CHK-FPGA-LINK: llvm-spirv{{.*}} "-o" "[[OUTPUT3:.+\.txt]]" "-spirv-max-version={{.*}}"{{.*}} "[[TABLEOUT]]" // CHK-FPGA-EARLY: aoc{{.*}} "-o" "[[OUTPUT4:.+\.aocr]]" "[[OUTPUT3]]" "-sycl" "-rtl" @@ -46,7 +46,7 @@ // CHK-FPGA-LINK-WIN-NOT: clang-offload-bundler{{.*}} // CHK-FPGA-LINK-WIN: spirv-to-ir-wrapper{{.*}} "[[OUTPUT1]]" "-o" "[[IROUTPUT1:.+\.bc]]" // CHK-FPGA-LINK-WIN: llvm-link{{.*}} "[[IROUTPUT1]]" "-o" "[[OUTPUT2_1:.+\.bc]]" -// CHK-FPGA-LINK-WIN: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=emulation"{{.*}} "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]" +// CHK-FPGA-LINK-WIN: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=emulation"{{.*}} "-o" "[[OUTPUT2:.+\.table]]" "[[OUTPUT2_1]]" // CHK-FPGA-LINK-WIN: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[OUTPUT2]]" // CHK-FPGA-LINK-WIN: llvm-spirv{{.*}} "-o" "[[OUTPUT3:.+\.txt]]" "-spirv-max-version={{.*}}"{{.*}} "[[TABLEOUT]]" // CHK-FPGA-LINK-WIN: aoc{{.*}} "-o" "[[OUTPUT5:.+\.aocr]]" "[[OUTPUT3]]" "-sycl" "-rtl" @@ -175,7 +175,7 @@ // CHK-FPGA: clang-offload-bundler{{.*}} "-type=o" "-targets=host-x86_64-unknown-linux-gnu,sycl-spir64_fpga-unknown-unknown" {{.*}} "-output=[[FINALLINK2x:.+\.o]]" "-output=[[OUTPUT1:.+\.o]]" "-unbundle" // CHK-FPGA: spirv-to-ir-wrapper{{.*}} "[[OUTPUT1]]" "-o" "[[IROUTPUT1:.+\.bc]]" // CHK-FPGA: llvm-link{{.*}} "[[IROUTPUT1]]" "-o" "[[OUTPUT2_BC:.+\.bc]]" -// CHK-FPGA: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=emulation"{{.*}} "-o" "[[OUTPUT3_TABLE:.+\.table]]" "[[OUTPUT2_BC]]" +// CHK-FPGA: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=emulation"{{.*}} "-o" "[[OUTPUT3_TABLE:.+\.table]]" "[[OUTPUT2_BC]]" // CHK-FPGA: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[OUTPUT3_TABLE]]" // CHK-FPGA: llvm-spirv{{.*}} "-o" "[[OUTPUT5:.+\.txt]]" "-spirv-max-version={{.*}}"{{.*}} "[[TABLEOUT]]" // CHK-FPGA: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-fpga_dep" {{.*}} "-output=[[DEPFILE:.+\.d]]" "-unbundle" @@ -237,7 +237,7 @@ // CHK-FPGA-AOCX-SRC: clang-offload-wrapper{{.*}} "-o=[[WRAPOUT:.+\.bc]]" {{.*}} "-target=spir64_fpga" "-kind=sycl" "--sym-prop-bc-files=[[SYM_AND_PROP]]" "-batch" "[[TABLEOUT]]" // CHK-FPGA-AOCX-SRC: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT:.+\.(o|obj)]]" "[[WRAPOUT]]" // CHK-FPGA-AOCX-SRC: llvm-link{{.*}} "[[DEVICEBC]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings" -// CHK-FPGA-AOCX-SRC: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=emulation"{{.*}} "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]] +// CHK-FPGA-AOCX-SRC: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=emulation"{{.*}} "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]] // CHK-FPGA-AOCX-SRC: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]" // CHK-FPGA-AOCX-SRC: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]" // CHK-FPGA-AOCX-SRC: aoc{{.*}} "-o" "[[AOCOUT:.+\.aocx]]" "[[LLVMSPVOUT]]" "-sycl" @@ -263,7 +263,7 @@ // CHK-FPGA-AOCX-OBJ: clang-offload-bundler{{.*}} "-type=o" {{.*}} "-output=[[HOSTOBJx:.+\.(o|obj)]]" "-output=[[DEVICEOBJ:.+\.(o|obj)]]" "-unbundle" // CHK-FPGA-AOCX-OBJ: spirv-to-ir-wrapper{{.*}} "[[DEVICEOBJ]]" "-o" "[[IROUTPUT:.+\.bc]]" // CHK-FPGA-AOCX-OBJ: llvm-link{{.*}} "[[IROUTPUT]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings" -// CHK-FPGA-AOCX-OBJ: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=emulation"{{.*}} "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]] +// CHK-FPGA-AOCX-OBJ: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=emulation"{{.*}} "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]] // CHK-FPGA-AOCX-OBJ: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]" // CHK-FPGA-AOCX-OBJ: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]" // CHK-FPGA-AOCX-OBJ: aoc{{.*}} "-o" "[[AOCOUT:.+\.aocx]]" "[[LLVMSPVOUT]]" "-sycl" @@ -283,7 +283,7 @@ // CHK-FPGA-AOCX-OBJ2: clang-offload-bundler{{.*}} "-type=o" {{.*}} "-output=[[HOSTOBJx:.+\.(o|obj)]]" "-output=[[DEVICEOBJ:.+\.(o|obj)]]" "-output=[[DEVICEOBJ2:.+\.(o|obj)]]" "-unbundle" // CHK-FPGA-AOCX-OBJ2: spirv-to-ir-wrapper{{.*}} "[[DEVICEOBJ]]" "-o" "[[IROUTPUT:.+\.bc]]" // CHK-FPGA-AOCX-OBJ2: llvm-link{{.*}} "[[IROUTPUT]]" "-o" "[[LLVMLINKOUT:.+\.bc]]" "--suppress-warnings" -// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=native"{{.*}} "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]" +// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=native"{{.*}} "-o" "[[POSTLINKOUT:.+\.table]]" "[[LLVMLINKOUT]]" // CHK-FPGA-AOCX-OBJ2: file-table-tform{{.*}} "-o" "[[TABLEOUT:.+\.txt]]" "[[POSTLINKOUT]]" // CHK-FPGA-AOCX-OBJ2: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT:.+\.txt]]" {{.*}} "[[TABLEOUT]]" // CHK-FPGA-AOCX-OBJ2: file-table-tform{{.*}} "-replace=Code,Code" "-o" "[[TFORM_OUT:.+\.table]]" "[[POSTLINKOUT]]" "[[LLVMSPVOUT]]" @@ -291,7 +291,7 @@ // CHK-FPGA-AOCX-OBJ2: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT:.+\.(o|obj)]]" "[[WRAPOUT]]" // CHK-FPGA-AOCX-OBJ2: spirv-to-ir-wrapper{{.*}} "[[DEVICEOBJ2]]" "-o" "[[IROUTPUT2:.+\.bc]]" // CHK-FPGA-AOCX-OBJ2: llvm-link{{.*}} "[[IROUTPUT2]]" "-o" "[[LLVMLINKOUT2:.+\.bc]]" "--suppress-warnings" -// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=emulation"{{.*}} "-o" "[[POSTLINKOUT2:.+\.table]]" "[[LLVMLINKOUT2]]" +// CHK-FPGA-AOCX-OBJ2: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=emulation"{{.*}} "-o" "[[POSTLINKOUT2:.+\.table]]" "[[LLVMLINKOUT2]]" // CHK-FPGA-AOCX-OBJ2: file-table-tform{{.*}} "-o" "[[TABLEOUT2:.+\.txt]]" "[[POSTLINKOUT2]]" // CHK-FPGA-AOCX-OBJ2: llvm-spirv{{.*}} "-o" "[[LLVMSPVOUT2:.+\.txt]]" {{.*}} "[[TABLEOUT2]]" // CHK-FPGA-AOCX-OBJ2: aoc{{.*}} "-o" "[[AOCOUT:.+\.aocx]]" "[[LLVMSPVOUT2]]" "-sycl" diff --git a/clang/test/Driver/sycl-offload-new-driver.c b/clang/test/Driver/sycl-offload-new-driver.c index 405bf64f5ccf6..90c4e7437ca64 100644 --- a/clang/test/Driver/sycl-offload-new-driver.c +++ b/clang/test/Driver/sycl-offload-new-driver.c @@ -60,7 +60,7 @@ // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ // RUN: -Xdevice-post-link -post-link-opt -### %s 2>&1 \ // RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_POSTLINK %s -// WRAPPER_OPTIONS_POSTLINK: clang-linker-wrapper{{.*}} "--sycl-post-link-options=-O2 -device-globals -post-link-opt" +// WRAPPER_OPTIONS_POSTLINK: clang-linker-wrapper{{.*}} "--sycl-post-link-options=-O2 -device-globals -properties -post-link-opt" // -fsycl-device-only behavior // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ diff --git a/clang/test/Driver/sycl-post-link-options-win.cpp b/clang/test/Driver/sycl-post-link-options-win.cpp index 65a802d1f0210..d68d4881c9404 100644 --- a/clang/test/Driver/sycl-post-link-options-win.cpp +++ b/clang/test/Driver/sycl-post-link-options-win.cpp @@ -3,7 +3,7 @@ // RUN: %clangxx -### --target=x86_64-pc-windows-msvc -fsycl \ // RUN: -Xdevice-post-link -O0 %s 2>&1 \ // RUN: | FileCheck -check-prefix OPTIONS_POSTLINK_JIT_OLD %s -// OPTIONS_POSTLINK_JIT_OLD: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=native" "-split=auto" "-emit-only-kernels-as-entry-points" "-emit-param-info" "-symbols" "-emit-exported-symbols" "-emit-imported-symbols" "-split-esimd" "-lower-esimd" "-O0" +// OPTIONS_POSTLINK_JIT_OLD: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=native" "-split=auto" "-emit-only-kernels-as-entry-points" "-emit-param-info" "-symbols" "-emit-exported-symbols" "-emit-imported-symbols" "-split-esimd" "-lower-esimd" "-O0" // RUN: %clang -cc1 %s -triple x86_64-pc-windows-msvc -emit-obj -o %t.elf.o // RUN: clang-offload-packager -o %t.out --image=file=%t.elf.o,kind=sycl,triple=spir64 diff --git a/clang/test/Driver/sycl-post-link-options.cpp b/clang/test/Driver/sycl-post-link-options.cpp index 4f81fb424ec7c..55a6b94feb86e 100644 --- a/clang/test/Driver/sycl-post-link-options.cpp +++ b/clang/test/Driver/sycl-post-link-options.cpp @@ -3,7 +3,7 @@ // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl -### \ // RUN: -Xdevice-post-link -O0 %s 2>&1 \ // RUN: | FileCheck -check-prefix OPTIONS_POSTLINK_JIT_OLD %s -// OPTIONS_POSTLINK_JIT_OLD: sycl-post-link{{.*}} "-O2" "-device-globals" "-spec-const=native" "-split=auto" "-emit-only-kernels-as-entry-points" "-emit-param-info" "-symbols" "-emit-exported-symbols" "-emit-imported-symbols" "-split-esimd" "-lower-esimd" "-O0" +// OPTIONS_POSTLINK_JIT_OLD: sycl-post-link{{.*}} "-O2" "-device-globals" "-properties" "-spec-const=native" "-split=auto" "-emit-only-kernels-as-entry-points" "-emit-param-info" "-symbols" "-emit-exported-symbols" "-emit-imported-symbols" "-split-esimd" "-lower-esimd" "-O0" // RUN: %clang -cc1 %s -triple x86_64-unknown-linux-gnu -emit-obj -o %t.elf.o // RUN: clang-offload-packager -o %t.out --image=file=%t.elf.o,kind=sycl,triple=spir64 @@ -11,6 +11,6 @@ // RUN: -fembed-offload-object=%t.out // RUN: clang-linker-wrapper --dry-run --host-triple=x86_64-unknown-linux-gnu \ // RUN: -sycl-device-library-location=%S/Inputs -sycl-device-libraries=libsycl-crt.new.o \ -// RUN: --sycl-post-link-options="-O2 -device-globals -O0" \ +// RUN: --sycl-post-link-options="-O2 -device-globals -properties -O0" \ // RUN: --linker-path=/usr/bin/ld %t.o -o a.out 2>&1 | FileCheck --check-prefix OPTIONS_POSTLINK_JIT_NEW %s -// OPTIONS_POSTLINK_JIT_NEW: sycl-post-link{{.*}} -spec-const=native -split=auto -emit-only-kernels-as-entry-points -emit-param-info -symbols -emit-exported-symbols -emit-imported-symbols -split-esimd -lower-esimd -O2 -device-globals -O0 +// OPTIONS_POSTLINK_JIT_NEW: sycl-post-link{{.*}} -spec-const=native -split=auto -emit-only-kernels-as-entry-points -emit-param-info -symbols -emit-exported-symbols -emit-imported-symbols -split-esimd -lower-esimd -O2 -device-globals -properties -O0 diff --git a/llvm/include/llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h b/llvm/include/llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h new file mode 100644 index 0000000000000..eaeecb44deb03 --- /dev/null +++ b/llvm/include/llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h @@ -0,0 +1,45 @@ +//===- ComputeModuleRuntimeInfo.h - compute runtime info for module -------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// Functions for computing module properties and symbols for SYCL modules. +//===----------------------------------------------------------------------===// + +#pragma once + +#include "llvm/ADT/SetVector.h" +#include "llvm/SYCLLowerIR/ModuleSplitter.h" +#include "llvm/Support/PropertySetIO.h" +#include +namespace llvm { + +class Function; +class Module; + +namespace sycl { + +struct GlobalBinImageProps { + bool EmitKernelParamInfo; + bool EmitProgramMetadata; + bool EmitExportedSymbols; + bool EmitImportedSymbols; + bool EmitDeviceGlobalPropSet; +}; +bool isModuleUsingAsan(const Module &M); +using PropSetRegTy = llvm::util::PropertySetRegistry; +using EntryPointSet = SetVector; + +PropSetRegTy computeModuleProperties(const Module &M, + const EntryPointSet &EntryPoints, + const GlobalBinImageProps &GlobProps, + bool SpecConstsMet, + bool IsSpecConstantDefault); + +std::string computeModuleSymbolTable(const Module &M, + const EntryPointSet &EntryPoints); + +} // namespace sycl +} // namespace llvm diff --git a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h index 8f485ad0c1667..2465bd212b6e0 100644 --- a/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h +++ b/llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h @@ -207,7 +207,7 @@ class ModuleDesc { const SYCLDeviceRequirements &getOrComputeDeviceRequirements() const { if (!Reqs.has_value()) - Reqs = computeDeviceRequirements(*this); + Reqs = computeDeviceRequirements(getModule(), entries()); return *Reqs; } @@ -306,6 +306,7 @@ struct ModuleSplitterSettings { Expected> splitSYCLModule(std::unique_ptr M, ModuleSplitterSettings Settings); +bool isESIMDFunction(const Function &F); bool canBeImportedFunction(const Function &F); } // namespace module_split diff --git a/llvm/tools/sycl-post-link/SYCLDeviceLibReqMask.h b/llvm/include/llvm/SYCLLowerIR/SYCLDeviceLibReqMask.h similarity index 100% rename from llvm/tools/sycl-post-link/SYCLDeviceLibReqMask.h rename to llvm/include/llvm/SYCLLowerIR/SYCLDeviceLibReqMask.h diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLDeviceRequirements.h b/llvm/include/llvm/SYCLLowerIR/SYCLDeviceRequirements.h index 8891f7f550c5f..8f67b115c43f9 100644 --- a/llvm/include/llvm/SYCLLowerIR/SYCLDeviceRequirements.h +++ b/llvm/include/llvm/SYCLLowerIR/SYCLDeviceRequirements.h @@ -8,6 +8,7 @@ #pragma once +#include "llvm/ADT/SetVector.h" #include "llvm/ADT/SmallString.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringRef.h" @@ -19,12 +20,10 @@ #include namespace llvm { - +class Function; +class Module; class StringRef; -namespace module_split { -class ModuleDesc; -} namespace util { class PropertyValue; } @@ -54,6 +53,7 @@ struct SYCLDeviceRequirements { }; SYCLDeviceRequirements -computeDeviceRequirements(const module_split::ModuleDesc &M); +computeDeviceRequirements(const Module &M, + const SetVector &EntryPoints); } // namespace llvm diff --git a/llvm/tools/sycl-post-link/SYCLKernelParamOptInfo.h b/llvm/include/llvm/SYCLLowerIR/SYCLKernelParamOptInfo.h similarity index 99% rename from llvm/tools/sycl-post-link/SYCLKernelParamOptInfo.h rename to llvm/include/llvm/SYCLLowerIR/SYCLKernelParamOptInfo.h index 00eeadedc1e00..c8331578a00c0 100644 --- a/llvm/tools/sycl-post-link/SYCLKernelParamOptInfo.h +++ b/llvm/include/llvm/SYCLLowerIR/SYCLKernelParamOptInfo.h @@ -12,6 +12,8 @@ // attached to kernel functions in a module. //===----------------------------------------------------------------------===// +#pragma once + #include "llvm/ADT/BitVector.h" #include "llvm/ADT/DenseMap.h" #include "llvm/ADT/StringRef.h" diff --git a/llvm/tools/sycl-post-link/SpecConstants.h b/llvm/include/llvm/SYCLLowerIR/SpecConstants.h similarity index 100% rename from llvm/tools/sycl-post-link/SpecConstants.h rename to llvm/include/llvm/SYCLLowerIR/SpecConstants.h diff --git a/llvm/tools/sycl-post-link/Support.h b/llvm/include/llvm/SYCLLowerIR/Support.h similarity index 100% rename from llvm/tools/sycl-post-link/Support.h rename to llvm/include/llvm/SYCLLowerIR/Support.h diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index 9160029e18c79..6bd5e9eb719b7 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -47,6 +47,7 @@ add_llvm_component_library(LLVMSYCLLowerIR RecordSYCLAspectNames.cpp CleanupSYCLMetadata.cpp CompileTimePropertiesPass.cpp + ComputeModuleRuntimeInfo.cpp DeviceGlobals.cpp ESIMD/LowerESIMDVLoadVStore.cpp ESIMD/LowerESIMDSlmReservation.cpp @@ -56,8 +57,11 @@ add_llvm_component_library(LLVMSYCLLowerIR LowerWGScope.cpp ModuleSplitter.cpp MutatePrintfAddrspace.cpp + SpecConstants.cpp SYCLAddOptLevelAttribute.cpp + SYCLDeviceLibReqMask.cpp SYCLDeviceRequirements.cpp + SYCLKernelParamOptInfo.cpp SYCLPropagateAspectsUsage.cpp SYCLPropagateJointMatrixUsage.cpp SYCLUtils.cpp diff --git a/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp b/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp new file mode 100644 index 0000000000000..643f2605e270f --- /dev/null +++ b/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp @@ -0,0 +1,379 @@ +//===--- ComputeModuleRuntimeInfo.cpp - compute runtime info for module ---===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// See comments in the header. +//===----------------------------------------------------------------------===// +#include "llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h" +#include "llvm/Demangle/Demangle.h" +#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" +#include "llvm/SYCLLowerIR/DeviceGlobals.h" +#include "llvm/SYCLLowerIR/HostPipes.h" +#include "llvm/SYCLLowerIR/ModuleSplitter.h" +#include "llvm/SYCLLowerIR/SYCLDeviceLibReqMask.h" +#include "llvm/SYCLLowerIR/SYCLKernelParamOptInfo.h" +#include "llvm/SYCLLowerIR/SYCLUtils.h" +#include "llvm/SYCLLowerIR/SpecConstants.h" +#include +#include +#ifndef NDEBUG +constexpr int DebugModuleProps = 0; +#endif + +namespace llvm::sycl { +bool isModuleUsingAsan(const Module &M) { + NamedMDNode *MD = M.getNamedMetadata("device.sanitizer"); + if (MD == nullptr) + return false; + assert(MD->getNumOperands() != 0); + auto *MDVal = cast(MD->getOperand(0)->getOperand(0)); + return MDVal->getString() == "asan"; +} + +// This function traverses over reversed call graph by BFS algorithm. +// It means that an edge links some function @func with functions +// which contain call of function @func. It starts from +// @StartingFunction and lifts up until it reach all reachable functions, +// or it reaches some function containing "referenced-indirectly" attribute. +// If it reaches "referenced-indirectly" attribute than it returns an empty +// Optional. +// Otherwise, it returns an Optional containing a list of reached +// SPIR kernel function's names. +std::optional> +traverseCGToFindSPIRKernels(const Function *StartingFunction) { + std::queue FunctionsToVisit; + std::unordered_set VisitedFunctions; + FunctionsToVisit.push(StartingFunction); + std::vector KernelNames; + + while (!FunctionsToVisit.empty()) { + const Function *F = FunctionsToVisit.front(); + FunctionsToVisit.pop(); + + auto InsertionResult = VisitedFunctions.insert(F); + // It is possible that we insert some particular function several + // times in functionsToVisit queue. + if (!InsertionResult.second) + continue; + + for (const auto *U : F->users()) { + const CallInst *CI = dyn_cast(U); + if (!CI) + continue; + + const Function *ParentF = CI->getFunction(); + + if (VisitedFunctions.count(ParentF)) + continue; + + if (ParentF->hasFnAttribute("referenced-indirectly")) + return {}; + + if (ParentF->getCallingConv() == CallingConv::SPIR_KERNEL) + KernelNames.push_back(ParentF->getName()); + + FunctionsToVisit.push(ParentF); + } + } + + return {std::move(KernelNames)}; +} +std::vector getKernelNamesUsingAssert(const Module &M) { + auto *DevicelibAssertFailFunction = M.getFunction("__devicelib_assert_fail"); + if (!DevicelibAssertFailFunction) + return {}; + + auto TraverseResult = + traverseCGToFindSPIRKernels(DevicelibAssertFailFunction); + + if (TraverseResult.has_value()) + return std::move(*TraverseResult); + + // Here we reached "referenced-indirectly", so we need to find all kernels and + // return them. + std::vector SPIRKernelNames; + for (const Function &F : M) { + if (F.getCallingConv() == CallingConv::SPIR_KERNEL) + SPIRKernelNames.push_back(F.getName()); + } + + return SPIRKernelNames; +} + +// Gets reqd_work_group_size information for function Func. +std::vector getKernelReqdWorkGroupSizeMetadata(const Function &Func) { + MDNode *ReqdWorkGroupSizeMD = Func.getMetadata("reqd_work_group_size"); + if (!ReqdWorkGroupSizeMD) + return {}; + size_t NumOperands = ReqdWorkGroupSizeMD->getNumOperands(); + assert(NumOperands >= 1 && NumOperands <= 3 && + "reqd_work_group_size does not have between 1 and 3 operands."); + std::vector OutVals; + OutVals.reserve(NumOperands); + for (const MDOperand &MDOp : ReqdWorkGroupSizeMD->operands()) + OutVals.push_back(mdconst::extract(MDOp)->getZExtValue()); + return OutVals; +} +// Gets work_group_num_dim information for function Func, conviniently 0 if +// metadata is not present. +uint32_t getKernelWorkGroupNumDim(const Function &Func) { + MDNode *MaxDimMD = Func.getMetadata("work_group_num_dim"); + if (!MaxDimMD) + return 0; + assert(MaxDimMD->getNumOperands() == 1 && "Malformed node."); + return mdconst::extract(MaxDimMD->getOperand(0))->getZExtValue(); +} + +PropSetRegTy computeModuleProperties(const Module &M, + const EntryPointSet &EntryPoints, + const GlobalBinImageProps &GlobProps, + bool SpecConstsMet, + bool IsSpecConstantDefault) { + + PropSetRegTy PropSet; + { + uint32_t MRMask = getSYCLDeviceLibReqMask(M); + std::map RMEntry = {{"DeviceLibReqMask", MRMask}}; + PropSet.add(PropSetRegTy::SYCL_DEVICELIB_REQ_MASK, RMEntry); + } + { + PropSet.add(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS, + computeDeviceRequirements(M, EntryPoints).asMap()); + } + if (SpecConstsMet) { + // extract spec constant maps per each module + SpecIDMapTy TmpSpecIDMap; + SpecConstantsPass::collectSpecConstantMetadata(M, TmpSpecIDMap); + PropSet.add(PropSetRegTy::SYCL_SPECIALIZATION_CONSTANTS, TmpSpecIDMap); + + // Add property with the default values of spec constants + std::vector DefaultValues; + SpecConstantsPass::collectSpecConstantDefaultValuesMetadata(M, + DefaultValues); + PropSet.add(PropSetRegTy::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES, "all", + DefaultValues); + } + if (GlobProps.EmitKernelParamInfo) { + // extract kernel parameter optimization info per module + ModuleAnalysisManager MAM; + // Register required analysis + MAM.registerPass([&] { return PassInstrumentationAnalysis(); }); + // Register the payload analysis + + MAM.registerPass([&] { return SYCLKernelParamOptInfoAnalysis(); }); + SYCLKernelParamOptInfo PInfo = + MAM.getResult(const_cast(M)); + + // convert analysis results into properties and record them + llvm::util::PropertySet &Props = + PropSet[PropSetRegTy::SYCL_KERNEL_PARAM_OPT_INFO]; + + for (const auto &NameInfoPair : PInfo) { + const llvm::BitVector &Bits = NameInfoPair.second; + if (Bits.empty()) + continue; // Nothing to add + + const llvm::ArrayRef Arr = Bits.getData(); + const unsigned char *Data = + reinterpret_cast(Arr.begin()); + llvm::util::PropertyValue::SizeTy DataBitSize = Bits.size(); + Props.insert(std::make_pair( + NameInfoPair.first, llvm::util::PropertyValue(Data, DataBitSize))); + } + } + if (GlobProps.EmitExportedSymbols) { + // extract exported functions if any and save them into property set + for (const auto *F : EntryPoints) { + // TODO FIXME some of SYCL/ESIMD functions maybe marked with __regcall CC, + // so they won't make it into the export list. Should the check be + // F->getCallingConv() != CallingConv::SPIR_KERNEL? + if (F->getCallingConv() == CallingConv::SPIR_FUNC) { + PropSet.add(PropSetRegTy::SYCL_EXPORTED_SYMBOLS, F->getName(), + /*PropVal=*/true); + } + } + } + + if (GlobProps.EmitImportedSymbols) { + // record imported functions in the property set + for (const auto &F : M) { + if ( // A function that can be imported may still be defined in one split + // image. Only add import property if this is not the image where the + // function is defined. + F.isDeclaration() && module_split::canBeImportedFunction(F)) { + + // StripDeadPrototypes is called during module splitting + // cleanup. At this point all function decls should have uses. + assert(!F.use_empty() && "Function F has no uses"); + PropSet.add(PropSetRegTy::SYCL_IMPORTED_SYMBOLS, F.getName(), + /*PropVal=*/true); + } + } + } + + // Metadata names may be composite so we keep them alive until the + // properties have been written. + SmallVector MetadataNames; + + if (GlobProps.EmitProgramMetadata) { + // Add reqd_work_group_size and work_group_num_dim information to + // program metadata. + for (const Function &Func : M.functions()) { + std::vector KernelReqdWorkGroupSize = + getKernelReqdWorkGroupSizeMetadata(Func); + if (!KernelReqdWorkGroupSize.empty()) { + MetadataNames.push_back(Func.getName().str() + "@reqd_work_group_size"); + PropSet.add(PropSetRegTy::SYCL_PROGRAM_METADATA, MetadataNames.back(), + KernelReqdWorkGroupSize); + } + + uint32_t WorkGroupNumDim = getKernelWorkGroupNumDim(Func); + if (WorkGroupNumDim) { + MetadataNames.push_back(Func.getName().str() + "@work_group_num_dim"); + PropSet.add(PropSetRegTy::SYCL_PROGRAM_METADATA, MetadataNames.back(), + WorkGroupNumDim); + } + } + + // Add global_id_mapping information with mapping between device-global + // unique identifiers and the variable's name in the IR. + for (auto &GV : M.globals()) { + if (!isDeviceGlobalVariable(GV)) + continue; + + StringRef GlobalID = getGlobalVariableUniqueId(GV); + MetadataNames.push_back(GlobalID.str() + "@global_id_mapping"); + PropSet.add(PropSetRegTy::SYCL_PROGRAM_METADATA, MetadataNames.back(), + GV.getName()); + } + } + bool SeenESIMDFunction = false; + bool SeenSYCLFunction = false; + for (const auto &F : M) { + if (llvm::module_split::isESIMDFunction(F)) + SeenESIMDFunction = true; + else if (utils::isSYCLExternalFunction(&F) && + !F.getName().starts_with("__itt")) + SeenSYCLFunction = true; + } + if (SeenESIMDFunction && !SeenSYCLFunction) + PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "isEsimdImage", true); + { + StringRef RegAllocModeAttr = "sycl-register-alloc-mode"; + uint32_t RegAllocModeVal; + + bool HasRegAllocMode = llvm::any_of(EntryPoints, [&](const Function *F) { + if (!F->hasFnAttribute(RegAllocModeAttr)) + return false; + const auto &Attr = F->getFnAttribute(RegAllocModeAttr); + RegAllocModeVal = getAttributeAsInteger(Attr); + return true; + }); + if (HasRegAllocMode) { + PropSet.add(PropSetRegTy::SYCL_MISC_PROP, RegAllocModeAttr, + RegAllocModeVal); + } + } + + { + StringRef GRFSizeAttr = "sycl-grf-size"; + uint32_t GRFSizeVal; + + bool HasGRFSize = llvm::any_of(EntryPoints, [&](const Function *F) { + if (!F->hasFnAttribute(GRFSizeAttr)) + return false; + const auto &Attr = F->getFnAttribute(GRFSizeAttr); + GRFSizeVal = getAttributeAsInteger(Attr); + return true; + }); + if (HasGRFSize) { + PropSet.add(PropSetRegTy::SYCL_MISC_PROP, GRFSizeAttr, GRFSizeVal); + } + } + + // FIXME: Remove 'if' below when possible + // GPU backend has a problem with accepting optimization level options in form + // described by Level Zero specification (-ze-opt-level=1) when 'invoke_simd' + // functionality is involved. JIT compilation results in the following error: + // error: VLD: Failed to compile SPIR-V with following error: + // invalid api option: -ze-opt-level=O1 + // -11 (PI_ERROR_BUILD_PROGRAM_FAILURE) + // 'if' below essentially preserves the behavior (presumably mistakenly) + // implemented in intel/llvm#8763: ignore 'optLevel' property for images which + // were produced my merge after ESIMD split + if (!SeenESIMDFunction || !SeenSYCLFunction) { + // Handle sycl-optlevel property + int OptLevel = -1; + for (const Function *F : EntryPoints) { + if (!F->hasFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL)) + continue; + + // getAsInteger returns true on error + if (!F->getFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL) + .getValueAsString() + .getAsInteger(10, OptLevel)) { + // It is expected that device-code split has separated kernels with + // different values of sycl-optlevel attribute. Therefore, it is enough + // to only look at the first function with such attribute to compute + // the property for the whole device image. + break; + } + } + + if (OptLevel != -1) + PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "optLevel", OptLevel); + } + { + std::vector FuncNames = getKernelNamesUsingAssert(M); + for (const StringRef &FName : FuncNames) + PropSet.add(PropSetRegTy::SYCL_ASSERT_USED, FName, true); + } + + { + if (isModuleUsingAsan(M)) + PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "asanUsed", true); + } + + if (GlobProps.EmitDeviceGlobalPropSet) { + // Extract device global maps per module + auto DevGlobalPropertyMap = collectDeviceGlobalProperties(M); + if (!DevGlobalPropertyMap.empty()) + PropSet.add(PropSetRegTy::SYCL_DEVICE_GLOBALS, DevGlobalPropertyMap); + } + + auto HostPipePropertyMap = collectHostPipeProperties(M); + if (!HostPipePropertyMap.empty()) { + PropSet.add(PropSetRegTy::SYCL_HOST_PIPES, HostPipePropertyMap); + } + + if (IsSpecConstantDefault) + PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "specConstsReplacedWithDefault", + 1); + + return PropSet; +} +std::string computeModuleSymbolTable(const Module &M, + const EntryPointSet &EntryPoints) { + +#ifndef NDEBUG + if (DebugModuleProps > 0) { + llvm::errs() << "ENTRY POINTS saving Sym table {\n"; + for (const auto *F : EntryPoints) { + llvm::errs() << " " << F->getName() << "\n"; + } + llvm::errs() << "}\n"; + } +#endif // NDEBUG + // Concatenate names of the input entry points with "\n". + std::string SymT; + + for (const auto *F : EntryPoints) { + SymT = (Twine(SymT) + Twine(F->getName()) + Twine("\n")).str(); + } + return SymT; +} + +} // namespace llvm::sycl diff --git a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp index ddac91c28875c..0c024610c8235 100644 --- a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp +++ b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp @@ -135,10 +135,6 @@ bool isEntryPoint(const Function &F, bool EmitOnlyKernelsAsEntryPoints) { return false; } -bool isESIMDFunction(const Function &F) { - return F.getMetadata(ESIMD_MARKER_MD) != nullptr; -} - // Represents "dependency" or "use" graph of global objects (functions and // global variables) in a module. It is used during device code split to // understand which global variables and functions (other than entry points) @@ -444,6 +440,10 @@ class ModuleSplitter : public ModuleSplitterBase { namespace llvm { namespace module_split { +bool isESIMDFunction(const Function &F) { + return F.getMetadata(ESIMD_MARKER_MD) != nullptr; +} + cl::OptionCategory &getModuleSplitCategory() { static cl::OptionCategory ModuleSplitCategory{"Module Split options"}; return ModuleSplitCategory; diff --git a/llvm/tools/sycl-post-link/SYCLDeviceLibReqMask.cpp b/llvm/lib/SYCLLowerIR/SYCLDeviceLibReqMask.cpp similarity index 99% rename from llvm/tools/sycl-post-link/SYCLDeviceLibReqMask.cpp rename to llvm/lib/SYCLLowerIR/SYCLDeviceLibReqMask.cpp index 1554e81751668..5f270baecec1d 100644 --- a/llvm/tools/sycl-post-link/SYCLDeviceLibReqMask.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLDeviceLibReqMask.cpp @@ -14,9 +14,9 @@ // SYCL runtime later. //===----------------------------------------------------------------------===// -#include "SYCLDeviceLibReqMask.h" -#include "llvm/TargetParser/Triple.h" +#include "llvm/SYCLLowerIR/SYCLDeviceLibReqMask.h" #include "llvm/IR/Module.h" +#include "llvm/TargetParser/Triple.h" #include #include diff --git a/llvm/lib/SYCLLowerIR/SYCLDeviceRequirements.cpp b/llvm/lib/SYCLLowerIR/SYCLDeviceRequirements.cpp index 60424c04027fa..cbc6ec2f847d9 100644 --- a/llvm/lib/SYCLLowerIR/SYCLDeviceRequirements.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLDeviceRequirements.cpp @@ -37,11 +37,12 @@ static llvm::StringRef ExtractStringFromMDNodeOperand(const MDNode *N, } SYCLDeviceRequirements -llvm::computeDeviceRequirements(const module_split::ModuleDesc &MD) { +llvm::computeDeviceRequirements(const Module &M, + const SetVector &EntryPoints) { SYCLDeviceRequirements Reqs; bool MultipleReqdWGSize = false; // Process all functions in the module - for (const Function &F : MD.getModule()) { + for (const Function &F : M) { if (auto *MDN = F.getMetadata("sycl_used_aspects")) { for (size_t I = 0, E = MDN->getNumOperands(); I < E; ++I) { StringRef AspectName = ""; @@ -98,7 +99,7 @@ llvm::computeDeviceRequirements(const module_split::ModuleDesc &MD) { } // Process just the entry points in the module - for (const Function *F : MD.entries()) { + for (const Function *F : EntryPoints) { if (auto *MDN = F->getMetadata("intel_reqd_sub_group_size")) { // There should only be at most one function with // intel_reqd_sub_group_size metadata when considering the entry diff --git a/llvm/tools/sycl-post-link/SYCLKernelParamOptInfo.cpp b/llvm/lib/SYCLLowerIR/SYCLKernelParamOptInfo.cpp similarity index 96% rename from llvm/tools/sycl-post-link/SYCLKernelParamOptInfo.cpp rename to llvm/lib/SYCLLowerIR/SYCLKernelParamOptInfo.cpp index 5a1c6f4e3e03b..427c4b6c47b5d 100644 --- a/llvm/tools/sycl-post-link/SYCLKernelParamOptInfo.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLKernelParamOptInfo.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#include "SYCLKernelParamOptInfo.h" +#include "llvm/SYCLLowerIR/SYCLKernelParamOptInfo.h" #include "llvm/IR/Constants.h" #include "llvm/Support/Casting.h" diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/lib/SYCLLowerIR/SpecConstants.cpp similarity index 99% rename from llvm/tools/sycl-post-link/SpecConstants.cpp rename to llvm/lib/SYCLLowerIR/SpecConstants.cpp index 3ef8d15d338bd..58f5a0d54b26e 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/lib/SYCLLowerIR/SpecConstants.cpp @@ -8,8 +8,8 @@ // See comments in the header. //===----------------------------------------------------------------------===// -#include "SpecConstants.h" -#include "Support.h" +#include "llvm/SYCLLowerIR/SpecConstants.h" +#include "llvm/SYCLLowerIR/Support.h" #include "llvm/ADT/APInt.h" #include "llvm/ADT/StringMap.h" diff --git a/llvm/test/tools/sycl-post-link/assert/indirect-with-split-2.ll b/llvm/test/tools/sycl-post-link/assert/indirect-with-split-2.ll index 65e1420d81356..e7011ae6de141 100644 --- a/llvm/test/tools/sycl-post-link/assert/indirect-with-split-2.ll +++ b/llvm/test/tools/sycl-post-link/assert/indirect-with-split-2.ll @@ -8,7 +8,7 @@ ; __devicelib_assert_fail, then all kernels in the module are conservatively ; marked as using asserts. -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop -check-prefixes=CHECK,CHECK0 \ ; RUN: --implicit-check-not TU1 ; RUN: FileCheck %s -input-file=%t_1.prop -check-prefixes=CHECK,CHECK1 \ diff --git a/llvm/test/tools/sycl-post-link/assert/indirect-with-split.ll b/llvm/test/tools/sycl-post-link/assert/indirect-with-split.ll index e7ba6c43bb240..639dd73359c1b 100644 --- a/llvm/test/tools/sycl-post-link/assert/indirect-with-split.ll +++ b/llvm/test/tools/sycl-post-link/assert/indirect-with-split.ll @@ -6,7 +6,7 @@ ; __devicelib_assert_fail, then all kernels in the module are conservatively ; marked as using asserts. -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes=CHECK,CHECK1 \ ; RUN: --implicit-check-not TU0 ; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes=CHECK,CHECK0 \ diff --git a/llvm/test/tools/sycl-post-link/assert/property-1.ll b/llvm/test/tools/sycl-post-link/assert/property-1.ll index 81e7c674187be..2df5e11dbbb07 100644 --- a/llvm/test/tools/sycl-post-link/assert/property-1.ll +++ b/llvm/test/tools/sycl-post-link/assert/property-1.ll @@ -2,16 +2,16 @@ ; property - it should include only kernels that call assertions in their call ; graph. -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop --implicit-check-not TheKernel2 ; -; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop --implicit-check-not TheKernel2 ; -; RUN: sycl-post-link -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop --implicit-check-not TheKernel2 ; -; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes=CHECK-K3 ; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes=CHECK-K1 ; RUN: FileCheck %s -input-file=%t_2.prop --check-prefixes=CHECK-K2 diff --git a/llvm/test/tools/sycl-post-link/assert/property-2.ll b/llvm/test/tools/sycl-post-link/assert/property-2.ll index 6f4bea447ab89..4fafcde79e829 100644 --- a/llvm/test/tools/sycl-post-link/assert/property-2.ll +++ b/llvm/test/tools/sycl-post-link/assert/property-2.ll @@ -2,7 +2,7 @@ ; property - it should include only kernels that call assertions in their call ; graph. -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop -check-prefix=PRESENCE-CHECK ; RUN: FileCheck %s -input-file=%t_0.prop -check-prefix=ABSENCE-CHECK diff --git a/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-1.ll b/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-1.ll index 0583cfde3af23..bbcde38cbadcb 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-1.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-1.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; By default auto mode is equal to source mode ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-TU0,CHECK ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-TU1,CHECK diff --git a/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-2.ll b/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-2.ll index 4ff2095f42bbb..0f7cf05ae86d5 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-2.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-2.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; ; This is the same as auto-module-split-1 test with the only difference is that ; @_Z3foov is marked with "referenced-indirectly" attribute. diff --git a/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-3.ll b/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-3.ll index a5c62a5912338..4dfbc4bfbd163 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-3.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-3.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; ; In precense of indirect calls we start matching functions using their ; signatures, i.e. we have an indirect call to i32(i32) function within diff --git a/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-func-ptr.ll b/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-func-ptr.ll index 730d9a5cd8efc..d36de8d1d3ff0 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-func-ptr.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/auto-module-split-func-ptr.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix=CHECK-SYM0 ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix=CHECK-SYM1 ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix=CHECK-IR0 diff --git a/llvm/test/tools/sycl-post-link/device-code-split/basic-module-split.ll b/llvm/test/tools/sycl-post-link/device-code-split/basic-module-split.ll index 48d58248d0095..f2898cbea387c 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/basic-module-split.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/basic-module-split.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-TU0,CHECK ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-TU1,CHECK ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-TU0-TXT diff --git a/llvm/test/tools/sycl-post-link/device-code-split/complex-indirect-call-chain.ll b/llvm/test/tools/sycl-post-link/device-code-split/complex-indirect-call-chain.ll index 064471405a58d..d2e2084ca2b34 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/complex-indirect-call-chain.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/complex-indirect-call-chain.ll @@ -1,7 +1,7 @@ ; The idea of the test is to ensure that sycl-post-link can trace through more ; complex call stacks involving several nested indirect calls -; RUN: sycl-post-link -split=auto -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK0 \ ; RUN: --implicit-check-not @foo --implicit-check-not @kernel_A \ ; RUN: --implicit-check-not @kernel_B --implicit-check-not @baz @@ -23,7 +23,7 @@ ; RUN: --implicit-check-not @BAZ --implicit-check-not @kernel_B \ ; RUN: --implicit-check-not @kernel_C ; -; RUN: sycl-post-link -split=source -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK0 \ ; RUN: --implicit-check-not @foo --implicit-check-not @kernel_A \ ; RUN: --implicit-check-not @kernel_B --implicit-check-not @baz @@ -45,7 +45,7 @@ ; RUN: --implicit-check-not @BAZ --implicit-check-not @kernel_B \ ; RUN: --implicit-check-not @kernel_C ; -; RUN: sycl-post-link -split=kernel -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK0 \ ; RUN: --implicit-check-not @foo --implicit-check-not @kernel_A \ ; RUN: --implicit-check-not @kernel_B --implicit-check-not @baz diff --git a/llvm/test/tools/sycl-post-link/device-code-split/one-kernel-per-module.ll b/llvm/test/tools/sycl-post-link/device-code-split/one-kernel-per-module.ll index 0197a2edd4a1b..26081964dd2e7 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/one-kernel-per-module.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/one-kernel-per-module.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefixes CHECK-MODULE0,CHECK ; RUN: FileCheck %s -input-file=%t.files_0.sym --check-prefixes CHECK-MODULE0-TXT ; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefixes CHECK-MODULE1,CHECK diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll index 51a2895f4d326..8c64029f5e1c2 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll @@ -7,7 +7,7 @@ ; that use aspects from kernels which doesn't use aspects regardless of device ; code split mode -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ @@ -35,7 +35,7 @@ ; RUN: FileCheck %s -input-file=%t2_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ @@ -63,7 +63,7 @@ ; RUN: FileCheck %s -input-file=%t2_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-2.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-2.ll index f4d66822b261c..4962b19721790 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-2.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-2.ll @@ -1,7 +1,7 @@ ; The test is intended to check that sycl-post-link correctly groups kernels ; by unique sets of aspects used in them -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M0-SYMS \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-3.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-3.ll index 523477a07573b..93dcf91b578ac 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-3.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-3.ll @@ -1,7 +1,7 @@ ; This test is intended to check that per-aspect device code split works as ; expected with SYCL_EXTERNAL functions -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M0-SYMS \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-4.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-4.ll index 9655a72e140e6..e2d32bce59c5d 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-4.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-4.ll @@ -1,7 +1,7 @@ ; This test is intended to check that we do not perform per-aspect split if ; it was disabled through one or another sycl-post-link option -; RUN: sycl-post-link -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK-IR ; diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-1.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-1.ll index 543a892415fa4..b7d0bcfceeccd 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-1.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-1.ll @@ -5,7 +5,7 @@ ; The test is intended to check that sycl-post-link correctly separates kernels ; that use different sycl_joint_matrix metadata -; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K2 \ @@ -25,7 +25,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K2 \ ; RUN: --implicit-check-not Kernel3 -; RUN: sycl-post-link -split=source -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K2 \ @@ -45,7 +45,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K2 \ ; RUN: --implicit-check-not Kernel3 -; RUN: sycl-post-link -split=kernel -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K2 \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-2.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-2.ll index 6c054fc579659..81fbd2a05df50 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-2.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-2.ll @@ -6,7 +6,7 @@ ; The test is intended to check that sycl-post-link correctly separates kernels ; that use different sycl_joint_matrix metadata and kernels without that metadata -; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K2 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel3 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K3 \ @@ -26,7 +26,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K3 \ ; RUN: --implicit-check-not Kernel2 -; RUN: sycl-post-link -split=source -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K2 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel3 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K3 \ @@ -46,7 +46,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K3 \ ; RUN: --implicit-check-not Kernel2 -; RUN: sycl-post-link -split=kernel -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K2 \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-3.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-3.ll index 9a365d17faedd..64180694fae72 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-3.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-3.ll @@ -1,7 +1,7 @@ ; This test is intended to check that we do not perform per-joint-matrix ; split if it was disabled through one or another sycl-post-link option -; RUN: sycl-post-link -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK-IR ; diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-1.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-1.ll index fd64b234b2c6f..132568926dbce 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-1.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-1.ll @@ -5,7 +5,7 @@ ; The test is intended to check that sycl-post-link correctly separates kernels ; that use different sycl_joint_matrix_mad metadata -; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K2 \ @@ -25,7 +25,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K2 \ ; RUN: --implicit-check-not Kernel3 -; RUN: sycl-post-link -split=source -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K2 \ @@ -45,7 +45,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K2 \ ; RUN: --implicit-check-not Kernel3 -; RUN: sycl-post-link -split=kernel -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K2 \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-2.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-2.ll index 4c4a4bc8a1a6e..ee8e9d1ead30a 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-2.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-2.ll @@ -6,7 +6,7 @@ ; The test is intended to check that sycl-post-link correctly separates kernels ; that use different sycl_joint_matrix_mad metadata and kernels without that metadata -; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K2 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel3 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K3 \ @@ -26,7 +26,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K3 \ ; RUN: --implicit-check-not Kernel2 -; RUN: sycl-post-link -split=source -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K2 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel3 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K3 \ @@ -46,7 +46,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K3 \ ; RUN: --implicit-check-not Kernel2 -; RUN: sycl-post-link -split=kernel -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K2 \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-3.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-3.ll index 251c10969c047..e3abc63e8c5ec 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-3.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-3.ll @@ -1,7 +1,7 @@ ; This test is intended to check that we do not perform per-joint-matrix-mad ; split if it was disabled through one or another sycl-post-link option -; RUN: sycl-post-link -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK-IR ; diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-4.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-4.ll index fe995542deba1..b94c42e1f4d4f 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-4.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-4.ll @@ -6,7 +6,7 @@ ; The test is intended to check that sycl-post-link correctly separates kernels ; that use different sycl_joint_matrix_mad metadata -; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K2 \ @@ -26,7 +26,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K2 \ ; RUN: --implicit-check-not Kernel3 -; RUN: sycl-post-link -split=source -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K1,CHECK-IR-K2 \ @@ -46,7 +46,7 @@ ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefixes CHECK-SYMS-K1,CHECK-SYMS-K2 \ ; RUN: --implicit-check-not Kernel3 -; RUN: sycl-post-link -split=kernel -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-K2 \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-5.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-5.ll index 25fd2e26f3ca4..0d9631af085d5 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-5.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-joint-matrix-mad-5.ll @@ -7,7 +7,7 @@ ; The test is intended to check that sycl-post-link correctly separates kernels ; that use different sycl_joint_matrix_mad metadata -; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3,CHECK-IR-K5 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 \ ; RUN: --implicit-check-not Kernel4 --implicit-check-not Kernel6 @@ -59,7 +59,7 @@ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 \ ; RUN: --implicit-check-not Kernel3 --implicit-check-not Kernel5 --implicit-check-not Kernel6 -; RUN: sycl-post-link -split=source -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K3,CHECK-IR-K5 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 \ ; RUN: --implicit-check-not Kernel4 --implicit-check-not Kernel6 @@ -111,7 +111,7 @@ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 \ ; RUN: --implicit-check-not Kernel3 --implicit-check-not Kernel5 --implicit-check-not Kernel6 -; RUN: sycl-post-link -split=kernel -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-K6 \ ; RUN: --implicit-check-not Kernel1 --implicit-check-not Kernel2 \ ; RUN: --implicit-check-not Kernel3 --implicit-check-not Kernel4 --implicit-check-not Kernel5 diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-1.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-1.ll index 393943b63db43..ca723e075d2c0 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-1.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-1.ll @@ -7,7 +7,7 @@ ; that use reqd_sub_group_size attributes from kernels which doesn't use them ; regardless of device code split mode -; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ @@ -35,7 +35,7 @@ ; RUN: FileCheck %s -input-file=%t2_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 -; RUN: sycl-post-link -split=kernel -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ @@ -63,7 +63,7 @@ ; RUN: FileCheck %s -input-file=%t2_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 -; RUN: sycl-post-link -split=source -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-2.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-2.ll index 1efeb364cb2e3..a943450a2a459 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-2.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-2.ll @@ -1,7 +1,7 @@ ; The test is intended to check that sycl-post-link correctly groups kernels ; by unique reqd_sub_group_size values used in them -; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M0-SYMS \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-1.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-1.ll index b156d71b1e3f6..4b105faa24073 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-1.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-1.ll @@ -7,7 +7,7 @@ ; that use reqd_work_group_size attributes from kernels which doesn't use them ; regardless of device code split mode -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ @@ -35,7 +35,7 @@ ; RUN: FileCheck %s -input-file=%t2_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 -; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ @@ -63,7 +63,7 @@ ; RUN: FileCheck %s -input-file=%t2_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 -; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll index c92ae8dbc9c03..d453ba763ec3c 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll @@ -1,7 +1,7 @@ ; The test is intended to check that sycl-post-link correctly groups kernels ; by unique reqd_work_group_size values used in them -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M0-SYMS \ diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-3.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-3.ll index ae6d79a041175..905eb6801f041 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-3.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-3.ll @@ -1,7 +1,7 @@ ; This test is intended to check that we do not perform per-reqd_work_group_size ; split if it was disabled through one or another sycl-post-link option -; RUN: sycl-post-link -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK-IR ; diff --git a/llvm/test/tools/sycl-post-link/device-code-split/split-with-func-ptrs.ll b/llvm/test/tools/sycl-post-link/device-code-split/split-with-func-ptrs.ll index af920f0ab373d..148f8e8545f75 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/split-with-func-ptrs.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/split-with-func-ptrs.ll @@ -3,14 +3,14 @@ ; modules. ; -- Per-source split -; RUN: sycl-post-link -split=source -emit-param-info -symbols -emit-exported-symbols -split-esimd -lower-esimd -O2 -spec-const=native -S < %s -o %tA.table +; RUN: sycl-post-link -properties -split=source -emit-param-info -symbols -emit-exported-symbols -split-esimd -lower-esimd -O2 -spec-const=native -S < %s -o %tA.table ; RUN: FileCheck %s -input-file=%tA_0.ll --check-prefixes CHECK-A0 ; RUN: FileCheck %s -input-file=%tA_1.ll --check-prefixes CHECK-A1 ; -- No split -; RUN: sycl-post-link -emit-param-info -symbols -emit-exported-symbols -split-esimd -lower-esimd -O2 -spec-const=native -S < %s -o %tB.table +; RUN: sycl-post-link -properties -emit-param-info -symbols -emit-exported-symbols -split-esimd -lower-esimd -O2 -spec-const=native -S < %s -o %tB.table ; RUN: FileCheck %s -input-file=%tB_0.ll --check-prefixes CHECK-B0 ; -- Per-kernel split -; RUN: sycl-post-link -split=kernel -emit-param-info -symbols -emit-exported-symbols -split-esimd -lower-esimd -O2 -spec-const=native -S < %s -o %tC.table +; RUN: sycl-post-link -properties -split=kernel -emit-param-info -symbols -emit-exported-symbols -split-esimd -lower-esimd -O2 -spec-const=native -S < %s -o %tC.table ; RUN: FileCheck %s -input-file=%tC_0.ll --check-prefixes CHECK-C0 ; RUN: FileCheck %s -input-file=%tC_1.ll --check-prefixes CHECK-C1 diff --git a/llvm/test/tools/sycl-post-link/device-code-split/split-with-kernel-declarations.ll b/llvm/test/tools/sycl-post-link/device-code-split/split-with-kernel-declarations.ll index 82213e4b3beeb..42dd9e980fe82 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/split-with-kernel-declarations.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/split-with-kernel-declarations.ll @@ -1,7 +1,7 @@ ; Purpose of this test is to check that sycl-post-link does not treat ; declarations as entry points. -; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-PER-SOURCE-TABLE ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-PER-SOURCE-SYM0 ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-PER-SOURCE-SYM1 @@ -11,7 +11,7 @@ ; RUN: FileCheck %s -input-file=%t1_0.sym --check-prefix CHECK-PER-SOURCE-SYM0 ; RUN: FileCheck %s -input-file=%t1_1.sym --check-prefix CHECK-PER-SOURCE-SYM1 ; -; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t2.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t2.table ; RUN: FileCheck %s -input-file=%t2.table --check-prefix CHECK-PER-KERNEL-TABLE ; RUN: FileCheck %s -input-file=%t2_0.sym --check-prefix CHECK-PER-KERNEL-SYM1 ; RUN: FileCheck %s -input-file=%t2_1.sym --check-prefix CHECK-PER-KERNEL-SYM2 diff --git a/llvm/test/tools/sycl-post-link/device-code-split/vtable.ll b/llvm/test/tools/sycl-post-link/device-code-split/vtable.ll index cb9fd1f77cf78..6d932aba577f9 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/vtable.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/vtable.ll @@ -39,7 +39,7 @@ ; variables storing vtable, are also included into the final module, even though ; they are not directly used in a kernel otherwise. ; -; RUN: sycl-post-link -split=auto -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll ; ; RUN: sycl-module-split -split=auto -S < %s -o %t2 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable.ll index a97d69ffbc7fe..cd07c948fb91d 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP source_filename = "test_global_variable.cpp" target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll index 986b8b1917cc7..2e2c6127083e1 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR ; ; Test checks that llvm.compiler.used is removed when all values in it are diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll index 4c786e0df9f11..80905a39b9518 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR ; ; Test checks that llvm.compiler.used is removed when all values in it are diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll index 03ccc2f5f696b..6ce11b4dcdf52 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals --split=source -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals --split=source -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD1 ; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD0 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll index 3b94a48cb3e90..406bf7edbf2b3 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals --split=source -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals --split=source -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD2 ; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD0 ; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD1 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll index 847ecbc2b102d..9934318e77793 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals --split=source -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals --split=source -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD2 ; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD0 ; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD1 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll index 9afd208726d79..e5d1b352ffca8 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals --split=source -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals --split=source -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD0 ; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD1 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_name_mapping_metadata.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_name_mapping_metadata.ll index a7546dafef2c8..f3d8c6e0bff07 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_name_mapping_metadata.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_name_mapping_metadata.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals --emit-program-metadata -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals --emit-program-metadata -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP ; This test is intended to check that the global_id_mapping program metadata properties are diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll index 6f17f406e5e01..b4d3bf3f557e1 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR ; ; Test checks that all device_global variables in llvm.compiler.used are removed diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll index 98668ef83369e..771f8ab1bada8 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR ; ; Test checks that all device_global variables in llvm.compiler.used are removed diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_no_property_set_header_for_an_empty_set.ll b/llvm/test/tools/sycl-post-link/device-globals/test_no_property_set_header_for_an_empty_set.ll index 06128c4210544..3d00379d2259b 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_no_property_set_header_for_an_empty_set.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_no_property_set_header_for_an_empty_set.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --device-globals -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --device-globals -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP ; This test is intended to check that sycl-post-link doesn't add the header for diff --git a/llvm/test/tools/sycl-post-link/device-requirements/aspects.ll b/llvm/test/tools/sycl-post-link/device-requirements/aspects.ll index 98e92650e516f..eb9f253df0b7a 100644 --- a/llvm/test/tools/sycl-post-link/device-requirements/aspects.ll +++ b/llvm/test/tools/sycl-post-link/device-requirements/aspects.ll @@ -13,10 +13,10 @@ ; }); ; } -; RUN: sycl-post-link -split=auto < %s -o %t.files.table +; RUN: sycl-post-link -properties -split=auto < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP-AUTO-SPLIT -; RUN: sycl-post-link -split=kernel < %s -o %t.files.table +; RUN: sycl-post-link -properties -split=kernel < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-1 ; RUN: FileCheck %s -input-file=%t.files_1.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-0 diff --git a/llvm/test/tools/sycl-post-link/device-requirements/fixed-target.ll b/llvm/test/tools/sycl-post-link/device-requirements/fixed-target.ll index a1b09c7b64a62..c7b7395ad0bd7 100644 --- a/llvm/test/tools/sycl-post-link/device-requirements/fixed-target.ll +++ b/llvm/test/tools/sycl-post-link/device-requirements/fixed-target.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=auto < %s -o %t.files.table +; RUN: sycl-post-link -properties -split=auto < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop ; CHECK: [SYCL/device requirements] diff --git a/llvm/test/tools/sycl-post-link/device-requirements/joint-matrix.ll b/llvm/test/tools/sycl-post-link/device-requirements/joint-matrix.ll index a710c0a9233e0..56e47665638f4 100644 --- a/llvm/test/tools/sycl-post-link/device-requirements/joint-matrix.ll +++ b/llvm/test/tools/sycl-post-link/device-requirements/joint-matrix.ll @@ -21,7 +21,7 @@ ; return 0; ; } -; RUN: sycl-post-link -split=kernel %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-0 ; RUN: FileCheck %s -input-file=%t_1.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-1 diff --git a/llvm/test/tools/sycl-post-link/device-requirements/reqd-sub-group-size.ll b/llvm/test/tools/sycl-post-link/device-requirements/reqd-sub-group-size.ll index df4d5682c0623..1f3a51849d1ab 100644 --- a/llvm/test/tools/sycl-post-link/device-requirements/reqd-sub-group-size.ll +++ b/llvm/test/tools/sycl-post-link/device-requirements/reqd-sub-group-size.ll @@ -22,7 +22,7 @@ ; return 0; ; } -; RUN: sycl-post-link -split=auto %s -o %t.table +; RUN: sycl-post-link -properties -split=auto %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefix CHECK-PROP-AUTO-SPLIT-0 ; RUN: FileCheck %s -input-file=%t_1.prop --check-prefix CHECK-PROP-AUTO-SPLIT-1 diff --git a/llvm/test/tools/sycl-post-link/device-requirements/reqd-work-group-size.ll b/llvm/test/tools/sycl-post-link/device-requirements/reqd-work-group-size.ll index 4a440332a3634..da6ffc46e49b8 100644 --- a/llvm/test/tools/sycl-post-link/device-requirements/reqd-work-group-size.ll +++ b/llvm/test/tools/sycl-post-link/device-requirements/reqd-work-group-size.ll @@ -21,7 +21,7 @@ ; return 0; ; } -; RUN: sycl-post-link -split=auto < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefix CHECK-PROP-AUTO-SPLIT-0 ; RUN: FileCheck %s -input-file=%t_1.prop --check-prefix CHECK-PROP-AUTO-SPLIT-1 diff --git a/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll b/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll index 002b14076dec2..c829f2dca9120 100644 --- a/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll +++ b/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll @@ -1,7 +1,7 @@ ; This test checks that the post-link tool properly generates "asanUsed=1" ; in [SYCL/misc properties] -; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop ; CHECK: [SYCL/misc properties] ; CHECK: asanUsed=1 diff --git a/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll b/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll index 49a94eefa2575..c6705ea972288 100644 --- a/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll +++ b/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll @@ -1,17 +1,17 @@ ; This test checks that the post-link tool generates list of exported symbols. ; ; Global scope -; RUN: sycl-post-link -symbols -emit-exported-symbols -S < %s -o %t.global.files.table +; RUN: sycl-post-link -properties -symbols -emit-exported-symbols -S < %s -o %t.global.files.table ; RUN: FileCheck %s -input-file=%t.global.files_0.prop --implicit-check-not="NotExported" --check-prefix=CHECK-GLOBAL-PROP ; ; Per-module split -; RUN: sycl-post-link -symbols -split=source -emit-exported-symbols -S < %s -o %t.per_module.files.table +; RUN: sycl-post-link -properties -symbols -split=source -emit-exported-symbols -S < %s -o %t.per_module.files.table ; RUN: FileCheck %s -input-file=%t.per_module.files_0.prop -implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP ; RUN: FileCheck %s -input-file=%t.per_module.files_1.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-0-PROP ; RUN: FileCheck %s -input-file=%t.per_module.files_2.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-2-PROP ; ; Per-kernel split -; RUN: sycl-post-link -symbols -split=kernel -emit-exported-symbols -S < %s -o %t.per_kernel.files.table +; RUN: sycl-post-link -properties -symbols -split=kernel -emit-exported-symbols -S < %s -o %t.per_kernel.files.table ; RUN: FileCheck %s -input-file=%t.per_kernel.files_0.prop --implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP ; RUN: FileCheck %s -input-file=%t.per_kernel.files_1.prop --implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP ; RUN: FileCheck %s -input-file=%t.per_kernel.files_2.prop --implicit-check-not="NotExported" --check-prefix=CHECK-PERKERNEL-0-PROP diff --git a/llvm/test/tools/sycl-post-link/emit_imported_symbols.ll b/llvm/test/tools/sycl-post-link/emit_imported_symbols.ll index ae824d293b9ea..bade08d34147e 100644 --- a/llvm/test/tools/sycl-post-link/emit_imported_symbols.ll +++ b/llvm/test/tools/sycl-post-link/emit_imported_symbols.ll @@ -6,7 +6,7 @@ ; Test with -split=kernel ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;; -; RUN: sycl-post-link -symbols -emit-imported-symbols -split=kernel -S < %s -o %t_kernel.table +; RUN: sycl-post-link -properties -symbols -emit-imported-symbols -split=kernel -S < %s -o %t_kernel.table ; RUN: FileCheck %s -input-file=%t_kernel_0.sym --check-prefixes CHECK-KERNEL-SYM-0 ; RUN: FileCheck %s -input-file=%t_kernel_1.sym --check-prefixes CHECK-KERNEL-SYM-1 @@ -41,11 +41,11 @@ ; Test with -split=source ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;; -; RUN: sycl-post-link -symbols -emit-imported-symbols -split=source -S < %s -o %t_source.table +; RUN: sycl-post-link -properties -symbols -emit-imported-symbols -split=source -S < %s -o %t_source.table ; RUN: FileCheck %s -input-file=%t_source_0.sym --check-prefixes CHECK-SOURCE-SYM-0 ; RUN: FileCheck %s -input-file=%t_source_0.prop --check-prefixes CHECK-SOURCE-IMPORTED-SYM-0 -; RUN: sycl-post-link -symbols -emit-imported-symbols -split=source -S < %s -o %t_source.table -O0 +; RUN: sycl-post-link -properties -symbols -emit-imported-symbols -split=source -S < %s -o %t_source.table -O0 ; RUN: FileCheck %s -input-file=%t_source_0.sym --check-prefixes CHECK-SOURCE-SYM-0 ; RUN: FileCheck %s -input-file=%t_source_0.prop --check-prefixes CHECK-SOURCE-IMPORTED-SYM-0 diff --git a/llvm/test/tools/sycl-post-link/emit_program_metadata.ll b/llvm/test/tools/sycl-post-link/emit_program_metadata.ll index 4dad16cacf80f..80dcdc74a2619 100644 --- a/llvm/test/tools/sycl-post-link/emit_program_metadata.ll +++ b/llvm/test/tools/sycl-post-link/emit_program_metadata.ll @@ -1,6 +1,6 @@ ; This test checks that the post-link tool generates SYCL program metadata. ; -; RUN: sycl-post-link -emit-program-metadata -device-globals -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -emit-program-metadata -device-globals -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files.table --check-prefixes CHECK-TABLE ; RUN: FileCheck %s -input-file=%t.files_0.prop --match-full-lines --check-prefixes CHECK-PROP diff --git a/llvm/test/tools/sycl-post-link/erase_used.ll b/llvm/test/tools/sycl-post-link/erase_used.ll index c8827bf7c70da..37aa550796f54 100644 --- a/llvm/test/tools/sycl-post-link/erase_used.ll +++ b/llvm/test/tools/sycl-post-link/erase_used.ll @@ -2,11 +2,11 @@ ; the output modules when splitting modules, creating a single row table, ; and outputing IR only ; -; RUN: sycl-post-link -split=kernel -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -split=kernel -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll ; RUN: FileCheck %s -input-file=%t.files_1.ll ; -; RUN: sycl-post-link -S -split=auto -symbols -split-esimd -lower-esimd -O2 -spec-const=emulation < %s -o %t.out.table +; RUN: sycl-post-link -properties -S -split=auto -symbols -split-esimd -lower-esimd -O2 -spec-const=emulation < %s -o %t.out.table ; RUN: FileCheck %s --input-file=%t.out_0.ll ; ; RUN: sycl-post-link -S -split=auto -ir-output-only < %s -o %t.out_ir_only.ll diff --git a/llvm/test/tools/sycl-post-link/erase_used_decl.ll b/llvm/test/tools/sycl-post-link/erase_used_decl.ll index 00e57c40da3b4..60eadd2a03c25 100644 --- a/llvm/test/tools/sycl-post-link/erase_used_decl.ll +++ b/llvm/test/tools/sycl-post-link/erase_used_decl.ll @@ -1,7 +1,7 @@ ; This test checks that the post-link tool doesn't incorrectly remove function ; declarations which are still in use while erasing the "llvm.used" global. ; -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll ; target triple = "spir64-unknown-unknown" diff --git a/llvm/test/tools/sycl-post-link/erase_used_decl_opaque.ll b/llvm/test/tools/sycl-post-link/erase_used_decl_opaque.ll index e9faef66c3463..4801c9945ba87 100644 --- a/llvm/test/tools/sycl-post-link/erase_used_decl_opaque.ll +++ b/llvm/test/tools/sycl-post-link/erase_used_decl_opaque.ll @@ -1,7 +1,7 @@ ; This test checks that the post-link tool doesn't incorrectly remove function ; declarations which are still in use while erasing the "llvm.used" global. ; -; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -split=auto -symbols -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll ; target triple = "spir64-unknown-unknown" diff --git a/llvm/test/tools/sycl-post-link/erase_used_opaque.ll b/llvm/test/tools/sycl-post-link/erase_used_opaque.ll index 0a562e8f2cf74..9f9158bdff087 100644 --- a/llvm/test/tools/sycl-post-link/erase_used_opaque.ll +++ b/llvm/test/tools/sycl-post-link/erase_used_opaque.ll @@ -2,11 +2,11 @@ ; the output modules when splitting modules, creating a single row table, ; and outputing IR only ; -; RUN: sycl-post-link -split=kernel -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -split=kernel -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll ; RUN: FileCheck %s -input-file=%t.files_1.ll ; -; RUN: sycl-post-link -S -split=auto -symbols -split-esimd -lower-esimd -O2 -spec-const=emulation < %s -o %t.out.table +; RUN: sycl-post-link -properties -S -split=auto -symbols -split-esimd -lower-esimd -O2 -spec-const=emulation < %s -o %t.out.table ; RUN: FileCheck %s --input-file=%t.out_0.ll ; ; RUN: sycl-post-link -S -split=auto -ir-output-only < %s -o %t.out_ir_only.ll diff --git a/llvm/test/tools/sycl-post-link/exclude_external_functions.ll b/llvm/test/tools/sycl-post-link/exclude_external_functions.ll index 7beff7f36113e..f9c7a3e58f376 100644 --- a/llvm/test/tools/sycl-post-link/exclude_external_functions.ll +++ b/llvm/test/tools/sycl-post-link/exclude_external_functions.ll @@ -2,7 +2,7 @@ ; dependencies to a function that can be imported do not cause the function ; to be added to a device image. -; RUN: sycl-post-link -symbols -support-dynamic-linking -split=kernel -S < %s -o %t.table +; RUN: sycl-post-link -properties -symbols -support-dynamic-linking -split=kernel -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYM-0 diff --git a/llvm/test/tools/sycl-post-link/exclude_external_functions_source.ll b/llvm/test/tools/sycl-post-link/exclude_external_functions_source.ll index a830ec31023c8..8c4cc75d15d6b 100644 --- a/llvm/test/tools/sycl-post-link/exclude_external_functions_source.ll +++ b/llvm/test/tools/sycl-post-link/exclude_external_functions_source.ll @@ -4,7 +4,7 @@ ; Also ensure that functions in the same source that can be imported do not get split into ; different images. -; RUN: sycl-post-link -symbols -support-dynamic-linking -split=source -S < %s -o %t.table +; RUN: sycl-post-link -properties -symbols -support-dynamic-linking -split=source -S < %s -o %t.table target triple = "spir64-unknown-unknown" diff --git a/llvm/test/tools/sycl-post-link/help.test b/llvm/test/tools/sycl-post-link/help.test index 6d41bc1f4fcac..2a74cd3b5797c 100644 --- a/llvm/test/tools/sycl-post-link/help.test +++ b/llvm/test/tools/sycl-post-link/help.test @@ -1,4 +1,4 @@ -// RUN: sycl-post-link --help | FileCheck %s +// RUN: sycl-post-link -properties --help | FileCheck %s CHECK: OVERVIEW: SYCL post-link device code processing tool. CHECK: This is a collection of utilities run on device code's LLVM IR before @@ -17,8 +17,8 @@ CHECK: for the SPIRV translator CHECK: Normally, the tool generates a number of files and "file table" CHECK: file listing all generated files in a table manner. For example, if CHECK: the input file 'example.bc' contains two kernels, then the command -CHECK: $ sycl-post-link --split=kernel --symbols --spec-const=native \ -CHECK: -o example.table example.bc +CHECK: $ sycl-post-link --properties --split=kernel --symbols \ +CHECK: --spec-const=native -o example.table example.bc CHECK: will produce 'example.table' file with the following content: CHECK: [Code|Properties|Symbols] CHECK: example_0.bc|example_0.prop|example_0.sym @@ -47,6 +47,7 @@ CHECK: -f - Enable binary output on terminals CHECK: --ir-output-only - Output single IR file CHECK: -o - Specifies an output file. Multiple output files can be specified. Additionally, a target may be specified alongside an output file, which has the effect that when module splitting is performed, the modules that are in that output table are filtered so those modules are compatible with the target. CHECK: --out-dir= - Directory where files listed in the result file table will be output +CHECK: --properties - generate module properties files CHECK: --spec-const= - lower and generate specialization constants information CHECK: =native - lower spec constants to native spirv instructions so that these values could be set at runtime CHECK: =emulation - remove specialization constants and replace it with emulation diff --git a/llvm/test/tools/sycl-post-link/multiple-filtered-outputs.ll b/llvm/test/tools/sycl-post-link/multiple-filtered-outputs.ll index 7c2ab6e91b925..f2abd47ecce0c 100644 --- a/llvm/test/tools/sycl-post-link/multiple-filtered-outputs.ll +++ b/llvm/test/tools/sycl-post-link/multiple-filtered-outputs.ll @@ -3,7 +3,7 @@ ; and that the output tables from inputs with target info have the modules ; that are not compatible with that target filtered out. -; RUN: sycl-post-link %s -symbols -split=auto \ +; RUN: sycl-post-link -properties %s -symbols -split=auto \ ; RUN: -o %t.table \ ; RUN: -o intel_gpu_pvc,%t-pvc.table \ ; RUN: -o intel_gpu_tgllp,%t-tgllp.table \ diff --git a/llvm/test/tools/sycl-post-link/no-args-to-eliminate.ll b/llvm/test/tools/sycl-post-link/no-args-to-eliminate.ll index 2b684f7f09a7e..274c5f0794400 100644 --- a/llvm/test/tools/sycl-post-link/no-args-to-eliminate.ll +++ b/llvm/test/tools/sycl-post-link/no-args-to-eliminate.ll @@ -1,7 +1,7 @@ ; This test ensures that sycl-post-link doesn't crash when kernel parameter ; optimization info metadata is empty ; -; RUN: sycl-post-link -emit-param-info -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -emit-param-info -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop ; ; CHECK: [SYCL/kernel param opt] diff --git a/llvm/test/tools/sycl-post-link/no-split-unused-func.ll b/llvm/test/tools/sycl-post-link/no-split-unused-func.ll index 09ad1d98afa22..fec85d4ad26e0 100644 --- a/llvm/test/tools/sycl-post-link/no-split-unused-func.ll +++ b/llvm/test/tools/sycl-post-link/no-split-unused-func.ll @@ -1,7 +1,7 @@ ; This test ensures that sycl-post-link will optimize away ; unused functions that are safe to remove even if there are no ; splits. -; RUN: sycl-post-link -split-esimd -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -split-esimd -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --implicit-check-not=foo ; CHECK: target datalayout diff --git a/llvm/test/tools/sycl-post-link/omit_kernel_args.ll b/llvm/test/tools/sycl-post-link/omit_kernel_args.ll index ef6d2c96b9392..0253e72c3addb 100644 --- a/llvm/test/tools/sycl-post-link/omit_kernel_args.ll +++ b/llvm/test/tools/sycl-post-link/omit_kernel_args.ll @@ -2,7 +2,7 @@ ; optimization info into a property file if the source IR contained ; corresponding metadata. ; -; RUN: sycl-post-link -emit-param-info -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -emit-param-info -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files.table --check-prefixes CHECK-TABLE ; RUN: FileCheck %s -input-file=%t.files_0.prop --match-full-lines --check-prefixes CHECK-PROP diff --git a/llvm/test/tools/sycl-post-link/registerallocmode.ll b/llvm/test/tools/sycl-post-link/registerallocmode.ll index a008d2593dd83..96035b4379c40 100644 --- a/llvm/test/tools/sycl-post-link/registerallocmode.ll +++ b/llvm/test/tools/sycl-post-link/registerallocmode.ll @@ -1,6 +1,6 @@ ; This test checks handling of RegisterAllocMode in SYCL post link -; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table ; CHECK: [Code|Properties|Symbols] diff --git a/llvm/test/tools/sycl-post-link/skip-properties-gen.ll b/llvm/test/tools/sycl-post-link/skip-properties-gen.ll new file mode 100644 index 0000000000000..6c400b38e6358 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/skip-properties-gen.ll @@ -0,0 +1,25 @@ +; This test verifies the behavior of the sycl-post-link tool without the -properties and -symbols options. +; In particular, we verify that the properties and symbols files are not added to the output table. +; +; RUN: sycl-post-link -split=source -S < %s -o %t.table +; RUN: FileCheck %s -input-file=%t.table -check-prefix=CHECK-TABLE + +; CHECK-TABLE: [Code] +; CHECK-TABLE: {{.*}}_0.ll +; CHECK-TABLE: {{.*}}_1.ll + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" +target triple = "spir64-unknown-unknown" + +; CHECK-FOO: define dso_local spir_func noundef void @foo +define dso_local spir_func noundef void @foo(i32 noundef %a, i32 noundef %b) local_unnamed_addr #0 { +entry: +ret void +} + +define dso_local spir_func noundef void @bar(i32 noundef %a, i32 noundef %b) #1 { +entry: +ret void +} +attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "sycl-module-id"="test.cpp" "sycl-grf-size"="128" } +attributes #1 = { convergent mustprogress noinline norecurse nounwind "sycl-module-id"="test.cpp" "sycl-grf-size"="256" } diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll index c7ab60148e1c5..1098628fa87f4 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll @@ -1,7 +1,7 @@ -; RUN: sycl-post-link --spec-const=native -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --spec-const=native -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst --spec-const=native -S < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst --spec-const=native -S < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; ; This test is intended to check that SpecConstantsPass is able to handle the ; situation where specialization constants with complex types such as arrays diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll index fb8a162d76c6c..21fdc1c835b79 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll @@ -1,5 +1,5 @@ ; RUN: sycl-post-link --ir-output-only --spec-const=native %s -S -o - | FileCheck %s -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst --spec-const=native %s -S 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst --spec-const=native %s -S 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; ; This test is intended to check that SpecConstantsPass is able to handle the ; situation where specialization constants have zeroinitializer in LLVM IR diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll index 113479edc448a..9efcf14a421a1 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020.ll @@ -1,11 +1,11 @@ -; RUN: sycl-post-link -spec-const=emulation < %s -S -o %t.table +; RUN: sycl-post-link -properties -spec-const=emulation < %s -S -o %t.table ; RUN: FileCheck %s -check-prefixes=CHECK,CHECK-DEF < %t_0.ll ; RUN: FileCheck %s --check-prefixes=CHECK-PROPS,CHECK-PROPS-DEF < %t_0.prop -; RUN: sycl-post-link -spec-const=native < %s -S -o %t.table +; RUN: sycl-post-link -properties -spec-const=native < %s -S -o %t.table ; RUN: FileCheck %s -check-prefixes=CHECK,CHECK-RT < %t_0.ll ; RUN: FileCheck %s --check-prefixes=CHECK-PROPS < %t_0.prop -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -spec-const=emulation < %s -S 2>&1 | FileCheck %s --check-prefixes=CHECK-LOG,CHECK-LOG-EMULATION %} -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -spec-const=native < %s -S 2>&1 | FileCheck %s --check-prefixes=CHECK-LOG,CHECK-LOG-NATIVE %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -spec-const=emulation < %s -S 2>&1 | FileCheck %s --check-prefixes=CHECK-LOG,CHECK-LOG-EMULATION %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -spec-const=native < %s -S 2>&1 | FileCheck %s --check-prefixes=CHECK-LOG,CHECK-LOG-NATIVE %} ; This test checks that the post link tool is able to correctly transform ; SYCL 2020 specialization constant intrinsics for different types in a device diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca-error.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca-error.ll index 07f74450d0375..daf0e0126dec5 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca-error.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca-error.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -spec-const=emulation %s 2>&1 | FileCheck %s +; RUN: sycl-post-link -properties -spec-const=emulation %s 2>&1 | FileCheck %s ; This test checks the `-spec-const` pass on SPIR-V targets and emulation mode, ; i.e., on AOT SPIR-V targets. In this scenario, 'llvm.sycl.alloca' intrinsics diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca.ll index 2bceb13b43843..f8d8694c62ff0 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-alloca.ll @@ -1,8 +1,8 @@ -; RUN: sycl-post-link -spec-const=native < %s -S -o %t.table +; RUN: sycl-post-link -properties -spec-const=native < %s -S -o %t.table ; RUN: FileCheck %s -check-prefixes=CHECK,CHECK-RT < %t_0.ll ; RUN: FileCheck %s --check-prefixes=CHECK-PROPS < %t_0.prop -; RUN: sycl-post-link -spec-const=emulation < %s -S -o %t.table +; RUN: sycl-post-link -properties -spec-const=emulation < %s -S -o %t.table ; RUN: FileCheck %s -check-prefixes=CHECK,CHECK-EMULATION < %t_0.ll ; This test checks that the post link tool is able to correctly transform diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll index 1511c68a7e381..4eb6862b2ca44 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll @@ -1,7 +1,7 @@ -; RUN: sycl-post-link --spec-const=native -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --spec-const=native -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst --spec-const=native -S < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst --spec-const=native -S < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; ; This test is intended to check that SpecConstantsPass is able to handle the ; situation where specialization constants with complex types such as structs diff --git a/llvm/test/tools/sycl-post-link/spec-constants/bool.ll b/llvm/test/tools/sycl-post-link/spec-constants/bool.ll index 4d1d63048b60b..7981450a12f00 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/bool.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/bool.ll @@ -2,8 +2,8 @@ ; RUN: FileCheck %s --input-file=%t.ll --implicit-check-not "call i8 bitcast" --check-prefixes=CHECK,CHECK-RT ; RUN: sycl-post-link -spec-const=emulation -S < %s --ir-output-only -o %t.ll ; RUN: FileCheck %s --input-file=%t.ll --check-prefixes=CHECK,CHECK-DEF -; RUN: %if asserts %{sycl-post-link -debug-only=SpecConst -spec-const=native -S < %s 2>&1 | FileCheck %s --check-prefixes=CHECK-LOG,CHECK-LOG-NATIVE %} -; RUN: %if asserts %{sycl-post-link -debug-only=SpecConst -spec-const=emulation -S < %s 2>&1 | FileCheck %s --check-prefixes=CHECK-LOG,CHECK-LOG-EMULATION %} +; RUN: %if asserts %{sycl-post-link -properties -debug-only=SpecConst -spec-const=native -S < %s 2>&1 | FileCheck %s --check-prefixes=CHECK-LOG,CHECK-LOG-NATIVE %} +; RUN: %if asserts %{sycl-post-link -properties -debug-only=SpecConst -spec-const=emulation -S < %s 2>&1 | FileCheck %s --check-prefixes=CHECK-LOG,CHECK-LOG-EMULATION %} ; CHECK-LABEL: void @kernel_A diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll index 538385b0401ef..2ef1141cbd740 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll @@ -1,6 +1,6 @@ ; RUN: sycl-post-link -spec-const=native --ir-output-only < %s -S -o - \ ; RUN: | FileCheck %s --implicit-check-not "call {{.*}} __sycl_getComposite2020SpecConstantValue" -; RUN: %if asserts %{ sycl-post-link -spec-const=native -debug-only=SpecConst < %s -S 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -spec-const=native -debug-only=SpecConst < %s -S 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; ; This test is intended to check that sycl-post-link tool is capable of handling ; composite specialization constants by lowering them into a set of SPIR-V diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-default-value-padding.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-default-value-padding.ll index 968e94547a194..bd376a260a719 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-default-value-padding.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/composite-default-value-padding.ll @@ -1,6 +1,6 @@ -; RUN: sycl-post-link -spec-const=emulation < %s -o %t.files.table +; RUN: sycl-post-link -properties -spec-const=emulation < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -spec-const=emulation < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -spec-const=emulation < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; ; This test checks that composite specialization constants with padding gets the ; correct padding in their default values to prevent values being inserted at diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll index 87715cb2db45a..a4bb392d7ff67 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/composite-no-sret.ll @@ -1,6 +1,6 @@ ; RUN: sycl-post-link -spec-const=native --ir-output-only < %s -S -o - \ ; RUN: | FileCheck %s --implicit-check-not "call {{.*}} __sycl_getCompositeSpecConstantValue" --implicit-check-not "call {{.*}} __sycl_getComposite2020SpecConstantValue" -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -spec-const=native < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -spec-const=native < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; CHECK: %[[#NS0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32 ; CHECK: %[[#NS1:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 1]], i32 42) diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-padding-desc.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-padding-desc.ll index 1064e662eff01..b525e8c5c2aaa 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-padding-desc.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/composite-padding-desc.ll @@ -1,6 +1,6 @@ -; RUN: sycl-post-link -spec-const=native < %s -o %t.files.table +; RUN: sycl-post-link -properties -spec-const=native < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -spec-const=native < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -spec-const=native < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; ; This test checks that composite specialization constants with implicit padding ; at the end of the composite type will have an additional padding descriptor at diff --git a/llvm/test/tools/sycl-post-link/spec-constants/default-value/SYCL-alloca.ll b/llvm/test/tools/sycl-post-link/spec-constants/default-value/SYCL-alloca.ll index c1fa304fbf159..30415be30a90c 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/default-value/SYCL-alloca.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/default-value/SYCL-alloca.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts +; RUN: sycl-post-link -properties -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts ; RUN: FileCheck %s -input-file %t_1.ll --implicit-check-not="SpecConst" ; This test checks that the post link tool is able to correctly transform diff --git a/llvm/test/tools/sycl-post-link/spec-constants/default-value/bool.ll b/llvm/test/tools/sycl-post-link/spec-constants/default-value/bool.ll index 2a8a82f262611..d96cfa1f333f2 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/default-value/bool.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/default-value/bool.ll @@ -1,8 +1,8 @@ ; Test checks handling of bool specialization constant. -; RUN: sycl-post-link -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts +; RUN: sycl-post-link -properties -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts ; RUN: FileCheck %s -input-file %t_1.ll --implicit-check-not="SpecConst" -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -split=auto -spec-const=native -S %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -split=auto -spec-const=native -S %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; CHECK: %bool1 = trunc i8 1 to i1 ; CHECK: %frombool = zext i1 %bool1 to i8 diff --git a/llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll b/llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll index 01af2cd741dde..d38e2469b7e2b 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll @@ -2,7 +2,7 @@ ; It checks scalar, sret and "return by value" versions of SpecConstant functions. ; Also test checks generated symbols. -; RUN: sycl-post-link -split=auto -spec-const=native -symbols -S -o %t.table %s -generate-device-image-default-spec-consts +; RUN: sycl-post-link -properties -split=auto -spec-const=native -symbols -S -o %t.table %s -generate-device-image-default-spec-consts ; RUN: FileCheck %s -input-file %t.table -check-prefix=CHECK-TABLE ; RUN: FileCheck %s -input-file %t_0.prop -check-prefix=CHECK-PROP0 ; RUN: FileCheck %s -input-file %t_1.prop -check-prefix=CHECK-PROP1 @@ -10,7 +10,7 @@ ; RUN: FileCheck %s -input-file %t_1.ll -check-prefix=CHECK-IR1 --implicit-check-not "SpecConstant" ; RUN: FileCheck %s -input-file %t_0.sym -check-prefix=CHECK-SYM0 ; RUN: FileCheck %s -input-file %t_1.sym -check-prefix=CHECK-SYM1 -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -split=auto -spec-const=native -symbols -S %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -split=auto -spec-const=native -symbols -S %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; CHECK-TABLE: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym ; CHECK-TABLE: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym diff --git a/llvm/test/tools/sycl-post-link/spec-constants/default-value/esimd.ll b/llvm/test/tools/sycl-post-link/spec-constants/default-value/esimd.ll index dbd995bd0ca44..5afb3b4aac1e4 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/default-value/esimd.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/default-value/esimd.ll @@ -1,10 +1,10 @@ ; Test checks generation of device image of esimd kernel. -; RUN: sycl-post-link -split=auto -split-esimd -lower-esimd -O2 -spec-const=native -o %t.table %s -generate-device-image-default-spec-consts +; RUN: sycl-post-link -properties -split=auto -split-esimd -lower-esimd -O2 -spec-const=native -o %t.table %s -generate-device-image-default-spec-consts ; RUN: FileCheck %s -input-file=%t.table -check-prefix=CHECK-TABLE ; RUN: FileCheck %s -input-file=%t_1.prop -check-prefix=CHECK-PROP ; RUN: FileCheck %s -input-file=%t_esimd_1.prop -check-prefix=CHECK-ESIMD-PROP -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -split=auto -split-esimd -lower-esimd -O2 -spec-const=native %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -split=auto -split-esimd -lower-esimd -O2 -spec-const=native %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; CHECK-TABLE: {{.*}}_esimd_0.bc|{{.*}}_esimd_0.prop ; CHECK-TABLE: {{.*}}_0.bc|{{.*}}_0.prop diff --git a/llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-kernel.ll b/llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-kernel.ll index e985687197a68..98d291b342e3c 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-kernel.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-kernel.ll @@ -1,6 +1,6 @@ ; Test checks generation of device images for splitted kernels. -; RUN: sycl-post-link -split=kernel -o %t.table %s -generate-device-image-default-spec-consts +; RUN: sycl-post-link -properties -split=kernel -o %t.table %s -generate-device-image-default-spec-consts ; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE ; RUN: cat %t_0.prop | FileCheck %s -check-prefix=CHECK-PROP0 ; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1 diff --git a/llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-source.ll b/llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-source.ll index ad38aa3ba7489..5cdc2d2aed1d5 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-source.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-source.ll @@ -1,6 +1,6 @@ ; Test checks generation of device images for splitted kernels by source. -; RUN: sycl-post-link -split=source -o %t.table %s -generate-device-image-default-spec-consts +; RUN: sycl-post-link -properties -split=source -o %t.table %s -generate-device-image-default-spec-consts ; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE ; RUN: cat %t_0.prop | FileCheck %s -check-prefix=CHECK-PROP0 ; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1 diff --git a/llvm/test/tools/sycl-post-link/spec-constants/default-value/struct-with-padding.ll b/llvm/test/tools/sycl-post-link/spec-constants/default-value/struct-with-padding.ll index 51dc1c557fd7f..8b9f3b72f3294 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/default-value/struct-with-padding.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/default-value/struct-with-padding.ll @@ -1,10 +1,10 @@ ; Test checks that struct with padding is handled correctly. -; RUN: sycl-post-link -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts +; RUN: sycl-post-link -properties -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts ; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE ; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1 ; RUN: cat %t_1.ll | FileCheck %s -check-prefix=CHECK-IR1 --implicit-check-not SpecConstant -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -split=auto -spec-const=native -S %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -split=auto -spec-const=native -S %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; CHECK-TABLE: {{.*}}_0.ll|{{.*}}_0.prop ; CHECK-TABLE: {{.*}}_1.ll|{{.*}}_1.prop diff --git a/llvm/test/tools/sycl-post-link/spec-constants/nested-struct.ll b/llvm/test/tools/sycl-post-link/spec-constants/nested-struct.ll index f0bd6b8df6c66..d68bf98094103 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/nested-struct.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/nested-struct.ll @@ -1,9 +1,9 @@ ; This test demonstrates that multiple padding elements can be ; inserted in the spec constant metadata -; RUN: sycl-post-link --spec-const=native -S %s -o %t.table +; RUN: sycl-post-link -properties --spec-const=native -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -spec-const=native < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -spec-const=native < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; CHECK: %[[#SCV1:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID1:]], i8 120) ; CHECK: %[[#SCV2:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID2:]], i8 121) diff --git a/llvm/test/tools/sycl-post-link/spec-constants/remove-dead-private-constants.ll b/llvm/test/tools/sycl-post-link/spec-constants/remove-dead-private-constants.ll index fb3b69228599e..0b96ef7172b1f 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/remove-dead-private-constants.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/remove-dead-private-constants.ll @@ -1,10 +1,10 @@ ; Test checks the content of simple generated device image. ; It checks for removal of unused private constants. -; RUN: sycl-post-link -split=auto -spec-const=native -symbols -S -o %t.table %s -generate-device-image-default-spec-consts +; RUN: sycl-post-link -properties -split=auto -spec-const=native -symbols -S -o %t.table %s -generate-device-image-default-spec-consts ; RUN: FileCheck %s -input-file %t_0.ll -check-prefix=CHECK-IR0 ; RUN: FileCheck %s -input-file %t_1.ll -check-prefix=CHECK-IR1 -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst -split=auto -spec-const=native -symbols -S %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -split=auto -spec-const=native -symbols -S %s -generate-device-image-default-spec-consts 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; CHECK-IR0-NOT: @__usid_str = private ; CHECK-IR1-NOT: @__usid_str = private diff --git a/llvm/test/tools/sycl-post-link/spec-constants/struct-with-padding-in-the-middle.ll b/llvm/test/tools/sycl-post-link/spec-constants/struct-with-padding-in-the-middle.ll index 57ddbfe23b55f..2dc002140362f 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/struct-with-padding-in-the-middle.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/struct-with-padding-in-the-middle.ll @@ -1,6 +1,6 @@ -; RUN: sycl-post-link --spec-const=native -S %s -o %t.table +; RUN: sycl-post-link -properties --spec-const=native -S %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst --spec-const=native -S %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst --spec-const=native -S %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; LLVM IR for this test is produced from the following SYCL code snippet: ; ; #include diff --git a/llvm/test/tools/sycl-post-link/spec-constants/struct-with-undef-padding-2.ll b/llvm/test/tools/sycl-post-link/spec-constants/struct-with-undef-padding-2.ll index 233933d87f6c2..460699ae68ed3 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/struct-with-undef-padding-2.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/struct-with-undef-padding-2.ll @@ -34,9 +34,9 @@ ; a situation, where spec constant default value contains less elements than ; spec constant type, due to padding inserted by a compiler. ; -; RUN: sycl-post-link --spec-const=native -S < %s -o %t.files.table +; RUN: sycl-post-link -properties --spec-const=native -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.ll -; RUN: %if asserts %{ sycl-post-link -debug-only=SpecConst --spec-const=native -S < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} +; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst --spec-const=native -S < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; ; CHECK: %[[#A:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID:]], float 0x40091EB860000000) ; CHECK: %[[#B:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID+1]], i32 42) diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/basic-esimd-lower.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/basic-esimd-lower.ll index 157742f6618f4..3773fd1048ba2 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/basic-esimd-lower.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/basic-esimd-lower.ll @@ -3,19 +3,19 @@ ; for ESIMD kernels in any case. ; No lowering -; RUN: sycl-post-link -split-esimd -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-NO-LOWERING ; Default lowering -; RUN: sycl-post-link -split-esimd -lower-esimd -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -lower-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-O2 ; -O2 lowering -; RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -lower-esimd -O2 -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-O2 ; -O0 lowering -; RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -lower-esimd -O0 -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-O0 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/basic-sycl-esimd-split.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/basic-sycl-esimd-split.ll index 68977a001fbf7..25258107ad0e0 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/basic-sycl-esimd-split.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/basic-sycl-esimd-split.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split-esimd -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-SYCL-IR ; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-ESIMD-IR diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/invoke-esimd-double.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/invoke-esimd-double.ll index a614184d2a83c..c7aeb4966fdbe 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/invoke-esimd-double.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/invoke-esimd-double.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --emit-only-kernels-as-entry-points -symbols -split=auto -S < %s -o %t.table +; RUN: sycl-post-link -properties --emit-only-kernels-as-entry-points -symbols -split=auto -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefixes CHECK-TABLE ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/lower-with-no-esimd-entry.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/lower-with-no-esimd-entry.ll index 63a5ffb1fe626..cdfb4bba20c04 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/lower-with-no-esimd-entry.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/lower-with-no-esimd-entry.ll @@ -1,7 +1,7 @@ ; This test checks to see if ESIMD lowering is performed even without the ; the presence of ESIMD entry points. -; RUN: sycl-post-link -symbols -lower-esimd -split=auto -S < %s -o %t.table +; RUN: sycl-post-link -properties -symbols -lower-esimd -split=auto -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefixes CHECK-TABLE ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYMS ; RUN: FileCheck %s -input-file=%t_0.ll diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/nbarriers-metadata.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/nbarriers-metadata.ll index 01655f8c5c70b..20ae18777b559 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/nbarriers-metadata.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/nbarriers-metadata.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split-esimd -lower-esimd -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -lower-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_esimd_0.ll target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/no-sycl-esimd-split-shared-func.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/no-sycl-esimd-split-shared-func.ll index 50cdcc5dec8e1..3fead6c4029c0 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/no-sycl-esimd-split-shared-func.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/no-sycl-esimd-split-shared-func.ll @@ -5,7 +5,7 @@ ; making sure no functions are shared by the callgraphs (currently required by ; IGC) -; RUN: sycl-post-link -lower-esimd -symbols -split=auto -S < %s -o %t.table +; RUN: sycl-post-link -properties -lower-esimd -symbols -split=auto -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table --check-prefixes CHECK-TABLE ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-HELPERS-SYM-1 ; RUN: FileCheck %s -input-file=%t_esimd_1.sym --check-prefixes CHECK-HELPERS-SYM-2 diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/no-sycl-esimd-split.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/no-sycl-esimd-split.ll index 0488297b7fabe..13485a015c679 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/no-sycl-esimd-split.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/no-sycl-esimd-split.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=source -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR-0 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR-1 diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/optnone.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/optnone.ll index aa0fb46bc8c03..14e64440ce97e 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/optnone.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/optnone.ll @@ -1,8 +1,8 @@ ; This ensures we remove optnone from ESIMD functions unless they are SIMT or we didn't split ESIMD code out. -; RUN: sycl-post-link -split-esimd -lower-esimd -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -lower-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK,CHECK-ESIMD-SPLIT -; RUN: sycl-post-link -lower-esimd -S < %s -o %t1.table +; RUN: sycl-post-link -properties -lower-esimd -S < %s -o %t1.table ; RUN: FileCheck %s -input-file=%t1_esimd_0.ll --check-prefixes CHECK,CHECK-NO-ESIMD-SPLIT target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" target triple = "spir64-unknown-unknown" diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-per-kernel.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-per-kernel.ll index 4782557b02615..eabed45ec75f0 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-per-kernel.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-per-kernel.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split-esimd -split=kernel -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -split=kernel -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-SYCL-IR-0 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-SYCL-IR-1 diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-per-source.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-per-source.ll index efc08cbc7191c..610e52b8476ce 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-per-source.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-per-source.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split-esimd -split=source -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -split=source -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-SYCL-IR-0 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-SYCL-IR-1 diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-symbols.ll b/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-symbols.ll index 4f7d001321077..cf0f42ab9a063 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-symbols.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-split-symbols.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split-esimd -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM ; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-SYM diff --git a/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-and-lower-esimd.ll b/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-and-lower-esimd.ll index 44f5d091926b8..55a3a7aad9c3b 100644 --- a/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-and-lower-esimd.ll +++ b/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-and-lower-esimd.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split=auto -split-esimd -lower-esimd -O0 -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=auto -split-esimd -lower-esimd -O0 -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_esimd_0.ll ; This test checks that unreferenced functions with sycl-module-id diff --git a/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-kernel.ll b/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-kernel.ll index 2f2a24c8cad39..86ec87a808c34 100644 --- a/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-kernel.ll +++ b/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-kernel.ll @@ -1,7 +1,7 @@ ; This test checks handling of unreferenced functions with sycl-module-id ; attribute with splitting in per-kernel mode. -; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR0 ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYM0 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR1 @@ -9,7 +9,7 @@ ; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-IR2 ; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-SYM2 -; RUN: sycl-post-link -split=kernel -emit-only-kernels-as-entry-points -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=kernel -emit-only-kernels-as-entry-points -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR1 ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYM1 ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR0 diff --git a/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source1.ll b/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source1.ll index 074ef073d25c4..461edbe4aa4c5 100644 --- a/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source1.ll +++ b/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source1.ll @@ -1,13 +1,13 @@ ; This test checks handling of unreferenced functions with sycl-module-id ; attribute with splitting in per-source mode. -; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR2 ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYM2 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR1 ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYM1 -; RUN: sycl-post-link -split=source -emit-only-kernels-as-entry-points -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -emit-only-kernels-as-entry-points -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYM2 ; RUN: FileCheck %s -input-file=%t.table --check-prefixes CHECK-TABLE ; CHECK-TABLE: [Code|Properties|Symbols] diff --git a/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source2.ll b/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source2.ll index ea93eb7536b67..c020fd8d7b0be 100644 --- a/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source2.ll +++ b/llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source2.ll @@ -1,13 +1,13 @@ ; This test checks handling of referenced SYCL_EXTERNAL functions with ; sycl-module-id attribute with splitting in per-source mode. -; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-IR0 ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYM0 ; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-IR1 ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYM1 -; RUN: sycl-post-link -split=source -emit-only-kernels-as-entry-points -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -emit-only-kernels-as-entry-points -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYM0 ; RUN: FileCheck %s -input-file=%t.table --check-prefixes CHECK-TABLE ; CHECK-TABLE: [Code|Properties|Symbols] diff --git a/llvm/test/tools/sycl-post-link/sycl-grf-size.ll b/llvm/test/tools/sycl-post-link/sycl-grf-size.ll index 9573b4bc1a4ab..28bb07e18fae7 100644 --- a/llvm/test/tools/sycl-post-link/sycl-grf-size.ll +++ b/llvm/test/tools/sycl-post-link/sycl-grf-size.ll @@ -1,6 +1,6 @@ ; This test checks handling of sycl-grf-size in SYCL post link -; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table ; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR --implicit-check-not='__ESIMD_kernel()' ; RUN: FileCheck %s -input-file=%t_esimd_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP diff --git a/llvm/test/tools/sycl-post-link/sycl-opt-level-external-funcs.ll b/llvm/test/tools/sycl-post-link/sycl-opt-level-external-funcs.ll index 053760fa3b86f..0f8ffe0f6644c 100644 --- a/llvm/test/tools/sycl-post-link/sycl-opt-level-external-funcs.ll +++ b/llvm/test/tools/sycl-post-link/sycl-opt-level-external-funcs.ll @@ -9,7 +9,7 @@ ; - module with 'bar' (as entry point) with 'optLevel' set to 2 (taken from ; 'bar') -; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-OPT-LEVEL-PROP-0 ; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-OPT-LEVEL-PROP-1 diff --git a/llvm/test/tools/sycl-post-link/sycl-opt-level.ll b/llvm/test/tools/sycl-post-link/sycl-opt-level.ll index 8967d3bc02621..071c76193f094 100644 --- a/llvm/test/tools/sycl-post-link/sycl-opt-level.ll +++ b/llvm/test/tools/sycl-post-link/sycl-opt-level.ll @@ -4,7 +4,7 @@ ; on their optimization levels. ; sycl-post-link adds 'optLevel' property to the device binary -; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: sycl-post-link -properties -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-OPT-LEVEL-PROP-0 ; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-OPT-LEVEL-PROP-1 diff --git a/llvm/test/tools/sycl-post-link/sycl-post-link-test.ll b/llvm/test/tools/sycl-post-link/sycl-post-link-test.ll index cfd6b2c644419..269276c7c0ff4 100644 --- a/llvm/test/tools/sycl-post-link/sycl-post-link-test.ll +++ b/llvm/test/tools/sycl-post-link/sycl-post-link-test.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S < %s -o %t.table +; RUN: sycl-post-link -properties -split-esimd -lower-esimd -O0 -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_esimd_0.ll ; This test checks that IR code below can be successfully processed by ; sycl-post-link. In this IR no extractelement instruction and no casting are used diff --git a/llvm/test/tools/sycl-post-link/sym_but_no_split.ll b/llvm/test/tools/sycl-post-link/sym_but_no_split.ll index ca624801f1832..7cbcda34f996d 100644 --- a/llvm/test/tools/sycl-post-link/sym_but_no_split.ll +++ b/llvm/test/tools/sycl-post-link/sym_but_no_split.ll @@ -2,7 +2,7 @@ ; table and a symbol file for an input module with two kernels when no code ; splitting is requested. ; -; RUN: sycl-post-link -symbols -spec-const=native -S < %s -o %t.files.table +; RUN: sycl-post-link -properties -symbols -spec-const=native -S < %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files.table --check-prefixes CHECK-TABLE ; RUN: FileCheck %s -input-file=%t.files_0.sym --match-full-lines --check-prefixes CHECK-SYM diff --git a/llvm/tools/sycl-post-link/CMakeLists.txt b/llvm/tools/sycl-post-link/CMakeLists.txt index cfb9b1a27560f..2a8f99795d599 100644 --- a/llvm/tools/sycl-post-link/CMakeLists.txt +++ b/llvm/tools/sycl-post-link/CMakeLists.txt @@ -24,9 +24,6 @@ include_directories( add_llvm_tool(sycl-post-link sycl-post-link.cpp - SpecConstants.cpp - SYCLDeviceLibReqMask.cpp - SYCLKernelParamOptInfo.cpp ADDITIONAL_HEADER_DIRS ${LLVMGenXIntrinsics_SOURCE_DIR}/GenXIntrinsics/include ${LLVMGenXIntrinsics_BINARY_DIR}/GenXIntrinsics/include diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index d2c29d10aea7f..252f070652c4c 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -13,11 +13,6 @@ // - specialization constant intrinsic transformation //===----------------------------------------------------------------------===// -#include "SYCLDeviceLibReqMask.h" -#include "SYCLKernelParamOptInfo.h" -#include "SpecConstants.h" -#include "Support.h" - #include "llvm/ADT/SmallSet.h" #include "llvm/ADT/StringRef.h" #include "llvm/Analysis/AssumptionCache.h" @@ -35,6 +30,7 @@ #include "llvm/Linker/Linker.h" #include "llvm/Passes/PassBuilder.h" #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" +#include "llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h" #include "llvm/SYCLLowerIR/DeviceConfigFile.hpp" #include "llvm/SYCLLowerIR/DeviceGlobals.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" @@ -42,9 +38,10 @@ #include "llvm/SYCLLowerIR/HostPipes.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" #include "llvm/SYCLLowerIR/ModuleSplitter.h" -#include "llvm/SYCLLowerIR/SYCLDeviceRequirements.h" #include "llvm/SYCLLowerIR/SYCLUtils.h" #include "llvm/SYCLLowerIR/SanitizeDeviceGlobal.h" +#include "llvm/SYCLLowerIR/SpecConstants.h" +#include "llvm/SYCLLowerIR/Support.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/InitLLVM.h" @@ -64,15 +61,13 @@ #include "llvm/Transforms/Utils/GlobalStatus.h" #include -#include #include -#include #include -#include #include #include using namespace llvm; +using namespace llvm::sycl; using string_vector = std::vector; @@ -201,6 +196,10 @@ cl::opt SplitMode( cl::opt DoSymGen{"symbols", cl::desc("generate exported symbol files"), cl::cat(PostLinkCat)}; +cl::opt DoPropGen{"properties", + cl::desc("generate module properties files"), + cl::cat(PostLinkCat)}; + enum SpecConstLowerMode { SC_NATIVE_MODE, SC_EMULATION_MODE }; cl::opt SpecConstLower{ @@ -251,14 +250,6 @@ cl::opt GenerateDeviceImageWithDefaultSpecConsts{ "replaced with default values from specialization id(s)."), cl::cat(PostLinkCat)}; -struct GlobalBinImageProps { - bool EmitKernelParamInfo; - bool EmitProgramMetadata; - bool EmitExportedSymbols; - bool EmitImportedSymbols; - bool EmitDeviceGlobalPropSet; -}; - struct IrPropSymFilenameTriple { std::string Ir; std::string Prop; @@ -273,111 +264,6 @@ void writeToFile(const std::string &Filename, const std::string &Content) { OS.close(); } -// This function traverses over reversed call graph by BFS algorithm. -// It means that an edge links some function @func with functions -// which contain call of function @func. It starts from -// @StartingFunction and lifts up until it reach all reachable functions, -// or it reaches some function containing "referenced-indirectly" attribute. -// If it reaches "referenced-indirectly" attribute than it returns an empty -// Optional. -// Otherwise, it returns an Optional containing a list of reached -// SPIR kernel function's names. -std::optional> -traverseCGToFindSPIRKernels(const Function *StartingFunction) { - std::queue FunctionsToVisit; - std::unordered_set VisitedFunctions; - FunctionsToVisit.push(StartingFunction); - std::vector KernelNames; - - while (!FunctionsToVisit.empty()) { - const Function *F = FunctionsToVisit.front(); - FunctionsToVisit.pop(); - - auto InsertionResult = VisitedFunctions.insert(F); - // It is possible that we insert some particular function several - // times in functionsToVisit queue. - if (!InsertionResult.second) - continue; - - for (const auto *U : F->users()) { - const CallInst *CI = dyn_cast(U); - if (!CI) - continue; - - const Function *ParentF = CI->getFunction(); - - if (VisitedFunctions.count(ParentF)) - continue; - - if (ParentF->hasFnAttribute("referenced-indirectly")) - return {}; - - if (ParentF->getCallingConv() == CallingConv::SPIR_KERNEL) - KernelNames.push_back(ParentF->getName()); - - FunctionsToVisit.push(ParentF); - } - } - - return {std::move(KernelNames)}; -} - -std::vector getKernelNamesUsingAssert(const Module &M) { - auto *DevicelibAssertFailFunction = M.getFunction("__devicelib_assert_fail"); - if (!DevicelibAssertFailFunction) - return {}; - - auto TraverseResult = - traverseCGToFindSPIRKernels(DevicelibAssertFailFunction); - - if (TraverseResult.has_value()) - return std::move(*TraverseResult); - - // Here we reached "referenced-indirectly", so we need to find all kernels and - // return them. - std::vector SPIRKernelNames; - for (const Function &F : M) { - if (F.getCallingConv() == CallingConv::SPIR_KERNEL) - SPIRKernelNames.push_back(F.getName()); - } - - return SPIRKernelNames; -} - -bool isModuleUsingAsan(const Module &M) { - NamedMDNode *MD = M.getNamedMetadata("device.sanitizer"); - if (MD == nullptr) - return false; - assert(MD->getNumOperands() != 0); - auto *MDVal = cast(MD->getOperand(0)->getOperand(0)); - return MDVal->getString() == "asan"; -} - -// Gets work_group_num_dim information for function Func, conviniently 0 if -// metadata is not present. -uint32_t getKernelWorkGroupNumDim(const Function &Func) { - MDNode *MaxDimMD = Func.getMetadata("work_group_num_dim"); - if (!MaxDimMD) - return 0; - assert(MaxDimMD->getNumOperands() == 1 && "Malformed node."); - return mdconst::extract(MaxDimMD->getOperand(0))->getZExtValue(); -} - -// Gets reqd_work_group_size information for function Func. -std::vector getKernelReqdWorkGroupSizeMetadata(const Function &Func) { - MDNode *ReqdWorkGroupSizeMD = Func.getMetadata("reqd_work_group_size"); - if (!ReqdWorkGroupSizeMD) - return {}; - size_t NumOperands = ReqdWorkGroupSizeMD->getNumOperands(); - assert(NumOperands >= 1 && NumOperands <= 3 && - "reqd_work_group_size does not have between 1 and 3 operands."); - std::vector OutVals; - OutVals.reserve(NumOperands); - for (const MDOperand &MDOp : ReqdWorkGroupSizeMD->operands()) - OutVals.push_back(mdconst::extract(MDOp)->getZExtValue()); - return OutVals; -} - // Creates a filename based on current output filename, given extension, // sequential ID and suffix. std::string makeResultFileName(Twine Ext, int I, StringRef Suffix) { @@ -420,221 +306,9 @@ std::string saveModuleIR(Module &M, int I, StringRef Suff) { std::string saveModuleProperties(module_split::ModuleDesc &MD, const GlobalBinImageProps &GlobProps, int I, StringRef Suff) { - using PropSetRegTy = llvm::util::PropertySetRegistry; - PropSetRegTy PropSet; - Module &M = MD.getModule(); - { - uint32_t MRMask = getSYCLDeviceLibReqMask(M); - std::map RMEntry = {{"DeviceLibReqMask", MRMask}}; - PropSet.add(PropSetRegTy::SYCL_DEVICELIB_REQ_MASK, RMEntry); - } - { - PropSet.add(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS, - MD.getOrComputeDeviceRequirements().asMap()); - } - if (MD.Props.SpecConstsMet) { - // extract spec constant maps per each module - SpecIDMapTy TmpSpecIDMap; - SpecConstantsPass::collectSpecConstantMetadata(M, TmpSpecIDMap); - PropSet.add(PropSetRegTy::SYCL_SPECIALIZATION_CONSTANTS, TmpSpecIDMap); - - // Add property with the default values of spec constants - std::vector DefaultValues; - SpecConstantsPass::collectSpecConstantDefaultValuesMetadata(M, - DefaultValues); - PropSet.add(PropSetRegTy::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES, "all", - DefaultValues); - } - if (GlobProps.EmitKernelParamInfo) { - // extract kernel parameter optimization info per module - ModuleAnalysisManager MAM; - // Register required analysis - MAM.registerPass([&] { return PassInstrumentationAnalysis(); }); - // Register the payload analysis - - MAM.registerPass([&] { return SYCLKernelParamOptInfoAnalysis(); }); - SYCLKernelParamOptInfo PInfo = - MAM.getResult(M); - - // convert analysis results into properties and record them - llvm::util::PropertySet &Props = - PropSet[PropSetRegTy::SYCL_KERNEL_PARAM_OPT_INFO]; - - for (const auto &NameInfoPair : PInfo) { - const llvm::BitVector &Bits = NameInfoPair.second; - if (Bits.empty()) - continue; // Nothing to add - - const llvm::ArrayRef Arr = Bits.getData(); - const unsigned char *Data = - reinterpret_cast(Arr.begin()); - llvm::util::PropertyValue::SizeTy DataBitSize = Bits.size(); - Props.insert(std::make_pair( - NameInfoPair.first, llvm::util::PropertyValue(Data, DataBitSize))); - } - } - if (GlobProps.EmitExportedSymbols) { - // extract exported functions if any and save them into property set - for (const auto *F : MD.entries()) { - // TODO FIXME some of SYCL/ESIMD functions maybe marked with __regcall CC, - // so they won't make it into the export list. Should the check be - // F->getCallingConv() != CallingConv::SPIR_KERNEL? - if (F->getCallingConv() == CallingConv::SPIR_FUNC) { - PropSet.add(PropSetRegTy::SYCL_EXPORTED_SYMBOLS, F->getName(), - /*PropVal=*/true); - } - } - } - - if (GlobProps.EmitImportedSymbols) { - // record imported functions in the property set - for (const auto &F : M) { - if ( // A function that can be imported may still be defined in one split - // image. Only add import property if this is not the image where the - // function is defined. - F.isDeclaration() && module_split::canBeImportedFunction(F)) { - - // StripDeadPrototypes is called during module splitting - // cleanup. At this point all function decls should have uses. - assert(!F.use_empty() && "Function F has no uses"); - - PropSet.add(PropSetRegTy::SYCL_IMPORTED_SYMBOLS, F.getName(), - /*PropVal=*/true); - } - } - } - - // Metadata names may be composite so we keep them alive until the - // properties have been written. - SmallVector MetadataNames; - - if (GlobProps.EmitProgramMetadata) { - // Add reqd_work_group_size and work_group_num_dim information to - // program metadata. - for (const Function &Func : M.functions()) { - std::vector KernelReqdWorkGroupSize = - getKernelReqdWorkGroupSizeMetadata(Func); - if (!KernelReqdWorkGroupSize.empty()) { - MetadataNames.push_back(Func.getName().str() + "@reqd_work_group_size"); - PropSet.add(PropSetRegTy::SYCL_PROGRAM_METADATA, MetadataNames.back(), - KernelReqdWorkGroupSize); - } - - uint32_t WorkGroupNumDim = getKernelWorkGroupNumDim(Func); - if (WorkGroupNumDim) { - MetadataNames.push_back(Func.getName().str() + "@work_group_num_dim"); - PropSet.add(PropSetRegTy::SYCL_PROGRAM_METADATA, MetadataNames.back(), - WorkGroupNumDim); - } - } - - // Add global_id_mapping information with mapping between device-global - // unique identifiers and the variable's name in the IR. - for (auto &GV : M.globals()) { - if (!isDeviceGlobalVariable(GV)) - continue; - - StringRef GlobalID = getGlobalVariableUniqueId(GV); - MetadataNames.push_back(GlobalID.str() + "@global_id_mapping"); - PropSet.add(PropSetRegTy::SYCL_PROGRAM_METADATA, MetadataNames.back(), - GV.getName()); - } - } - if (MD.isESIMD()) { - PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "isEsimdImage", true); - } - { - StringRef RegAllocModeAttr = "sycl-register-alloc-mode"; - uint32_t RegAllocModeVal; - - bool HasRegAllocMode = llvm::any_of(MD.entries(), [&](const Function *F) { - if (!F->hasFnAttribute(RegAllocModeAttr)) - return false; - const auto &Attr = F->getFnAttribute(RegAllocModeAttr); - RegAllocModeVal = getAttributeAsInteger(Attr); - return true; - }); - if (HasRegAllocMode) { - PropSet.add(PropSetRegTy::SYCL_MISC_PROP, RegAllocModeAttr, - RegAllocModeVal); - } - } - - { - StringRef GRFSizeAttr = "sycl-grf-size"; - uint32_t GRFSizeVal; - - bool HasGRFSize = llvm::any_of(MD.entries(), [&](const Function *F) { - if (!F->hasFnAttribute(GRFSizeAttr)) - return false; - const auto &Attr = F->getFnAttribute(GRFSizeAttr); - GRFSizeVal = getAttributeAsInteger(Attr); - return true; - }); - if (HasGRFSize) { - PropSet.add(PropSetRegTy::SYCL_MISC_PROP, GRFSizeAttr, GRFSizeVal); - } - } - - // FIXME: Remove 'if' below when possible - // GPU backend has a problem with accepting optimization level options in form - // described by Level Zero specification (-ze-opt-level=1) when 'invoke_simd' - // functionality is involved. JIT compilation results in the following error: - // error: VLD: Failed to compile SPIR-V with following error: - // invalid api option: -ze-opt-level=O1 - // -11 (PI_ERROR_BUILD_PROGRAM_FAILURE) - // 'if' below essentially preserves the behavior (presumably mistakenly) - // implemented in intel/llvm#8763: ignore 'optLevel' property for images which - // were produced my merge after ESIMD split - if (MD.getEntryPointGroup().Props.HasESIMD != - module_split::SyclEsimdSplitStatus::SYCL_AND_ESIMD) { - // Handle sycl-optlevel property - int OptLevel = -1; - for (const Function *F : MD.entries()) { - if (!F->hasFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL)) - continue; - - // getAsInteger returns true on error - if (!F->getFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL) - .getValueAsString() - .getAsInteger(10, OptLevel)) { - // It is expected that device-code split has separated kernels with - // different values of sycl-optlevel attribute. Therefore, it is enough - // to only look at the first function with such attribute to compute - // the property for the whole device image. - break; - } - } - - if (OptLevel != -1) - PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "optLevel", OptLevel); - } - { - std::vector FuncNames = getKernelNamesUsingAssert(M); - for (const StringRef &FName : FuncNames) - PropSet.add(PropSetRegTy::SYCL_ASSERT_USED, FName, true); - } - - { - if (isModuleUsingAsan(M)) - PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "asanUsed", true); - } - - if (GlobProps.EmitDeviceGlobalPropSet) { - // Extract device global maps per module - auto DevGlobalPropertyMap = collectDeviceGlobalProperties(M); - if (!DevGlobalPropertyMap.empty()) - PropSet.add(PropSetRegTy::SYCL_DEVICE_GLOBALS, DevGlobalPropertyMap); - } - - auto HostPipePropertyMap = collectHostPipeProperties(M); - if (!HostPipePropertyMap.empty()) { - PropSet.add(PropSetRegTy::SYCL_HOST_PIPES, HostPipePropertyMap); - } - - if (MD.isSpecConstantDefault()) - PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "specConstsReplacedWithDefault", - 1); + const auto &PropSet = computeModuleProperties( + MD.getModule(), MD.entries(), GlobProps, MD.Props.SpecConstsMet, + MD.isSpecConstantDefault()); std::error_code EC; std::string SCFile = makeResultFileName(".prop", I, Suff); @@ -646,24 +320,9 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, } // Saves specified collection of symbols to a file. -std::string saveModuleSymbolTable(const module_split::EntryPointSet &Es, int I, +std::string saveModuleSymbolTable(const module_split::ModuleDesc &MD, int I, StringRef Suffix) { -#ifndef NDEBUG - if (DebugPostLink > 0) { - llvm::errs() << "ENTRY POINTS saving Sym table {\n"; - for (const auto *F : Es) { - llvm::errs() << " " << F->getName() << "\n"; - } - llvm::errs() << "}\n"; - } -#endif // NDEBUG - // Concatenate names of the input entry points with "\n". - std::string SymT; - - for (const auto *F : Es) { - SymT = (Twine(SymT) + Twine(F->getName()) + Twine("\n")).str(); - } - // Save to file. + auto SymT = computeModuleSymbolTable(MD.getModule(), MD.entries()); std::string OutFileName = makeResultFileName(".sym", I, Suffix); writeToFile(OutFileName, SymT); return OutFileName; @@ -758,11 +417,12 @@ IrPropSymFilenameTriple saveModule(module_split::ModuleDesc &MD, int I, GlobalBinImageProps Props = {EmitKernelParamInfo, EmitProgramMetadata, EmitExportedSymbols, EmitImportedSymbols, DeviceGlobals}; - Res.Prop = saveModuleProperties(MD, Props, I, Suffix); + if (DoPropGen) + Res.Prop = saveModuleProperties(MD, Props, I, Suffix); if (DoSymGen) { // save the names of the entry points - the symbol table - Res.Sym = saveModuleSymbolTable(MD.entries(), I, Suffix); + Res.Sym = saveModuleSymbolTable(MD, I, Suffix); } return Res; } @@ -1061,11 +721,14 @@ std::vector> processInputModule(std::unique_ptr M) { // Construct the resulting table which will accumulate all the outputs. SmallVector ColumnTitles{ - StringRef(COL_CODE), StringRef(COL_PROPS)}; + StringRef(COL_CODE)}; - if (DoSymGen) { + if (DoPropGen) + ColumnTitles.push_back(COL_PROPS); + + if (DoSymGen) ColumnTitles.push_back(COL_SYM); - } + Expected> TableE = util::SimpleTable::create(ColumnTitles); CHECK_AND_EXIT(TableE.takeError()); @@ -1255,8 +918,8 @@ int main(int argc, char **argv) { "Normally, the tool generates a number of files and \"file table\"\n" "file listing all generated files in a table manner. For example, if\n" "the input file 'example.bc' contains two kernels, then the command\n" - " $ sycl-post-link --split=kernel --symbols --spec-const=native \\\n" - " -o example.table example.bc\n" + " $ sycl-post-link --properties --split=kernel --symbols \\\n" + " --spec-const=native -o example.table example.bc\n" "will produce 'example.table' file with the following content:\n" " [Code|Properties|Symbols]\n" " example_0.bc|example_0.prop|example_0.sym\n" @@ -1282,7 +945,7 @@ int main(int argc, char **argv) { bool DoGenerateDeviceImageWithDefaulValues = GenerateDeviceImageWithDefaultSpecConsts.getNumOccurrences() > 0; - if (!DoSplit && !DoSpecConst && !DoSymGen && !DoParamInfo && + if (!DoSplit && !DoSpecConst && !DoSymGen && !DoPropGen && !DoParamInfo && !DoProgMetadata && !DoSplitEsimd && !DoExportedSyms && !DoImportedSyms && !DoDeviceGlobals && !DoLowerEsimd) { errs() << "no actions specified; try --help for usage info\n"; @@ -1303,6 +966,11 @@ int main(int argc, char **argv) { << IROutputOnly.ArgStr << "\n"; return 1; } + if (IROutputOnly && DoPropGen) { + errs() << "error: -" << DoPropGen.ArgStr << " can't be used with -" + << IROutputOnly.ArgStr << "\n"; + return 1; + } if (IROutputOnly && DoParamInfo) { errs() << "error: -" << EmitKernelParamInfo.ArgStr << " can't be used with" << " -" << IROutputOnly.ArgStr << "\n"; diff --git a/sycl/test/basic_tests/SYCL-2020-spec-const-ids-order.cpp b/sycl/test/basic_tests/SYCL-2020-spec-const-ids-order.cpp index 0aefe26bd826c..2e3bac17a03c2 100644 --- a/sycl/test/basic_tests/SYCL-2020-spec-const-ids-order.cpp +++ b/sycl/test/basic_tests/SYCL-2020-spec-const-ids-order.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-device-only -c -o %t.bc %s -// RUN: sycl-post-link %t.bc -spec-const=emulation -S -o %t-split1.txt +// RUN: sycl-post-link -properties %t.bc -spec-const=emulation -S -o %t-split1.txt // RUN: cat %t-split1_0.ll | FileCheck %s -check-prefixes=CHECK-IR // RUN: cat %t-split1_0.prop | FileCheck %s -check-prefixes=CHECK-PROP // diff --git a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp index b9d29507a8eac..136043dba5fd6 100644 --- a/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp +++ b/sycl/test/basic_tests/SYCL-2020-spec-constants.cpp @@ -1,7 +1,7 @@ // RUN: %clangxx -fsycl -fsycl-device-only -c -o %t.bc %s -// RUN: %if asserts %{sycl-post-link -debug-only=SpecConst %t.bc -spec-const=native -o %t-split1.txt 2>&1 | FileCheck %s -check-prefixes=CHECK-LOG %} %else %{sycl-post-link %t.bc -spec-const=native -o %t-split1.txt 2>&1 %} +// RUN: %if asserts %{sycl-post-link -properties -debug-only=SpecConst %t.bc -spec-const=native -o %t-split1.txt 2>&1 | FileCheck %s -check-prefixes=CHECK-LOG %} %else %{sycl-post-link -properties %t.bc -spec-const=native -o %t-split1.txt 2>&1 %} // RUN: cat %t-split1_0.prop | FileCheck %s -check-prefixes=CHECK,CHECK-RT -// RUN: sycl-post-link %t.bc -spec-const=emulation -o %t-split2.txt +// RUN: sycl-post-link -properties %t.bc -spec-const=emulation -o %t-split2.txt // RUN: cat %t-split2_0.prop | FileCheck %s -check-prefixes=CHECK,CHECK-DEF // RUN: llvm-spirv -o %t-split1_0.spv -spirv-max-version=1.1 -spirv-ext=+all %t-split1_0.bc // RUN: llvm-spirv -o - --to-text %t-split1_0.spv | FileCheck %s -check-prefixes=CHECK-SPV diff --git a/sycl/test/basic_tests/sycl-kernel-save-user-names.cpp b/sycl/test/basic_tests/sycl-kernel-save-user-names.cpp index 4663aa2f55baf..5c3c0a771ecd8 100644 --- a/sycl/test/basic_tests/sycl-kernel-save-user-names.cpp +++ b/sycl/test/basic_tests/sycl-kernel-save-user-names.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-device-only -fno-discard-value-names -fno-sycl-early-optimizations -o %t.bc %s -// RUN: sycl-post-link %t.bc -spec-const=emulation -o %t.table +// RUN: sycl-post-link -properties %t.bc -spec-const=emulation -o %t.table // RUN: llvm-spirv -o %t.spv -spirv-max-version=1.3 -spirv-ext=+all %t.bc // RUN: llvm-spirv -o %t.rev.bc -r %t.spv // RUN: llvm-dis %t.rev.bc -o=- | FileCheck %s diff --git a/sycl/test/check_device_code/esimd/NBarrierAttr.cpp b/sycl/test/check_device_code/esimd/NBarrierAttr.cpp index 7f9c7a6114e33..174b870c9b965 100644 --- a/sycl/test/check_device_code/esimd/NBarrierAttr.cpp +++ b/sycl/test/check_device_code/esimd/NBarrierAttr.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -O2 -fsycl -fsycl-device-only -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -lower-esimd -O2 -S %t -o %t.table +// RUN: sycl-post-link -properties -lower-esimd -O2 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll // Checks NBarrier ESIMD intrinsic translation. diff --git a/sycl/test/check_device_code/esimd/dpas.cpp b/sycl/test/check_device_code/esimd/dpas.cpp index 0892200641bb6..7e2c2682d3988 100644 --- a/sycl/test/check_device_code/esimd/dpas.cpp +++ b/sycl/test/check_device_code/esimd/dpas.cpp @@ -1,6 +1,6 @@ // RUN: %clangxx -fsycl -c -Xclang -emit-llvm %s -o %t // RUN: %clangxx -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll #include diff --git a/sycl/test/check_device_code/esimd/fp16_converts.cpp b/sycl/test/check_device_code/esimd/fp16_converts.cpp index 83b1b86f96f67..61ed76a211837 100644 --- a/sycl/test/check_device_code/esimd/fp16_converts.cpp +++ b/sycl/test/check_device_code/esimd/fp16_converts.cpp @@ -5,7 +5,7 @@ // Checks that lowerESIMD pass builds proper vc-intrinsics // RUN: %clangxx -O2 -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll #include diff --git a/sycl/test/check_device_code/esimd/genx_func_attr.cpp b/sycl/test/check_device_code/esimd/genx_func_attr.cpp index 00fb5d7171431..f92c5f40c36a3 100644 --- a/sycl/test/check_device_code/esimd/genx_func_attr.cpp +++ b/sycl/test/check_device_code/esimd/genx_func_attr.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -O2 -fsycl -fsycl-device-only -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -lower-esimd -O2 -S %t -o %t.table +// RUN: sycl-post-link -properties -lower-esimd -O2 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll // Checks ESIMD intrinsic translation. diff --git a/sycl/test/check_device_code/esimd/glob.cpp b/sycl/test/check_device_code/esimd/glob.cpp index b61a6e59830ec..2f005a7c71340 100644 --- a/sycl/test/check_device_code/esimd/glob.cpp +++ b/sycl/test/check_device_code/esimd/glob.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -O2 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll // This test checks that globals with register attribute are allowed in ESIMD diff --git a/sycl/test/check_device_code/esimd/intrins_trans.cpp b/sycl/test/check_device_code/esimd/intrins_trans.cpp index f42a49bfdd87e..57543e6025702 100644 --- a/sycl/test/check_device_code/esimd/intrins_trans.cpp +++ b/sycl/test/check_device_code/esimd/intrins_trans.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx %clang_O0 -fsycl -fsycl-device-only -fno-sycl-esimd-force-stateless-mem -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx %clang_O0 -fsycl -fsycl-device-only -fsycl-esimd-force-stateless-mem -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=true -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=true -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD intrinsic translation with opaque pointers. diff --git a/sycl/test/check_device_code/esimd/lower-external-funcs.cpp b/sycl/test/check_device_code/esimd/lower-external-funcs.cpp index 580d1bff5b121..7a8e0a492c5fa 100644 --- a/sycl/test/check_device_code/esimd/lower-external-funcs.cpp +++ b/sycl/test/check_device_code/esimd/lower-external-funcs.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm -x c++ %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -O2 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll // This test checks that unreferenced SYCL_EXTERNAL functions are not dropped diff --git a/sycl/test/check_device_code/esimd/lsc.cpp b/sycl/test/check_device_code/esimd/lsc.cpp index 131952b2ee03a..e1e157225fe77 100644 --- a/sycl/test/check_device_code/esimd/lsc.cpp +++ b/sycl/test/check_device_code/esimd/lsc.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx -O0 -fsycl -fno-sycl-esimd-force-stateless-mem -fsycl-device-only -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx -O0 -fsycl -fsycl-esimd-force-stateless-mem -fsycl-device-only -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD intrinsic translation. diff --git a/sycl/test/check_device_code/esimd/memory_properties_atomic_update.cpp b/sycl/test/check_device_code/esimd/memory_properties_atomic_update.cpp index 403448e3677f0..d960dabc0b515 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_atomic_update.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_atomic_update.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fno-sycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD memory functions accepting compile time properties for diff --git a/sycl/test/check_device_code/esimd/memory_properties_copytocopyfrom.cpp b/sycl/test/check_device_code/esimd/memory_properties_copytocopyfrom.cpp index 6f9ab411b616d..05d953357bade 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_copytocopyfrom.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_copytocopyfrom.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fno-sycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD copy_to and copy_from functions accepting compile time diff --git a/sycl/test/check_device_code/esimd/memory_properties_gather.cpp b/sycl/test/check_device_code/esimd/memory_properties_gather.cpp index 5828ac743c24e..03b08b9e8e8b8 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_gather.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_gather.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fno-sycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD memory functions accepting compile time properties for gather diff --git a/sycl/test/check_device_code/esimd/memory_properties_load_store.cpp b/sycl/test/check_device_code/esimd/memory_properties_load_store.cpp index c1e465536268f..5dacc5e5c99b2 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_load_store.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_load_store.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fno-sycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD memory functions accepting compile time properties for diff --git a/sycl/test/check_device_code/esimd/memory_properties_prefetch_2d.cpp b/sycl/test/check_device_code/esimd/memory_properties_prefetch_2d.cpp index c9c5f33854057..03945af5a97b5 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_prefetch_2d.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_prefetch_2d.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fno-sycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD memory functions accepting compile time properties for prefetch diff --git a/sycl/test/check_device_code/esimd/memory_properties_scatter.cpp b/sycl/test/check_device_code/esimd/memory_properties_scatter.cpp index 2bf1b0f4019c7..e97532d528549 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_scatter.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_scatter.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fno-sycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-esimd-force-stateless-mem -D__ESIMD_GATHER_SCATTER_LLVM_IR -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD memory functions accepting compile time properties for scatter diff --git a/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp b/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp index 39ecacb2fa663..f12404afc81fe 100644 --- a/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp +++ b/sycl/test/check_device_code/esimd/slm_init_specconst_size.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -O2 -fsycl -fsycl-device-only -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -O2 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll // Checks that we set 0 as VCSLMSize when slm_init is used with // non-constant operand, like with specialization constant. diff --git a/sycl/test/check_device_code/native_cpu/offload-wrapper.cpp b/sycl/test/check_device_code/native_cpu/offload-wrapper.cpp index 50bab6804255e..98d4ce88e623f 100644 --- a/sycl/test/check_device_code/native_cpu/offload-wrapper.cpp +++ b/sycl/test/check_device_code/native_cpu/offload-wrapper.cpp @@ -1,7 +1,7 @@ // This test checks the output for the clang-offload-wrapper for the Native CPU // target: // RUN: %clangxx -fsycl-device-only -fsycl-targets=native_cpu %s -o %t.bc -// RUN: sycl-post-link -emit-param-info -symbols -emit-exported-symbols -O2 -spec-const=native -device-globals -o %t.table %t.bc +// RUN: sycl-post-link -properties -emit-param-info -symbols -emit-exported-symbols -O2 -spec-const=native -device-globals -o %t.table %t.bc // RUN: clang-offload-wrapper -o=%t_wrap.bc -host=x86_64-unknown-linux-gnu -target=native_cpu -kind=sycl -batch %t.table // RUN: llvm-dis %t_wrap.bc -o - | FileCheck %s diff --git a/sycl/test/esimd/sycl_half_basic_ops.cpp b/sycl/test/esimd/sycl_half_basic_ops.cpp index ae357936a6f54..224286c95213a 100644 --- a/sycl/test/esimd/sycl_half_basic_ops.cpp +++ b/sycl/test/esimd/sycl_half_basic_ops.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-device-only -S %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -S %t -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll // The test checks that there are no unexpected extra conversions or intrinsic diff --git a/sycl/test/esimd/sycl_half_math_ops.cpp b/sycl/test/esimd/sycl_half_math_ops.cpp index 38e645c703ebc..10a81279b535c 100644 --- a/sycl/test/esimd/sycl_half_math_ops.cpp +++ b/sycl/test/esimd/sycl_half_math_ops.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-device-only -S %s -o %t.ll -// RUN: sycl-post-link -split-esimd -lower-esimd -S %t.ll -o %t.table +// RUN: sycl-post-link -properties -split-esimd -lower-esimd -S %t.ll -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll // The test checks that there are no unexpected extra conversions or intrinsic diff --git a/sycl/test/extensions/private_alloca.cpp b/sycl/test/extensions/private_alloca.cpp index 3e3d3f91ddfc5..22f834dc5217c 100644 --- a/sycl/test/extensions/private_alloca.cpp +++ b/sycl/test/extensions/private_alloca.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-device-only -c -o %t.bc %s -// RUN: %if asserts %{sycl-post-link -debug-only=SpecConst %t.bc -spec-const=native -o %t.txt 2>&1 | FileCheck %s -check-prefixes=CHECK-LOG %} %else %{sycl-post-link %t.bc -spec-const=native -o %t.txt 2>&1 %} +// RUN: %if asserts %{sycl-post-link -properties -debug-only=SpecConst %t.bc -spec-const=native -o %t.txt 2>&1 | FileCheck %s -check-prefixes=CHECK-LOG %} %else %{sycl-post-link %t.bc -properties -spec-const=native -o %t.txt 2>&1 %} // RUN: cat %t_0.prop | FileCheck %s -check-prefixes=CHECK,CHECK-RT // RUN: llvm-spirv -o %t_0.spv -spirv-max-version=1.1 -spirv-ext=+all %t_0.bc // RUN: llvm-spirv -o - --to-text %t_0.spv | FileCheck %s -check-prefixes=CHECK-SPV diff --git a/sycl/test/optional_kernel_features/fp-accuracy.c b/sycl/test/optional_kernel_features/fp-accuracy.c index 80acc2baa893f..f1107acd59d72 100644 --- a/sycl/test/optional_kernel_features/fp-accuracy.c +++ b/sycl/test/optional_kernel_features/fp-accuracy.c @@ -3,7 +3,7 @@ // 1. Accuracy is specified for particular math functions. // RUN: %clangxx %s -o %test_func.bc -ffp-accuracy=high:sin,sqrt -ffp-accuracy=medium:cos -ffp-accuracy=low:tan -ffp-accuracy=cuda:exp,acos -ffp-accuracy=sycl:log,asin -fno-math-errno -fsycl -fsycl-device-only -// RUN: sycl-post-link -split=auto -symbols %test_func.bc -o %test_func.table +// RUN: sycl-post-link -properties -split=auto -symbols %test_func.bc -o %test_func.table // RUN: FileCheck %s -input-file=%test_func.table --check-prefixes CHECK-FUNC-TABLE // RUN: FileCheck %s -input-file=%test_func_0.sym --check-prefixes CHECK-FUNC-M0-SYMS // RUN: FileCheck %s -input-file=%test_func_1.sym --check-prefixes CHECK-FUNC-M1-SYMS @@ -14,14 +14,14 @@ // 2. Accuracy is specified for TU. // RUN: %clangxx %s -o %test_tu.bc -ffp-accuracy=high -fno-math-errno -fsycl -fsycl-device-only -// RUN: sycl-post-link -split=auto -symbols %test_tu.bc -o %test_tu.table +// RUN: sycl-post-link -properties -split=auto -symbols %test_tu.bc -o %test_tu.table // RUN: FileCheck %s -input-file=%test_tu.table --check-prefixes CHECK-TU-TABLE // RUN: FileCheck %s -input-file=%test_tu_0.sym --check-prefixes CHECK-TU-M0-SYMS // RUN: FileCheck %s -input-file=%test_tu_1.sym --check-prefixes CHECK-TU-M1-SYMS // 3. Mixed case. // RUN: %clangxx %s -o %test_mix.bc -ffp-accuracy=medium -ffp-accuracy=high:sin,sqrt -ffp-accuracy=medium:cos -ffp-accuracy=cuda:exp -ffp-accuracy=sycl:log -fno-math-errno -fsycl -fsycl-device-only -// RUN: sycl-post-link -split=auto -symbols %test_mix.bc -o %test_mix.table +// RUN: sycl-post-link -properties -split=auto -symbols %test_mix.bc -o %test_mix.table // RUN: FileCheck %s -input-file=%test_mix.table --check-prefixes CHECK-MIX-TABLE // RUN: FileCheck %s -input-file=%test_mix_0.sym --check-prefixes CHECK-MIX-M0-SYMS // RUN: FileCheck %s -input-file=%test_mix_1.sym --check-prefixes CHECK-MIX-M1-SYMS diff --git a/sycl/test/optional_kernel_features/fp-accuracy.cpp b/sycl/test/optional_kernel_features/fp-accuracy.cpp index 80acc2baa893f..f1107acd59d72 100644 --- a/sycl/test/optional_kernel_features/fp-accuracy.cpp +++ b/sycl/test/optional_kernel_features/fp-accuracy.cpp @@ -3,7 +3,7 @@ // 1. Accuracy is specified for particular math functions. // RUN: %clangxx %s -o %test_func.bc -ffp-accuracy=high:sin,sqrt -ffp-accuracy=medium:cos -ffp-accuracy=low:tan -ffp-accuracy=cuda:exp,acos -ffp-accuracy=sycl:log,asin -fno-math-errno -fsycl -fsycl-device-only -// RUN: sycl-post-link -split=auto -symbols %test_func.bc -o %test_func.table +// RUN: sycl-post-link -properties -split=auto -symbols %test_func.bc -o %test_func.table // RUN: FileCheck %s -input-file=%test_func.table --check-prefixes CHECK-FUNC-TABLE // RUN: FileCheck %s -input-file=%test_func_0.sym --check-prefixes CHECK-FUNC-M0-SYMS // RUN: FileCheck %s -input-file=%test_func_1.sym --check-prefixes CHECK-FUNC-M1-SYMS @@ -14,14 +14,14 @@ // 2. Accuracy is specified for TU. // RUN: %clangxx %s -o %test_tu.bc -ffp-accuracy=high -fno-math-errno -fsycl -fsycl-device-only -// RUN: sycl-post-link -split=auto -symbols %test_tu.bc -o %test_tu.table +// RUN: sycl-post-link -properties -split=auto -symbols %test_tu.bc -o %test_tu.table // RUN: FileCheck %s -input-file=%test_tu.table --check-prefixes CHECK-TU-TABLE // RUN: FileCheck %s -input-file=%test_tu_0.sym --check-prefixes CHECK-TU-M0-SYMS // RUN: FileCheck %s -input-file=%test_tu_1.sym --check-prefixes CHECK-TU-M1-SYMS // 3. Mixed case. // RUN: %clangxx %s -o %test_mix.bc -ffp-accuracy=medium -ffp-accuracy=high:sin,sqrt -ffp-accuracy=medium:cos -ffp-accuracy=cuda:exp -ffp-accuracy=sycl:log -fno-math-errno -fsycl -fsycl-device-only -// RUN: sycl-post-link -split=auto -symbols %test_mix.bc -o %test_mix.table +// RUN: sycl-post-link -properties -split=auto -symbols %test_mix.bc -o %test_mix.table // RUN: FileCheck %s -input-file=%test_mix.table --check-prefixes CHECK-MIX-TABLE // RUN: FileCheck %s -input-file=%test_mix_0.sym --check-prefixes CHECK-MIX-M0-SYMS // RUN: FileCheck %s -input-file=%test_mix_1.sym --check-prefixes CHECK-MIX-M1-SYMS