[SYCL] Change unique-stable-name to just use the lambda's mangled name. #4457

erichkeane · 2021-09-01T19:52:40Z

Currently, whether a lambda is used in a kernel is used to determine the
result of __builtin_sycl_unique_stable_name.This leads to a sitaution
where using a lambda in a kernel that is lexically higher than an
existing kernel will change the result of the builtin, causing a build
failure.

This patch replaces this with JUST the lambda's itanium-mangled name.
This has the limitation in that it is not stable/unique in light of
certain macro-expansions that differ between device and host, however
wording is being put into the SYCL standard to make those cases UB.

This DOES have one bit of fallout, which is to make sure that the name
on a Windows host does NOT have its lambda mangling 'number', which we
are doing by storing the value in the Device Lamba Mangling Number
(previously done for CUDA only for similar reasons!).

Additionally, this patch removes the 'mark_kernel_name' builtin, as this
results in it being a no-op.

Currently, whether a lambda is used in a kernel is used to determine the result of `__builtin_sycl_unique_stable_name`.This leads to a sitaution where using a lambda in a kernel that is lexically higher than an existing kernel will change the result of the builtin, causing a build failure. This patch replaces this with JUST the lambda's itanium-mangled name. This has the limitation in that it is not stable/unique in light of certain macro-expansions that differ between device and host, however wording is being put into the SYCL standard to make those cases UB. This DOES have one bit of fallout, which is to make sure that the name on a Windows host does NOT have its lambda mangling 'number', which we are doing by storing the value in the Device Lamba Mangling Number (previously done for CUDA only for similar reasons!). Additionally, this patch removes the 'mark_kernel_name' builtin, as this results in it being a no-op.

erichkeane · 2021-09-01T19:53:48Z

Attn @rolandschulz @AlexeySachkov @rarutyun

clang/test/SemaSYCL/unique_stable_name.cpp

sycl/include/CL/sycl/kernel.hpp

premanandrao

Looks good overall, just minor nits.

clang/lib/AST/ItaniumMangle.cpp

clang/test/SemaSYCL/unique-stable-id.cpp

clang/test/SemaSYCL/unique_stable_name.cpp

clang/lib/AST/Expr.cpp

clang/lib/AST/ItaniumMangle.cpp

The Unique-stable-** now make sure to use 'device mangling number' always all of the time, rather than counting on Itanium to always set it to 0, which would result in us using the 'lambda number'. This patch instead changes our ItaniumCXXABI to always set the device mangling number to match the lambda mangling number, so we can always use the device mangling number if it exists. This patch also fixes all of the review comments so far.

…_name minor conflict in : clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp

…cient for his request for this test

sycl/test/regression/unnamed-lambda-split-order.cpp

@AlexeySachkov

…mbda.cpp instead as @AlexeySachkov requested offline

dm-vodopyanov

sycl/ changes LGTM

sycl/test/regression/unnamed-lambda.cpp

premanandrao · 2021-09-02T17:47:50Z

clang/lib/AST/ItaniumMangle.cpp

-                        ? *DeviceNumber
-                        : Lambda->getLambdaManglingNumber();
+  unsigned Number =
+      DeviceNumber ? *DeviceNumber : Lambda->getLambdaManglingNumber();

  assert(Number > 0 && "Lambda should be mangled as an unnamed class");
  if (Number > 1)


What happens for 3 here? Would it not conflict with 1?

No, because '3' would mangle the number '1', '1' would mangle 'nothing'.

Thanks, so there is {"unnamed class", nothing, 0,1,...} for {0, 1, 2, 3, ...} respectively. Did I get that right?

exactly. Isn't Itanium fun :)?

premanandrao

LGTM

AaronBallman · 2021-09-02T18:17:45Z

sycl/test/regression/unnamed-lambda.cpp

+  sycl::queue q;
+  q.parallel_for(10, [](auto i) {});
+  q.parallel_for(10, w);
+}


Do you think there would be value in adding a test that uses macros in a problematic way to demonstrate that we don't assert/crash/do bad things? Or do you think we have sufficient coverage for that already?

I don't see value in seeing if the compiler would crash/etc. But we DO know that the runtime behavior of those programs would be bad. I have no idea how to test those situations.

I don't want to see if we do crash; I want to prove that we don't crash in a known-problematic situation.

Perhaps instead of a FE test, it'd make sense for a SYCL end-to-end test? I don't want to codify "this is the behavior you can rely on when you do this stuff", but I do think there's some benefit to showing "we know this code pattern is a sharp edge and we explicitly don't care what happens so long as 'compiler crash' isn't the behavior."

I see. The expected outcome is that the 'wrong' kernel would be called, or the runtime would fail (depending on how this is configured). We're intentionally making this "UB" in the standard to cover both of those cases.

Would it be ok if @AlexeySachkov or @romanovvlad could write a test in a separate patch?

It'd be totally fine to add a test in a separate review; no need to block this review on it. I agree on the expected outcome -- the value I was seeing in the test is showing that we don't care about the runtime behavior (you get what you get), but that we don't expect the FE to explode on it. However, if others don't think that's a particularly valuable test, I don't insist on one being added at all.

Not sure if it makes a lot of sense to test UB. AFAIK users do not look at our tests to see what is code is legal.

The goal isn't to "test UB", it's to ensure our product doesn't crash just because of UB. We added new assertions, we changed old ones, and we did it on a fast-turnaround patch on something we're now on our third iteration of due to issues discovered in the field. As I said, I don't insist on the test, but hopefully we're at least all talking about the same thing being tested when we make that decision.

This is a test for unnamed kernels, so if we don't have unnamed kernels, it isn't meaningful.

AlexeySachkov

sycl/ changes LGTM

erichkeane · 2021-09-02T21:18:50Z

@ping all! Looks like I need a CFE and a intel/llvm-reviews-runtime reviewer. All but the one bot have finished, so just waiting on that one now.

AaronBallman

LGTM!

AaronBallman · 2021-09-02T21:22:11Z

@ping all! Looks like I need a CFE and a intel/llvm-reviews-runtime reviewer. All but the one bot have finished, so just waiting on that one now.

Oops, it seemed my previous approval dropped. I re-approved for CFE.

erichkeane · 2021-09-02T21:57:31Z

@dm-vodopyanov : Looks like I need a re-+1 from you or someone else from intel/llvm-reviews-runtime when you get a chance. Thanks!

erichkeane · 2021-09-03T01:33:13Z

@bader : The failure is another one of the 'out of disk space errors': /localdisk2/sycl_ci/buildbot/worker/Lit_With_Cuda/llvm.src/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp:8989:1: fatal error: error writing to /tmp/ccuEHlOL.s: No space left on device

Looks like it has been failed for hours. I've submitted a ticket internally to hopefully get this taken care of, but I thought you might have some ideas on who to ping as well.

Thanks!

bader · 2021-09-03T05:54:39Z

Looks like it has been failed for hours. I've submitted a ticket internally to hopefully get this taken care of, but I thought you might have some ideas on who to ping as well.

AFAIK, @tfzhu is the contact point for CI system.

erichkeane requested review from AaronBallman, elizabethandrews, premanandrao and a team as code owners September 1, 2021 19:52

erichkeane requested a review from dm-vodopyanov September 1, 2021 19:52

AlexeySachkov reviewed Sep 1, 2021

View reviewed changes

clang/test/SemaSYCL/unique_stable_name.cpp Show resolved Hide resolved

jbrodman reviewed Sep 1, 2021

View reviewed changes

sycl/include/CL/sycl/kernel.hpp Outdated Show resolved Hide resolved

Fix typo in kernel.hpp

34ae59b

premanandrao reviewed Sep 2, 2021

View reviewed changes

clang/lib/AST/ItaniumMangle.cpp Outdated Show resolved Hide resolved

clang/test/SemaSYCL/unique-stable-id.cpp Outdated Show resolved Hide resolved

clang/test/SemaSYCL/unique_stable_name.cpp Outdated Show resolved Hide resolved

AaronBallman reviewed Sep 2, 2021

View reviewed changes

clang/lib/AST/Expr.cpp Outdated Show resolved Hide resolved

clang/lib/AST/ItaniumMangle.cpp Outdated Show resolved Hide resolved

clang/lib/AST/ItaniumMangle.cpp Outdated Show resolved Hide resolved

Erich Keane added 4 commits September 2, 2021 07:41

Merge remote-tracking branch 'SYCL_public/sycl' into USN_Just_mangled…

6953012

…_name minor conflict in : clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp

Finish my comment, fix the clang-format with it

dff7734

Add sycl regression test from @AlexeySachkov 's repo, hopefully suffi…

1ffbae7

…cient for his request for this test

dm-vodopyanov reviewed Sep 2, 2021

View reviewed changes

sycl/test/regression/unnamed-lambda-split-order.cpp Outdated Show resolved Hide resolved

Add comment to the regression suite test case, put it into unnamed-la…

e37ef8c

…mbda.cpp instead as @AlexeySachkov requested offline

dm-vodopyanov previously approved these changes Sep 2, 2021

View reviewed changes

premanandrao reviewed Sep 2, 2021

View reviewed changes

Fix typo in comment

7700932

erichkeane dismissed dm-vodopyanov’s stale review via 7700932 September 2, 2021 17:52

premanandrao previously approved these changes Sep 2, 2021

View reviewed changes

AaronBallman previously approved these changes Sep 2, 2021

View reviewed changes

Move the test back, since it doesn't work in 1.2.1 mode

4eaeb32

This is a test for unnamed kernels, so if we don't have unnamed kernels, it isn't meaningful.

erichkeane dismissed stale reviews from AaronBallman and premanandrao via 4eaeb32 September 2, 2021 18:38

Make the test use 'main' so that it links

7ea633d

AlexeySachkov approved these changes Sep 2, 2021

View reviewed changes

AaronBallman approved these changes Sep 2, 2021

View reviewed changes

dm-vodopyanov approved these changes Sep 3, 2021

View reviewed changes

bader merged commit 27c632e into intel:sycl Sep 3, 2021

joeatodd mentioned this pull request Oct 13, 2021

Build error for sycl::get_kernel_id for functor containing lambda #4756

Closed

masterleinad mentioned this pull request Oct 13, 2021

Use sycl::info::kernel_device_specific::preferred_work_group_size_multiple kokkos/kokkos#4405

Closed

[SYCL] Change unique-stable-name to just use the lambda's mangled name. #4457

[SYCL] Change unique-stable-name to just use the lambda's mangled name. #4457

Uh oh!

Conversation

erichkeane commented Sep 1, 2021

Uh oh!

erichkeane commented Sep 1, 2021

Uh oh!

Uh oh!

Uh oh!

premanandrao left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

dm-vodopyanov left a comment

Choose a reason for hiding this comment

Uh oh!

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!

premanandrao 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!

AlexeySachkov left a comment

Choose a reason for hiding this comment

Uh oh!

erichkeane commented Sep 2, 2021

Uh oh!

AaronBallman left a comment

Choose a reason for hiding this comment

Uh oh!

AaronBallman commented Sep 2, 2021

Uh oh!

erichkeane commented Sep 2, 2021

Uh oh!

erichkeane commented Sep 3, 2021

Uh oh!

bader commented Sep 3, 2021

Uh oh!

Uh oh!