[SYCL][NewOffloadModel] add support of -no-sycl-rdc in clang-linker-wrapper (#22118)

Author: maksimsab

clang/test/Driver/clang-linker-wrapper.cpp

@@ -372,3 +372,18 @@
 // RUN:                      %t.wrong_spirv_subarch.o -o %t.out 2>&1 --linker-path="/usr/bin/ld" | FileCheck %s --check-prefix=CHECK-ERROR-WITH-WRONG-SPIRV-SUBARCH
 
 // CHECK-ERROR-WITH-WRONG-SPIRV-SUBARCH: SYCL device compilation is not supported for the target: spirv64v1.0
+
+// Check -no-sycl-rdc
+// RUN: %clang %s -fsycl -fsycl-targets=spir64-unknown-unknown -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t1.o
+// RUN: %clang %s -fsycl -fsycl-targets=spir64-unknown-unknown -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t2.o
+//
+// RUN: clang-linker-wrapper "--host-triple=x86_64-unknown-linux-gnu" "--linker-path=/usr/bin/ld" "--" "-o" "a.out"  %t1.o t2.o --dry-run 2>&1 | FileCheck -check-prefix=CHECK-RDC %s
+// CHECK-RDC: spirv-to-ir-wrapper{{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global
+// CHECK-RDC-NEXT: llvm-link{{.*}} --suppress-warnings [[FIRSTLLVMLINKIN]].bc -o [[LLVMLINKOUT:.*]].bc
+// CHECK-RDC-NEXT: sycl-post-link{{.*}} -o [[SYCLPOSTLINKOUT:.*]].table [[LLVMLINKOUT]].bc
+
+// RUN: clang-linker-wrapper -no-sycl-rdc "--host-triple=x86_64-unknown-linux-gnu" "--linker-path=/usr/bin/ld" "--" "-o" "a.out"  %t1.o %t2.o --dry-run 2>&1 | FileCheck -check-prefix=CHECK-NO-RDC %s
+// CHECK-NO-RDC-NOT: llvm-link
+// CHECK-NO-RDC: llvm-spirv{{.*}} -o [[FIRST_SPIRV_OUTPUT:.*]].spv
+// CHECK-NO-RDC: llvm-spirv{{.*}} -o [[SECOND_SPIRV_OUTPUT:.*]].spv
+// CHECK-NO-RDC: offload-wrapper:{{.*}} input: [[FIRST_SPIRV_OUTPUT]].spv, {{.*}} input: [[SECOND_SPIRV_OUTPUT]].spv

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

@@ -1981,13 +1981,38 @@ Expected<std::vector<module_split::SplitModule>> runSYCLOffloadingPipeline(
     ArrayRef<StringRef> InputModules, const ArgList &LinkerArgs,
     const std::pair<std::string, std::string> &CompileLinkOptions,
     function_ref<void(StringRef)> WrappedOutputCallback) {
-  Expected<StringRef> LinkedModuleOrErr =
-      sycl::linkDevice(InputModules, LinkerArgs);
-  if (!LinkedModuleOrErr)
-    return LinkedModuleOrErr.takeError();
+  // Note: pipeline can skip linking due to -fno-sycl-rdc option.
+  // In that case, we apply sycl processing to several modules.
+  std::vector<StringRef> Modules;
+  if (LinkerArgs.hasArg(OPT_no_sycl_rdc)) {
+    // No need to perform any linking.
+    Modules = std::vector<StringRef>(InputModules.begin(), InputModules.end());
+  } else {
+    Expected<StringRef> OutputOrErr =
+        sycl::linkDevice(InputModules, LinkerArgs);
+    if (!OutputOrErr)
+      return OutputOrErr.takeError();
+
+    Modules.push_back(*OutputOrErr);
+  }
+
+  std::vector<module_split::SplitModule> OutputModules;
+  // TODO: parallelization of the loop below would lead to big performance
+  // improvement.
+  for (StringRef Module : Modules) {
+    // Note: sycl-post-link can produce more modules than incoming due to module
+    // split.
+    Expected<std::vector<module_split::SplitModule>> ModulesOrErr =
+        postLinkProcessModule(Module, LinkerArgs, CompileLinkOptions,
+                              WrappedOutputCallback);
+    if (!ModulesOrErr)
+      return ModulesOrErr.takeError();
+
+    for (module_split::SplitModule &M : *ModulesOrErr)
+      OutputModules.push_back(std::move(M));
+  }
 
-  return postLinkProcessModule(*LinkedModuleOrErr, LinkerArgs,
-                               CompileLinkOptions, WrappedOutputCallback);
+  return OutputModules;
 }
 
 } // namespace sycl

clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td

@@ -213,6 +213,10 @@ Flags<[WrapperOnlyOption]>,  HelpText<"Link SYCL device code using thinLTO">;
 def sycl_embed_ir : Flag<["--", "-"], "sycl-embed-ir">,
 Flags<[WrapperOnlyOption]>,  HelpText<"Embed LLVM IR for runtime kernel fusion">;
 
+def no_sycl_rdc : Flag<["--", "-"], "no-sycl-rdc">,
+  Flags<[WrapperOnlyOption]>,
+  HelpText<"Disable RDC mode (Relocatable device code) for SYCL linking">;
+
 def sycl_device_link
     : Flag<["--", "-"], "sycl-device-link">,
       Flags<[WrapperOnlyOption]>,