[SYCL] Generate imported symbol files in sycl-post-link

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

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

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 \

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

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[];

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

llvm/tools/sycl-post-link/CMakeLists.txt

@@ -1,6 +1,7 @@
 set(LLVM_LINK_COMPONENTS
   BitWriter
   Core
+  Demangle
   IPO
   IRPrinter
   IRReader

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<bool> EmitExportedSymbols{"emit-exported-symbols",
                                   cl::desc("emit exported symbols"),
                                   cl::cat(PostLinkCat)};
 
+cl::opt<bool> EmitImportedSymbols{"emit-imported-symbols",
+                                  cl::desc("emit imported symbols"),
+                                  cl::cat(PostLinkCat)};
+
 cl::opt<bool> 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<std::string, 4> 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";