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

asudarsa · 2024-06-14T05:07:08Z

In an earlier PR (#14101), we added support to generate sycl-post-link options (that partly depend on target triple) in clang Driver and then update the options using target triple information in the clang-linker-wrapper.
This PR cleans up that implementation by fully generating these options inside the clang-linker-wrapper. This requires some more user-specified options to be passed into the clang-linker-wrapper.

Thanks

asudarsa · 2024-06-14T05:09:35Z

@sarnex

I have made changes based on PR 14101 review comments, Can you please take a look?

Thanks

asudarsa · 2024-06-14T05:13:32Z

clang/lib/Driver/ToolChains/Clang.cpp

-  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


Main changes in this function are removal of NewOffloadDriver boolean and some comment updates.

asudarsa · 2024-06-14T05:14:23Z

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

-  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


Removed all string parsing logic.

mdtoguchi · 2024-06-14T16:21:44Z

clang/lib/Driver/ToolChains/Clang.cpp

+  // Handle -Xdevice-post-link
+  TC.TranslateTargetOpt(T, TCArgs, CmdArgs, options::OPT_Xdevice_post_link,
+                        options::OPT_Xdevice_post_link_EQ,
+                        JA.getOffloadingArch());


The options passed by the user should override any options we set in the driver, so this should be moved to after getTripleBasedSYCLPostLinkOpts

... I just noticed that this used to be later in the function - why the move?

Ah....Ok....I just moved it to match the order in new-offloading model. I will move this down in both places. Thanks. Good catch.

mdtoguchi · 2024-06-14T16:26:24Z

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

+                                      [=](sycl::id<1> idx) {});
+  });
+  return 0;
+}


These two tests don't look particularly different. At a high level, these don't need to be scrutinized with the REQUIRES: restriction and should work as a single test with just different --target values to exercise the given target.

The test failed under windows. So, I split it into two.

I have two tests. But I have removed the need to compile them (replaced with dry runs)

sarnex

Thank you! Some comments below:

sarnex · 2024-06-14T20:49:24Z

clang/lib/Driver/ToolChains/Clang.cpp

  if (!TCArgs.hasArg(options::OPT_fsycl_device_code_split_EQ) &&
-      (NewOffloadDriver || !(Triple.getArchName() == "spir64_fpga")))
+      (!(Triple.getArchName() == "spir64_fpga")))


Can we make this

Suggested change

(!(Triple.getArchName() == "spir64_fpga")))

Triple.getArchName() != "spir64_fpga"

Would like to keep one pair of surrounding parenthesis...

sarnex · 2024-06-14T20:50:59Z

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

+                               const llvm::Triple Triple) {
+  const llvm::Triple HostTriple(Args.getLastArgValue(OPT_host_triple_EQ));
+  bool SYCLNativeCPU = (HostTriple == Triple);
+  bool SpecConsts = !(Triple.isNVPTX() || Triple.isAMDGCN() ||


Nit: I find it easier to read if we distrubute the '!', so !Triple.isNVPTS() && !Triple.isAMDGCN() ...

sarnex · 2024-06-14T20:52:02Z

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

+                               const llvm::Triple Triple) {
+  const llvm::Triple HostTriple(Args.getLastArgValue(OPT_host_triple_EQ));
+  bool SYCLNativeCPU = (HostTriple == Triple);
+  bool SpecConsts = !(Triple.isNVPTX() || Triple.isAMDGCN() ||


Can we rename this to be more descriptive? If this is what it is really for, maybe this

Suggested change

bool SpecConsts = !(Triple.isNVPTX() || Triple.isAMDGCN() ||

bool SpecConstsSupported = !(Triple.isNVPTX() || Triple.isAMDGCN() ||

Also we should do the same in the Clang.cpp function

sarnex · 2024-06-14T20:53:29Z

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

+  // added based on the user setting of -fsycl-device-code-split.
+  bool NoSplit = true;
+  for (auto Arg : PostLinkArgs)
+    if (Arg.contains("-split=")) {


I wonder if instead we can rely on the user settings and use Args.hasArg or something rather than a string check based on what the other function did.

This is one scenario where I felt the current approach is simpler than passing -fsycl-device-code-split option along to the clang-linker-wrapper. It opens up additional issues like passing -fsycl-device-code-split as well as -split=? etc.

the tradeoff is then we have multiple places changing for the tool option name itself and not the clang-linker-wrapper arg, so if we want to change the tool option or something then we have multiple places to update, but i don't think that's a huge deal

sarnex · 2024-06-14T20:56:07Z

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

+    PostLinkArgs.push_back("-split-esimd");
+  PostLinkArgs.push_back("-lower-esimd");
+
+  bool IsAOT = Triple.isNVPTX() || Triple.isAMDGCN() ||


Can we use isSPIRAOT to get rid of some of these checks?

bool isSPIRAOT() const { return isSPIR() && (getSubArch() == Triple::SPIRSubArch_fpga || getSubArch() == Triple::SPIRSubArch_gen || getSubArch() == Triple::SPIRSubArch_x86_64); }

Good one. Will do

sarnex · 2024-06-14T20:58:25Z

clang/lib/Driver/ToolChains/Clang.cpp

-  // 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


Suggested change

// to user supplied options. This function is invoked only for old offloading

// to user supplied options. This function is invoked only for the old offloading

sarnex · 2024-06-14T20:58:33Z

clang/lib/Driver/ToolChains/Clang.cpp

-  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


Suggested change

// model. For new offloading model, a slightly modified version of this

// model. For the new offloading model, a slightly modified version of this

I remember my first thesis report where my professor marked out atleast a 100 extraneous 'the's. Since then I have been stringent with them. Sorry.

haha dont worry it's super minor, i just have an eye for typos

sarnex · 2024-06-14T20:59:02Z

clang/lib/Driver/ToolChains/Clang.cpp

+// 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.


Do we have any plans to make this a single function that both can call somehow? Having two functions that need to be kept in sync is not ideal.

+1 on this. I will work with Mike Toguchi to identify a common location for this.

sarnex · 2024-06-14T21:04:52Z

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.


Suggested change

// on target triple.

// on the target triple.

sarnex · 2024-06-14T21:10:38Z

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

+// 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,


Sorry I might have missed something, but can you remind me why we need two functions, one that cares about the triple and one that doesn't? It seem like it would be simpler if we had a single function because then we won't need to rely on things the previous one did.

It's due to how clang-linker-wrapper processes images. clang-linker-wrapper receives a set of images each with its own target triple. So, we do not have a singular target triple at the call site of clang-linker-wrapper. The options which do not depend on the target triple can be computed once and passed into the linker wrapper. The options that depend on the target triple will need to be computed for each device image group ( group corresponds to a target triple) inside the clang-linker-wrapper. There are few other ways to implement this. one option will be to pass all required information to the clang-linker-wrapper and compute all of them in one place. Passing in all the required information seemed overkill. Current approach is agreeable, I Hope.

Thanks

assuming it doesnt get too much more complex i think it's okay

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

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

…-post-link command line Signed-off-by: Arvind Sudarsanam <[email protected]>

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

asudarsa requested review from a team as code owners June 14, 2024 05:07

asudarsa had a problem deploying to WindowsCILock June 14, 2024 05:07 — with GitHub Actions Failure

asudarsa commented Jun 14, 2024

View reviewed changes

asudarsa temporarily deployed to WindowsCILock June 14, 2024 05:39 — with GitHub Actions Inactive

asudarsa temporarily deployed to WindowsCILock June 14, 2024 08:51 — with GitHub Actions Inactive

asudarsa temporarily deployed to WindowsCILock June 14, 2024 09:24 — with GitHub Actions Inactive

mdtoguchi reviewed Jun 14, 2024

View reviewed changes

asudarsa had a problem deploying to WindowsCILock June 14, 2024 20:09 — with GitHub Actions Error

asudarsa temporarily deployed to WindowsCILock June 14, 2024 20:21 — with GitHub Actions Inactive

asudarsa requested a review from mdtoguchi June 14, 2024 20:45

sarnex reviewed Jun 14, 2024

View reviewed changes

asudarsa had a problem deploying to WindowsCILock June 14, 2024 21:42 — with GitHub Actions Error

asudarsa requested a review from sarnex June 14, 2024 21:46

asudarsa temporarily deployed to WindowsCILock June 14, 2024 21:46 — with GitHub Actions Inactive

mdtoguchi approved these changes Jun 14, 2024

View reviewed changes

asudarsa temporarily deployed to WindowsCILock June 14, 2024 22:49 — with GitHub Actions Inactive

sarnex approved these changes Jun 17, 2024

View reviewed changes

asudarsa had a problem deploying to WindowsCILock June 17, 2024 14:01 — with GitHub Actions Error

asudarsa added 6 commits June 17, 2024 07:38

[New offload model] Cleanup the way sycl-post-link options are genera…

a14cac0

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

Fix windows test fail

ce71e8c

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

Fix tests and move user specified options towards the end of the sycl…

68fa250

…-post-link command line Signed-off-by: Arvind Sudarsanam <[email protected]>

Test update

ac744f5

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

Address review comments

4f1eb5e

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

Test updates after latest merge with SYCL branch tip

164b5d5

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

asudarsa force-pushed the cleanup_sycl_post_link_options branch from 7ff9292 to 164b5d5 Compare June 17, 2024 15:03

asudarsa temporarily deployed to WindowsCILock June 17, 2024 15:05 — with GitHub Actions Inactive

asudarsa temporarily deployed to WindowsCILock June 17, 2024 16:23 — with GitHub Actions Inactive

asudarsa merged commit 3e474e0 into intel:sycl Jun 18, 2024
14 checks passed

asudarsa mentioned this pull request Aug 28, 2024

[SYCL][Doc] Release Notes for Jul'24 release #15138

Merged

	(!(Triple.getArchName() == "spir64_fpga")))
	Triple.getArchName() != "spir64_fpga"

	bool SpecConsts = !(Triple.isNVPTX() \|\| Triple.isAMDGCN() \|\|
	bool SpecConstsSupported = !(Triple.isNVPTX() \|\| Triple.isAMDGCN() \|\|

	// to user supplied options. This function is invoked only for old offloading
	// to user supplied options. This function is invoked only for the old offloading

	// model. For new offloading model, a slightly modified version of this
	// model. For the new offloading model, a slightly modified version of this

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

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

Uh oh!

Conversation

asudarsa commented Jun 14, 2024

Uh oh!

asudarsa commented Jun 14, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

mdtoguchi Jun 14, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

sarnex left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

asudarsa commented Jun 14, 2024 •

edited

Loading

mdtoguchi Jun 14, 2024 •

edited

Loading