[New offload model] Cleanup the way sycl-post-link options are generated.

Signed-off-by: Arvind Sudarsanam <[email protected]>

Author: asudarsa

clang/lib/Driver/ToolChains/Clang.cpp

@@ -10648,31 +10648,33 @@ static void getNonTripleBasedSYCLPostLinkOpts(const ToolChain &TC,
     addArgs(PostLinkArgs, TCArgs, {"-lower-esimd-force-stateless-mem=false"});
 }
 
-// Add any sycl-post-link options that rely on a specific Triple.
-static void
-getTripleBasedSYCLPostLinkOpts(const ToolChain &TC, const JobAction &JA,
-                               const llvm::opt::ArgList &TCArgs,
-                               llvm::Triple Triple, ArgStringList &PostLinkArgs,
-                               bool SpecConsts, types::ID OutputType) {
-  bool NewOffloadDriver = TC.getDriver().getUseNewOffloadingDriver();
-  // Note: Do not use Triple when NewOffloadDriver is 'true'.
-  if (!NewOffloadDriver && (OutputType == types::TY_LLVM_BC)) {
+// Add any sycl-post-link options that rely on a specific Triple in addition
+// to user supplied options. This function is invoked only for old offloading
+// model. For new offloading model, a slightly modified version of this
+// function is called inside clang-linker-wrapper.
+// NOTE: Any changes made here should be reflected in the similarly named
+// function in clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp.
+static void getTripleBasedSYCLPostLinkOpts(const ToolChain &TC,
+                                           const llvm::opt::ArgList &TCArgs,
+                                           ArgStringList &PostLinkArgs,
+                                           llvm::Triple Triple, bool SpecConsts,
+                                           types::ID OutputType) {
+  if (OutputType == types::TY_LLVM_BC) {
     // single file output requested - this means only perform necessary IR
     // transformations (like specialization constant intrinsic lowering) and
     // output LLVMIR
     addArgs(PostLinkArgs, TCArgs, {"-ir-output-only"});
   }
-  // specialization constants processing is mandatory
   if (SpecConsts)
     addArgs(PostLinkArgs, TCArgs, {"-spec-const=native"});
   else
     addArgs(PostLinkArgs, TCArgs, {"-spec-const=emulation"});
 
   // See if device code splitting is requested.  The logic here works along side
-  // the behavior in setOtherSYCLPostLinkOpts, where the option is added based
-  // on the user setting of-fsycl-device-code-split.
+  // the behavior in getNonTripleBasedSYCLPostLinkOpts, where the option is
+  // added based on the user setting of -fsycl-device-code-split.
   if (!TCArgs.hasArg(options::OPT_fsycl_device_code_split_EQ) &&
-      (NewOffloadDriver || !(Triple.getArchName() == "spir64_fpga")))
+      (!(Triple.getArchName() == "spir64_fpga")))
     addArgs(PostLinkArgs, TCArgs, {"-split=auto"});
 
   // On Intel targets we don't need non-kernel functions as entry points,
@@ -10683,18 +10685,17 @@ getTripleBasedSYCLPostLinkOpts(const ToolChain &TC, const JobAction &JA,
                        options::OPT_fsycl_remove_unused_external_funcs,
                        false) &&
        !isSYCLNativeCPU(TC)) &&
-      (NewOffloadDriver || (!Triple.isNVPTX() && !Triple.isAMDGPU())))
+      !Triple.isNVPTX() && !Triple.isAMDGPU())
     addArgs(PostLinkArgs, TCArgs, {"-emit-only-kernels-as-entry-points"});
 
-  if (!NewOffloadDriver && !Triple.isAMDGCN())
+  if (!Triple.isAMDGCN())
     addArgs(PostLinkArgs, TCArgs, {"-emit-param-info"});
   // Enable program metadata
-  if ((!NewOffloadDriver && (Triple.isNVPTX() || Triple.isAMDGCN())) ||
-      isSYCLNativeCPU(TC))
+  if (Triple.isNVPTX() || Triple.isAMDGCN() || isSYCLNativeCPU(TC))
     addArgs(PostLinkArgs, TCArgs, {"-emit-program-metadata"});
   if (OutputType != types::TY_LLVM_BC) {
     assert(OutputType == types::TY_Tempfiletable);
-    bool SplitEsimdByDefault = !NewOffloadDriver && Triple.isSPIROrSPIRV();
+    bool SplitEsimdByDefault = Triple.isSPIROrSPIRV();
     bool SplitEsimd = TCArgs.hasFlag(
         options::OPT_fsycl_device_code_split_esimd,
         options::OPT_fno_sycl_device_code_split_esimd, SplitEsimdByDefault);
@@ -10714,7 +10715,7 @@ getTripleBasedSYCLPostLinkOpts(const ToolChain &TC, const JobAction &JA,
   if (TCArgs.hasFlag(options::OPT_fsycl_add_default_spec_consts_image,
                      options::OPT_fno_sycl_add_default_spec_consts_image,
                      false) &&
-      (IsAOT || NewOffloadDriver))
+      IsAOT)
     addArgs(PostLinkArgs, TCArgs,
             {"-generate-device-image-default-spec-consts"});
 }
@@ -10737,8 +10738,16 @@ void SYCLPostLink::ConstructJob(Compilation &C, const JobAction &JA,
   ArgStringList CmdArgs;
 
   llvm::Triple T = getToolChain().getTriple();
+  const toolchains::SYCLToolChain &TC =
+      static_cast<const toolchains::SYCLToolChain &>(getToolChain());
+
+  // Handle -Xdevice-post-link
+  TC.TranslateTargetOpt(T, TCArgs, CmdArgs, options::OPT_Xdevice_post_link,
+                        options::OPT_Xdevice_post_link_EQ,
+                        JA.getOffloadingArch());
+
   getNonTripleBasedSYCLPostLinkOpts(getToolChain(), JA, TCArgs, CmdArgs);
-  getTripleBasedSYCLPostLinkOpts(getToolChain(), JA, TCArgs, T, CmdArgs,
+  getTripleBasedSYCLPostLinkOpts(getToolChain(), TCArgs, CmdArgs, T,
                                  SYCLPostLink->getRTSetsSpecConstants(),
                                  SYCLPostLink->getTrueType());
 
@@ -10751,14 +10760,6 @@ void SYCLPostLink::ConstructJob(Compilation &C, const JobAction &JA,
 
   addArgs(CmdArgs, TCArgs, {"-o", OutputArg});
 
-  const toolchains::SYCLToolChain &TC =
-      static_cast<const toolchains::SYCLToolChain &>(getToolChain());
-
-  // Handle -Xdevice-post-link
-  TC.TranslateTargetOpt(T, TCArgs, CmdArgs, options::OPT_Xdevice_post_link,
-                        options::OPT_Xdevice_post_link_EQ,
-                        JA.getOffloadingArch());
-
   // Add input file
   assert(Inputs.size() == 1 && Inputs.front().isFilename() &&
          "single input file expected");
@@ -11116,17 +11117,7 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
         appendOption(PostLinkOptString, A);
     }
     ArgStringList PostLinkArgs;
-    bool IsSYCLNativeCPU = driver::isSYCLNativeCPU(Args);
-    types::ID OutputType = TargetTriple.isSPIROrSPIRV() || IsSYCLNativeCPU
-                               ? types::TY_Tempfiletable
-                               : types::TY_LLVM_BC;
-    bool SpecConsts = TargetTriple.isSPIROrSPIRV();
     getNonTripleBasedSYCLPostLinkOpts(getToolChain(), JA, Args, PostLinkArgs);
-    // Some options like -spec-consts=* depend on target triple as well as some
-    // user options. So, these options are partly computed here and then
-    // updated inside the clang-linker-wrapper.
-    getTripleBasedSYCLPostLinkOpts(getToolChain(), JA, Args, TargetTriple,
-                                   PostLinkArgs, SpecConsts, OutputType);
     for (const auto &A : PostLinkArgs)
       appendOption(PostLinkOptString, A);
     if (!PostLinkOptString.empty())

clang/test/Driver/sycl-post-link-options.cpp

@@ -0,0 +1,21 @@
+/// Verify same set of sycl-post-link options generated for old and new offloading model
+// Test for JIT compilation
+// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \
+// RUN:          -Xdevice-post-link -O0 -v %s 2>&1 \
+// RUN:   | FileCheck -check-prefix OPTIONS_POSTLINK_JIT %s
+// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver \
+// RUN:          -Xdevice-post-link -O0 -v %s 2>&1 \
+// RUN:   | FileCheck -check-prefix OPTIONS_POSTLINK_JIT %s
+// OPTIONS_POSTLINK_JIT: sycl-post-link{{.*}} -O0 -O2 -device-globals -spec-const=native -split=auto -emit-only-kernels-as-entry-points -emit-param-info -symbols -emit-exported-symbols -split-esimd -lower-esimd
+
+#include <sycl/sycl.hpp>
+using namespace sycl;
+
+int main(void) {
+  sycl::queue queue;
+  sycl::event event = queue.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for<class set_range>(sycl::range<1>{16},
+                                      [=](sycl::id<1> idx) {});
+  });
+  return 0;
+}

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

@@ -522,78 +522,72 @@ static Expected<StringRef> convertSPIRVToIR(StringRef Filename,
   return *TempFileOrErr;
 }
 
-// Update sycl-post-link options based on target triple.
-static void updateCmdArgs(SmallVector<StringRef, 8> &CmdArgs,
-                          llvm::Triple Triple) {
-  // Get an argument in CmdArgs that contains Str. If there is no such
-  // argument, an empty argument is returned
-  auto getArg = [&](const StringRef &Str) {
-    for (auto Arg : CmdArgs)
-      if (Arg.contains(Str))
-        return Arg;
-    return StringRef("");
-  };
-  // Add a new argument Arg to CmdArgs if not present already.
-  auto addArg = [&](const StringRef &Arg) {
-    if (getArg(Arg).empty())
-      CmdArgs.push_back(Arg);
-  };
-  // Replace an argument in CmdArgs that contains Str with NewArg. If no such
-  // argument is present, add the NewArg to CmdArgs.
-  auto replaceOrAddArg = [&](const StringRef &NewArg, const StringRef &Str) {
-    for (auto &Arg : CmdArgs)
-      if (Arg.contains(Str)) {
-        Arg = NewArg;
-        return;
-      }
-    CmdArgs.push_back(NewArg);
-  };
-  // Remove argument containing Str from CmdArgs.
-  auto removeArg = [&](const StringRef &Str) {
-    CmdArgs.erase(
-        std::remove_if(CmdArgs.begin(), CmdArgs.end(),
-                       [&](StringRef Arg) { return Arg.contains(Str); }),
-        CmdArgs.end());
-  };
-
-  // specialization constants processing.
-  bool IsAOTGPU = Triple.isNVPTX() || Triple.isAMDGCN() || Triple.isSPIRAOT();
-  if (!IsAOTGPU)
-    replaceOrAddArg("-spec-const=native", "-spec-const");
-  else
-    replaceOrAddArg("-spec-const=emulation", "-spec-const");
-
-  // -emit-only-kernels-as-entry-points is set by the user and is enabled only
-  // for Intel targets.
-  auto EmitOnlyKernelsAsEntryPointsArg =
-      getArg("-emit-only-kernels-as-entry-points");
-  if ((!EmitOnlyKernelsAsEntryPointsArg.empty()) && !Triple.isNVPTX() &&
-      !Triple.isAMDGPU())
-    addArg("-emit-only-kernels-as-entry-points");
+// Add any sycl-post-link options that rely on a specific Triple in addition
+// to user supplied options.
+// NOTE: Any changes made here should be reflected in the similarly named
+// function in clang/lib/Driver/ToolChains/Clang.cpp.
+static void
+getTripleBasedSYCLPostLinkOpts(const ArgList &Args,
+                               SmallVector<StringRef, 8> &PostLinkArgs,
+                               const llvm::Triple Triple) {
+  const llvm::Triple HostTriple(Args.getLastArgValue(OPT_host_triple_EQ));
+  bool SYCLNativeCPU = (HostTriple == Triple);
+  bool SpecConsts = !(Triple.isNVPTX() || Triple.isAMDGCN() ||
+                      Triple.isSPIRAOT() || SYCLNativeCPU);
+  if (SpecConsts)
+    PostLinkArgs.push_back("-spec-const=native");
   else
-    removeArg("-emit-only-kernels-as-entry-points");
-
-  if (!(Triple.isAMDGCN()))
-    addArg("-emit-param-info");
-
-  if (Triple.isNVPTX() || Triple.isAMDGCN())
-    addArg("-emit-program-metadata");
-
-  if (Triple.isSPIROrSPIRV()) {
-    addArg("-symbols");
-    addArg("-emit-exported-symbols");
-    addArg("-split-esimd");
-    addArg("-lower-esimd");
-  }
-
-  // Here, IsAOT includes x86_64 device as well.
-  bool IsAOT =
-      IsAOTGPU || Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64;
-  auto GenDeviceImageArg = getArg("-generate-device-image-default-spec-consts");
-  if ((!GenDeviceImageArg.empty()) && IsAOT)
-    addArg("-generate-device-image-default-spec-consts");
-  else
-    removeArg("-generate-device-image-default-spec-consts");
+    PostLinkArgs.push_back("-spec-const=emulation");
+
+  // See if device code splitting is requested.  The logic here works along side
+  // the behavior in getNonTripleBasedSYCLPostLinkOpts, where the option is
+  // added based on the user setting of -fsycl-device-code-split.
+  bool NoSplit = true;
+  for (auto Arg : PostLinkArgs)
+    if (Arg.contains("-split=")) {
+      NoSplit = false;
+      break;
+    }
+  if (NoSplit && (Triple.getSubArch() != llvm::Triple::SPIRSubArch_fpga))
+    PostLinkArgs.push_back("-split=auto");
+
+  // On Intel targets we don't need non-kernel functions as entry points,
+  // because it only increases amount of code for device compiler to handle,
+  // without any actual benefits.
+  // TODO: Try to extend this feature for non-Intel GPUs.
+  if ((!Args.hasFlag(OPT_no_sycl_remove_unused_external_funcs,
+                     OPT_sycl_remove_unused_external_funcs, false) &&
+       !SYCLNativeCPU) &&
+      !Triple.isNVPTX() && !Triple.isAMDGPU())
+    PostLinkArgs.push_back("-emit-only-kernels-as-entry-points");
+
+  if (!Triple.isAMDGCN())
+    PostLinkArgs.push_back("-emit-param-info");
+  // Enable program metadata
+  if (Triple.isNVPTX() || Triple.isAMDGCN() || SYCLNativeCPU)
+    PostLinkArgs.push_back("-emit-program-metadata");
+
+  bool SplitEsimdByDefault = Triple.isSPIROrSPIRV();
+  bool SplitEsimd =
+      Args.hasFlag(OPT_sycl_device_code_split_esimd,
+                   OPT_no_sycl_device_code_split_esimd, SplitEsimdByDefault);
+
+  // Symbol file and specialization constant info generation is mandatory -
+  // add options unconditionally
+  PostLinkArgs.push_back("-symbols");
+  PostLinkArgs.push_back("-emit-exported-symbols");
+  if (SplitEsimd)
+    PostLinkArgs.push_back("-split-esimd");
+  PostLinkArgs.push_back("-lower-esimd");
+
+  bool IsAOT = Triple.isNVPTX() || Triple.isAMDGCN() ||
+               Triple.getSubArch() == llvm::Triple::SPIRSubArch_fpga ||
+               Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen ||
+               Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64;
+  if (Args.hasFlag(OPT_sycl_add_default_spec_consts_image,
+                   OPT_no_sycl_add_default_spec_consts_image, false) &&
+      IsAOT)
+    PostLinkArgs.push_back("-generate-device-image-default-spec-consts");
 }
 
 // Run sycl-post-link tool
@@ -619,7 +613,7 @@ static Expected<StringRef> runSYCLPostLink(ArrayRef<StringRef> InputFiles,
   SYCLPostLinkOptions.split(CmdArgs, " ", /* MaxSplit = */ -1,
                             /* KeepEmpty = */ false);
   const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ));
-  updateCmdArgs(CmdArgs, Triple);
+  getTripleBasedSYCLPostLinkOpts(Args, CmdArgs, Triple);
   CmdArgs.push_back("-o");
   CmdArgs.push_back(*TempFileOrErr);
   for (auto &File : InputFiles)

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

@@ -170,3 +170,18 @@ def sycl_post_link_options_EQ : Joined<["--", "-"], "sycl-post-link-options=">,
 def llvm_spirv_options_EQ : Joined<["--", "-"], "llvm-spirv-options=">,
   Flags<[WrapperOnlyOption]>,
   HelpText<"Options that will control llvm-spirv step">;
+
+// Extra SYCL options to help generate sycl-post-link options that also depend
+// on target triple.
+def sycl_remove_unused_external_funcs : Flag<["--", "-"], "sycl-remove-unused-external-funcs">,
+  Flags<[WrapperOnlyOption, HelpHidden]>;
+def no_sycl_remove_unused_external_funcs : Flag<["--", "-"], "no-sycl-remove-unused-external-funcs">,
+  Flags<[WrapperOnlyOption, HelpHidden]>;
+def sycl_device_code_split_esimd : Flag<["--", "-"], "sycl-device-code-split-esimd">,
+  Flags<[WrapperOnlyOption, HelpHidden]>;
+def no_sycl_device_code_split_esimd : Flag<["--", "-"], "no-sycl-device-code-split-esimd">,
+  Flags<[WrapperOnlyOption, HelpHidden]>;
+def sycl_add_default_spec_consts_image : Flag<["--", "-"], "sycl-add-default-spec-consts-image">,
+  Flags<[WrapperOnlyOption, HelpHidden]>;
+def no_sycl_add_default_spec_consts_image : Flag<["--", "-"], "no-sycl-add-default-spec-consts-image">,
+  Flags<[WrapperOnlyOption, HelpHidden]>;