[SYCL][New offload] Add support for -fsycl-link in new offload model (#19434)

This PR adds support for -fsycl-link in the compilation flow for the new
offloading model. Specifically, we pass a new option
"--sycl-device-link" to clang-linker-wrapper. We DO NOT call the final
host linker if the option is set.

Thanks

---------

Signed-off-by: Arvind Sudarsanam <arvind.sudarsanam@intel.com>

Author: asudarsa

clang/lib/Driver/ToolChains/Clang.cpp

@@ -11222,6 +11222,10 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
 
   // Add any SYCL offloading specific options to the clang-linker-wrapper
   if (C.hasOffloadToolChain<Action::OFK_SYCL>()) {
+
+    if (Args.hasArg(options::OPT_fsycl_link_EQ))
+      CmdArgs.push_back(Args.MakeArgString("--sycl-device-link"));
+
     // -sycl-device-libraries=<comma separated list> contains all of the SYCL
     // device specific libraries that are needed. This generic list will be
     // populated with device binaries for all target triples in the current

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

@@ -130,3 +130,15 @@
 // CHK-CMDS-DEVICE-LIB-DIR-NEXT: "{{.*}}llvm-link.exe" {{.*}} --suppress-warnings
 // CHK-CMDS-DEVICE-LIB-DIR-NEXT: "{{.*}}llvm-link.exe" -only-needed {{.*}} --suppress-warnings
 // CHK-CMDS-DEVICE-LIB-DIR-NEXT: "{{.*}}sycl-post-link.exe"{{.*}} --device-lib-dir={{.*}}/Inputs/SYCL/lib {{.*}} SYCL_POST_LINK_OPTIONS {{.*}}
+
+// Verify that host linker is not called when --sycl-device-link is passed to clang-linker-wrapper
+// RUN: clang-linker-wrapper --sycl-device-link -sycl-device-libraries=%t.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-pc-windows-msvc" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.exe" HOST_LIB_PATH HOST_STAT_LIB %t.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-DEVLINK-CMDS %s
+// CHK-DEVLINK-CMDS: "{{.*}}spirv-to-ir-wrapper.exe" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}llvm-link.exe" [[FIRSTLLVMLINKIN:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}llvm-link.exe" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}sycl-post-link.exe"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}llvm-spirv.exe"{{.*}} LLVM_SPIRV_OPTIONS -o {{.*}}
+// CHK-DEVLINK-CMDS-NEXT: offload-wrapper: input: {{.*}}, output: [[WRAPPEROUT:.*]].bc
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}clang.exe"{{.*}} -c -o {{.*}} [[WRAPPEROUT]].bc
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}copy"{{.*}} {{.*}} a.exe
+// CHK-DEVLINK-CMDS-NOT: "{{.*}}ld"

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

@@ -216,3 +216,15 @@
 // CHK-CMDS-NATIVE-CPU-NEXT:  offload-wrapper: input: [[OUT1]], output: [[OUT2:.*\.bc]]
 // CHK-CMDS-NATIVE-CPU-NEXT: "{{.*}}clang" --target=x86_64-unknown-linux-gnu -c -o [[OUT3:.*\.o]] [[OUT2]]
 // CHK-CMDS-NATIVE-CPU-NEXT: "{{.*}}ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[OUT1]] [[OUT3]] {{.*\.o}}
+
+// Verify that host linker is not called when --sycl-device-link is passed to clang-linker-wrapper
+// RUN: clang-linker-wrapper --sycl-device-link -sycl-device-libraries=%t.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-DEVLINK-CMDS %s
+// CHK-DEVLINK-CMDS: "{{.*}}spirv-to-ir-wrapper" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}llvm-link" [[FIRSTLLVMLINKIN]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}llvm-link" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}sycl-post-link"{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}llvm-spirv"{{.*}} LLVM_SPIRV_OPTIONS -o {{.*}}
+// CHK-DEVLINK-CMDS-NEXT: offload-wrapper: input: {{.*}}, output: [[WRAPPEROUT:.*]].bc
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}clang"{{.*}} -c -o [[CLANGOUT:.*]] [[WRAPPEROUT]].bc
+// CHK-DEVLINK-CMDS-NEXT: "{{.*}}cp"{{.*}} [[CLANGOUT]] a.out
+// CHK-DEVLINK-CMDS-NOT: "{{.*}}/ld"

clang/test/Driver/sycl-offload-new-driver.c

@@ -223,3 +223,8 @@
 // CHK-MACRO-GFX90A: clang{{.*}} "-fsycl-is-device"{{.*}} "-D__SYCL_TARGET_AMD_GPU_GFX90A__"{{.*}}
 // CHK-MACRO-GFX90C: clang{{.*}} "-fsycl-is-device"{{.*}} "-D__SYCL_TARGET_AMD_GPU_GFX90C__"{{.*}}
 
+/// Check that -sycl-device-link is passed to clang-linker-wrapper tool
+// RUN: %clangxx -fsycl -### --offload-new-driver \
+// RUN:          -fsycl-link %s 2>&1 \
+// RUN:  | FileCheck -check-prefix CHECK_SYCL_DEVICE_LINKING %s
+// CHECK_SYCL_DEVICE_LINKING: clang-linker-wrapper{{.*}} "--sycl-device-link"

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

@@ -1227,12 +1227,26 @@ packageSYCLBIN(SYCLBIN::BundleState State,
   return *OutFileOrErr;
 }
 
+Error copyFileToFinalExecutable(StringRef File, const ArgList &Args) {
+  if (Verbose || DryRun) {
+    llvm::Triple Triple(Args.getLastArgValue(OPT_host_triple_EQ,
+                                             sys::getDefaultTargetTriple()));
+    StringRef CopyCommand = Triple.isOSWindows() ? "copy" : "cp";
+    llvm::errs() << "\"" << CopyCommand << "\" " << File << " "
+                 << ExecutableName << "\n";
+  }
+  // TODO: check if copy can be replaced by rename.
+  if (std::error_code EC = sys::fs::copy_file(File, ExecutableName))
+    return createFileError(ExecutableName, EC);
+  return Error::success();
+}
+
 Error mergeSYCLBIN(ArrayRef<StringRef> Files, const ArgList &Args) {
   // Fast path for the general case where there's only one file. In this case we
   // do not need to parse it and can instead simply copy it.
   if (Files.size() == 1) {
-    if (std::error_code EC = sys::fs::copy_file(Files[0], ExecutableName))
-      return createFileError(ExecutableName, EC);
+    if (Error Err = copyFileToFinalExecutable(Files[0], Args))
+      reportError(std::move(Err));
     return Error::success();
   }
   // TODO: Merge SYCLBIN files here and write to ExecutableName output.
@@ -2227,7 +2241,7 @@ Expected<SmallVector<StringRef>> linkAndWrapDeviceFiles(
 
       // Store the offloading image for each linked output file.
       for (OffloadKind Kind = OFK_OpenMP; Kind != OFK_LAST;
-        Kind = static_cast<OffloadKind>((uint16_t)(Kind) << 1)) {
+           Kind = static_cast<OffloadKind>((uint16_t)(Kind) << 1)) {
         if ((ActiveOffloadKindMask & Kind) == 0)
           continue;
         llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> FileOrErr =
@@ -2789,6 +2803,14 @@ int main(int Argc, char **Argv) {
     if (OutputSYCLBIN) {
       if (Error Err = sycl::mergeSYCLBIN(*FilesOrErr, Args))
         reportError(std::move(Err));
+    } else if (Args.hasArg(OPT_sycl_device_link)) {
+      // Skip host linker if --sycl-device-link option is set.
+      // Just copy the output of device linking and wrapping action.
+      if (FilesOrErr->size() != 1)
+        reportError(
+            createStringError("Expect single output from the device linker."));
+      if (Error Err = sycl::copyFileToFinalExecutable((*FilesOrErr)[0], Args))
+        reportError(std::move(Err));
     } else {
       // Run the host linking job with the rendered arguments.
       if (Error Err = runLinker(*FilesOrErr, Args))

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

@@ -228,6 +228,11 @@ 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 sycl_device_link
+    : Flag<["--", "-"], "sycl-device-link">,
+      Flags<[WrapperOnlyOption]>,
+      HelpText<"Link SYCL device code only. Discard host code.">;
+
 def sycl_dump_device_code_EQ : Joined<["--", "-"], "sycl-dump-device-code=">,
   Flags<[WrapperOnlyOption]>,
   HelpText<"Directory to dump offloading images to.">;