diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 60b04903f25c..4616c9e71714 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -10702,6 +10702,7 @@ getTripleBasedSYCLPostLinkOpts(const ToolChain &TC, const JobAction &JA, // add options unconditionally addArgs(PostLinkArgs, TCArgs, {"-symbols"}); addArgs(PostLinkArgs, TCArgs, {"-emit-exported-symbols"}); + addArgs(PostLinkArgs, TCArgs, {"-emit-imported-symbols"}); if (SplitEsimd) addArgs(PostLinkArgs, TCArgs, {"-split-esimd"}); addArgs(PostLinkArgs, TCArgs, {"-lower-esimd"}); diff --git a/clang/test/Driver/sycl-device-lib.cpp b/clang/test/Driver/sycl-device-lib.cpp index d478c022a7e5..df90b2987220 100644 --- a/clang/test/Driver/sycl-device-lib.cpp +++ b/clang/test/Driver/sycl-device-lib.cpp @@ -185,7 +185,7 @@ // RUN: | FileCheck %s -check-prefix=SYCL_LLVM_LINK_NO_DEVICE_LIB // SYCL_LLVM_LINK_NO_DEVICE_LIB: clang{{.*}} "-cc1" {{.*}} "-fsycl-is-device" // SYCL_LLVM_LINK_NO_DEVICE_LIB-NOT: llvm-link{{.*}} "-only-needed" -// SYCL_LLVM_LINK_NO_DEVICE_LIB: sycl-post-link{{.*}} "-symbols" "-emit-exported-symbols"{{.*}} "-o" "{{.*}}.table" "{{.*}}.bc" +// SYCL_LLVM_LINK_NO_DEVICE_LIB: sycl-post-link{{.*}} "-symbols" "-emit-exported-symbols" "-emit-imported-symbols"{{.*}} "-o" "{{.*}}.table" "{{.*}}.bc" /// ########################################################################### /// test llvm-link behavior for special user input whose filename resembles SYCL device library diff --git a/clang/test/Driver/sycl-offload-new-driver.c b/clang/test/Driver/sycl-offload-new-driver.c index 8d672efe9264..74a3dce4c9f4 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=-post-link-opt -O2 -device-globals -spec-const=native -split=auto -emit-only-kernels-as-entry-points -symbols -emit-exported-symbols -lower-esimd" +// WRAPPER_OPTIONS_POSTLINK: clang-linker-wrapper{{.*}} "--sycl-post-link-options=-post-link-opt -O2 -device-globals -spec-const=native -split=auto -emit-only-kernels-as-entry-points -symbols -emit-exported-symbols -emit-imported-symbols -lower-esimd" // -fsycl-device-only behavior // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index 93e045256ed9..bbda6c548825 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -205,6 +205,7 @@ class PropertySetRegistry { static constexpr char SYCL_MISC_PROP[] = "SYCL/misc properties"; static constexpr char SYCL_ASSERT_USED[] = "SYCL/assert used"; static constexpr char SYCL_EXPORTED_SYMBOLS[] = "SYCL/exported symbols"; + static constexpr char SYCL_IMPORTED_SYMBOLS[] = "SYCL/imported symbols"; static constexpr char SYCL_DEVICE_GLOBALS[] = "SYCL/device globals"; static constexpr char SYCL_DEVICE_REQUIREMENTS[] = "SYCL/device requirements"; static constexpr char SYCL_HOST_PIPES[] = "SYCL/host pipes"; diff --git a/llvm/lib/Support/PropertySetIO.cpp b/llvm/lib/Support/PropertySetIO.cpp index 96593d4aa26b..f14f8cd5b16c 100644 --- a/llvm/lib/Support/PropertySetIO.cpp +++ b/llvm/lib/Support/PropertySetIO.cpp @@ -202,6 +202,7 @@ constexpr char PropertySetRegistry::SYCL_PROGRAM_METADATA[]; constexpr char PropertySetRegistry::SYCL_MISC_PROP[]; constexpr char PropertySetRegistry::SYCL_ASSERT_USED[]; constexpr char PropertySetRegistry::SYCL_EXPORTED_SYMBOLS[]; +constexpr char PropertySetRegistry::SYCL_IMPORTED_SYMBOLS[]; constexpr char PropertySetRegistry::SYCL_DEVICE_GLOBALS[]; constexpr char PropertySetRegistry::SYCL_DEVICE_REQUIREMENTS[]; constexpr char PropertySetRegistry::SYCL_HOST_PIPES[]; diff --git a/llvm/test/tools/sycl-post-link/emit_imported_symbols.ll b/llvm/test/tools/sycl-post-link/emit_imported_symbols.ll new file mode 100644 index 000000000000..ae824d293b9e --- /dev/null +++ b/llvm/test/tools/sycl-post-link/emit_imported_symbols.ll @@ -0,0 +1,113 @@ +; This test checks that the -emit-imported-symbols option generates a list of imported symbols +; Function names were chosen so that no function with a 'inside' in their function name is imported +; + +;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;; +; Test with -split=kernel +;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;; + +; RUN: sycl-post-link -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 +; RUN: FileCheck %s -input-file=%t_kernel_2.sym --check-prefixes CHECK-KERNEL-SYM-2 + +; RUN: FileCheck %s -input-file=%t_kernel_0.prop --check-prefixes CHECK-KERNEL-IMPORTED-SYM-0 +; RUN: FileCheck %s -input-file=%t_kernel_1.prop --check-prefixes CHECK-KERNEL-IMPORTED-SYM-1 +; RUN: FileCheck %s -input-file=%t_kernel_2.prop --check-prefixes CHECK-KERNEL-IMPORTED-SYM-2 + +; CHECK-KERNEL-SYM-0: middle +; CHECK-KERNEL-IMPORTED-SYM-0: [SYCL/imported symbols] +; CHECK-KERNEL-IMPORTED-SYM-0-NEXT: childD +; CHECK-KERNEL-IMPORTED-SYM-0-EMPTY: + +; CHECK-KERNEL-SYM-1: foo +; CHECK-KERNEL-IMPORTED-SYM-1: [SYCL/imported symbols] +; CHECK-KERNEL-IMPORTED-SYM-1-NEXT: childA +; CHECK-KERNEL-IMPORTED-SYM-1-NEXT: childC +; CHECK-KERNEL-IMPORTED-SYM-1-NEXT: childD +; CHECK-KERNEL-IMPORTED-SYM-1-EMPTY: + + +; CHECK-KERNEL-SYM-2: bar +; CHECK-KERNEL-IMPORTED-SYM-2: [SYCL/imported symbols] +; CHECK-KERNEL-IMPORTED-SYM-2-NEXT: childB +; CHECK-KERNEL-IMPORTED-SYM-2-NEXT: childC +; CHECK-KERNEL-IMPORTED-SYM-2-NEXT: childD +; CHECK-KERNEL-IMPORTED-SYM-2-NEXT: _Z7outsidev +; CHECK-KERNEL-IMPORTED-SYM-2-EMPTY: + +;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;; +; Test with -split=source +;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;; + +; RUN: sycl-post-link -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: 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 + +; CHECK-SOURCE-SYM-0-DAG: foo +; CHECK-SOURCE-SYM-0-DAG: bar +; CHECK-SOURCE-SYM-0-DAG: middle + +; CHECK-SOURCE-IMPORTED-SYM-0: [SYCL/imported symbols] +; CHECK-SOURCE-IMPORTED-SYM-0-NEXT: childA +; CHECK-SOURCE-IMPORTED-SYM-0-NEXT: childB +; CHECK-SOURCE-IMPORTED-SYM-0-NEXT: childC +; CHECK-SOURCE-IMPORTED-SYM-0-NEXT: childD +; CHECK-SOURCE-IMPORTED-SYM-0-NEXT: _Z7outsidev +; CHECK-SOURCE-IMPORTED-SYM-0-EMPTY: + +target triple = "spir64-unknown-unknown" + +@llvm.used = appending global [2 x ptr] [ptr @foo, ptr @bar], section "llvm.metadata" + +define weak_odr spir_kernel void @foo() #0 { + call void @childA() + call void @childC() + call void @middle() + ret void +} + +define weak_odr spir_kernel void @bar() #0 { + ;; Functions that are not SYCL External (i.e. they have no sycl-module-id) cannot be imported + call spir_func void @__itt_offload_wi_start_wrapper() + + call void @childB() + call void @childC() + call void @middle() + ;; LLVM intrinsics cannot be imported + %dummy = call i8 @llvm.bitreverse.i8(i8 0) + ;; Functions with a demangled name prefixed with a '__' are not imported + call void @_Z8__insidev() + call void @_Z7outsidev() + + ;; Functions that are not SYCL External (i.e. they have no sycl-module-id) cannot be imported + call spir_func void @__itt_offload_wi_finish_wrapper() + ret void +} + +define void @middle() #0 { + call void @childD() + ret void +} + +declare void @childA() #1 +declare void @childB() #1 +declare void @childC() #1 +declare void @childD() #1 + +declare void @_Z7outsidev() #1 +;; Verify unused functions are not imported +declare void @insideUnusedFunction() #1 +declare void @_Z8__insidev() #1 +declare i8 @llvm.bitreverse.i8(i8) + +declare spir_func void @__itt_offload_wi_start_wrapper() +declare spir_func void @__itt_offload_wi_finish_wrapper() + +attributes #0 = { "sycl-module-id"="a.cpp" } +attributes #1 = { "sycl-module-id"="external.cpp" } diff --git a/llvm/tools/sycl-post-link/CMakeLists.txt b/llvm/tools/sycl-post-link/CMakeLists.txt index cfb9b1a27560..aa98f4942edb 100644 --- a/llvm/tools/sycl-post-link/CMakeLists.txt +++ b/llvm/tools/sycl-post-link/CMakeLists.txt @@ -1,6 +1,7 @@ set(LLVM_LINK_COMPONENTS BitWriter Core + Demangle IPO IRPrinter IRReader diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 0d060e0c9aaf..9afa25c3a655 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -25,6 +25,7 @@ #include "llvm/Analysis/TargetLibraryInfo.h" #include "llvm/Analysis/TargetTransformInfo.h" #include "llvm/Bitcode/BitcodeWriterPass.h" +#include "llvm/Demangle/Demangle.h" #include "llvm/GenXIntrinsics/GenXSPIRVWriterAdaptor.h" #include "llvm/IR/Dominators.h" #include "llvm/IR/LLVMContext.h" @@ -228,6 +229,10 @@ cl::opt EmitExportedSymbols{"emit-exported-symbols", cl::desc("emit exported symbols"), cl::cat(PostLinkCat)}; +cl::opt EmitImportedSymbols{"emit-imported-symbols", + cl::desc("emit imported symbols"), + cl::cat(PostLinkCat)}; + cl::opt EmitOnlyKernelsAsEntryPoints{ "emit-only-kernels-as-entry-points", cl::desc("Consider only sycl_kernel functions as entry points for " @@ -250,6 +255,7 @@ struct GlobalBinImageProps { bool EmitKernelParamInfo; bool EmitProgramMetadata; bool EmitExportedSymbols; + bool EmitImportedSymbols; bool EmitDeviceGlobalPropSet; }; @@ -411,6 +417,25 @@ std::string saveModuleIR(Module &M, int I, StringRef Suff) { return OutFilename; } +bool isImportedFunction(const Function &F) { + if (!F.isDeclaration() || F.isIntrinsic() || + !llvm::sycl::utils::isSYCLExternalFunction(&F)) + return false; + + // 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"); + + bool ReturnValue = true; + if (char *NameStr = itaniumDemangle(F.getName())) { + StringRef DemangledName(NameStr); + if (DemangledName.starts_with("__")) + ReturnValue = false; + free(NameStr); + } + return ReturnValue; +} + std::string saveModuleProperties(module_split::ModuleDesc &MD, const GlobalBinImageProps &GlobProps, int I, StringRef Suff) { @@ -474,10 +499,21 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, // 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(), true); + 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 (isImportedFunction(F)) + 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; @@ -730,7 +766,8 @@ IrPropSymFilenameTriple saveModule(module_split::ModuleDesc &MD, int I, Res.Ir = saveModuleIR(MD.getModule(), I, Suffix); } GlobalBinImageProps Props = {EmitKernelParamInfo, EmitProgramMetadata, - EmitExportedSymbols, DeviceGlobals}; + EmitExportedSymbols, EmitImportedSymbols, + DeviceGlobals}; Res.Prop = saveModuleProperties(MD, Props, I, Suffix); if (DoSymGen) { @@ -1249,13 +1286,14 @@ int main(int argc, char **argv) { bool DoParamInfo = EmitKernelParamInfo.getNumOccurrences() > 0; bool DoProgMetadata = EmitProgramMetadata.getNumOccurrences() > 0; bool DoExportedSyms = EmitExportedSymbols.getNumOccurrences() > 0; + bool DoImportedSyms = EmitImportedSymbols.getNumOccurrences() > 0; bool DoDeviceGlobals = DeviceGlobals.getNumOccurrences() > 0; bool DoGenerateDeviceImageWithDefaulValues = GenerateDeviceImageWithDefaultSpecConsts.getNumOccurrences() > 0; if (!DoSplit && !DoSpecConst && !DoSymGen && !DoParamInfo && - !DoProgMetadata && !DoSplitEsimd && !DoExportedSyms && !DoDeviceGlobals && - !DoLowerEsimd) { + !DoProgMetadata && !DoSplitEsimd && !DoExportedSyms && !DoImportedSyms && + !DoDeviceGlobals && !DoLowerEsimd) { errs() << "no actions specified; try --help for usage info\n"; return 1; } @@ -1289,6 +1327,11 @@ int main(int argc, char **argv) { << " -" << IROutputOnly.ArgStr << "\n"; return 1; } + if (IROutputOnly && DoImportedSyms) { + errs() << "error: -" << EmitImportedSymbols.ArgStr << " can't be used with" + << " -" << IROutputOnly.ArgStr << "\n"; + return 1; + } if (IROutputOnly && DoGenerateDeviceImageWithDefaulValues) { errs() << "error: -" << GenerateDeviceImageWithDefaultSpecConsts.ArgStr << " can't be used with -" << IROutputOnly.ArgStr << "\n";