[SYCL] Unique stable name rebase #3835

erichkeane · 2021-05-27T15:18:45Z

This reverts our existing implementation, updates the library uses of this, and adds the community patch back:

8347ee4

The original version of this was reverted, and @rjmcall provided some advice to architect a new solution. This is that solution. This implements a builtin to provide a unique name that is stable across compilations of this TU for the purposes of implementing the library component of the unnamed kernel feature of SYCL. It does this by running the Itanium mangler with a few modifications. Because it is somewhat common to wrap non-kernel-related lambdas in macros that aren't present on the device (such as for logging), this uniquely generates an ID for all lambdas involved in the naming of a kernel. It uses the lambda-mangling number to do this, except replaces this with its own number (starting at 10000 for readabililty reasons) for lambdas used to name a kernel. Additionally, this implements itself as constexpr with a slight catch: if a name would be invalidated by the use of this lambda in a later kernel invocation, it is diagnosed as an error (see the Sema tests). Differential Revision: https://reviews.llvm.org/D103112

…ect spelling

erichkeane · 2021-05-27T15:19:38Z

Note that I noticed a bunch of check-all failures locally that i'm going to work through as well. There shouldn't be much to review here, as the community reviewed the big patch, and this mostly just removes the old one and adds that. There are a few library changes that are just changing the name.

s-kanaev

RT changes LGTM

… crashes

I discovered when merging the __builtin_sycl_unique_stable_name into my downstream that it is actually possible for the cc1 invocation to have more than 1 Sema instance, if you pass it multiple input files, each gets its own Sema instance and thus ASTContext instance. The result was that the call to Filter the SYCL kernels was using an ItaniumMangleContext stored via a 'magic static', so it had an invalid reference to ASTContext when processing the 2nd failure. The failure is unfortunately flakey/transient, but the test that fails was added anyway. The magic-static was switched to a unique_ptr member variable in ASTContext that is initialized when needed.

erichkeane · 2021-05-27T22:17:24Z

@AlexeySachkov : Can you look into the level-zero issue that showed up on the build-bots? I look at it when it came up on the linux bot, but don't have a good idea on how to start debugging it.

AlexeySachkov · 2021-05-31T14:17:02Z

@AlexeySachkov : Can you look into the level-zero issue that showed up on the build-bots? I look at it when it came up on the linux bot, but don't have a good idea on how to start debugging it.

@erichkeane, the failure is not directly related to level-zero, but it is related to "unnamed lambda" functionality. In Jenkins/Precommit there are ~18 failures and all of them because kernel wasn't found in device images.

I've used minimalistic (could be reduced more, of course) example to reproduce:

#include <CL/sycl.hpp>

constexpr sycl::specialization_id<int> test_id_1{42};

int main() {

  sycl::queue q;
  {
    sycl::buffer<double, 1> Buf{sycl::range{1}};
    q.submit([&](sycl::handler &cgh) {
      auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
      cgh.set_specialization_constant<test_id_1>(1);
      cgh.single_task<class Kernel1>([=](sycl::kernel_handler kh) {
        Acc[0] = kh.get_specialization_constant<test_id_1>();
      });
    });
    auto Acc = Buf.get_access<sycl::access::mode::read>();
    assert(Acc[0] == 1);
  }
}

$ clang++ -fsycl -fsycl-unnamed-lambda t.cpp
$ ./a.out
terminate called after throwing an instance of 'cl::sycl::exception'
  what():  The kernel bundle does not contain the kernel identified by kernelId. -30 (CL_INVALID_VALUE)
Aborted (core dumped)

KernelInfo::getName() returns an empty string. From what I see it seems like the template arguments of KernelInfoData structure (defined in kernel_desc.hpp) are different between the integration header and instantiation of that structure.

From integration header:

 // names of all kernels defined in the corresponding source
 static constexpr
 const char* const kernel_names[] = {
   "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE10000_clES2_EUlNS0_14kernel_handlerEE10000_
 template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', '4', 'm', 'a', 'i', 'n', 'E', 'N', 'K', 'U', 'l', 'R', 'N', '2', 'c'
 , 'l', '4', 's', 'y', 'c', 'l', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '1', '0', '0', '0', '0', '_', 'c', 'l', 'E', 'S',
 '2', '_', 'E', 'U', 'l', 'N', 'S', '0', '_', '1', '4', 'k', 'e', 'r', 'n', 'e', 'l', '_', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E
 ', '1', '0', '0', '0', '0', '_'> {

During instantiation (I've commented getName method default implementation to trigger the error)

include/sycl/CL/sycl/handler.hpp:540:13: error: no member named
      'getName' in 'sycl::detail::KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', '4', 'm', 'a', 'i', 'n', 'E', 'N', 'K', 'U', 'l', 'R',
      'N', '2', 'c', 'l', '4', 's', 'y', 'c', 'l', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', '2', '_',
      'E', 'U', 'l', 'N', 'S', '0', '_', '1', '4', 'k', 'e', 'r', 'n', 'e', 'l', '_', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_'>'
// Or: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_EUlNS0_14kernel_handlerEE_

erichkeane · 2021-05-31T23:01:34Z

Interesting.... The 'host' side is obviously wrong there, since it is a lambda without the unnamed-lambda spelling. Could this be a missed "sycl_kernel" attribute? One of the gotchas with the new mechanism is that the sycl_kernel attribute needs to be present on both sides.

AaronBallman

LGTM

erichkeane · 2021-06-01T14:51:33Z

1 more patch :) @romanovvlad please take a look at the warnings.cpp test from here: 45d489e to figure out if you want to make other different changes to your RT commits here.

Thanks!

romanovvlad · 2021-06-01T16:29:11Z

1 more patch :) @romanovvlad please take a look at the warnings.cpp test from here: 45d489e to figure out if you want to make other different changes to your RT commits here.

Thanks!

Will try to do tomorrow.

In the case where the device is an itanium target, and the host is a windows target, we were getting the names wrong, since in the itanium case we filter by lambda-signature. The fix is to always filter by the signature rather than just on non-windows builds. I considered doing the reverse (that is, checking the aux-triple), but doing so would result in duplicate lambda mangling numbers (from linux reusing the same number for different signatures).

…me_rebase

AaronBallman

LGTM!

s-kanaev

Other than the questions, RT changes LGTM

s-kanaev · 2021-06-03T11:13:03Z

sycl/include/CL/sycl/handler.hpp

@@ -825,9 +825,8 @@ class __SYCL_EXPORT handler {

      range<Dims> AdjustedRange = NumWorkItems;
      AdjustedRange.set_range_dim0(NewValX);
-#ifdef __SYCL_DEVICE_ONLY__
      kernel_parallel_for_wrapper<NameWT, TransformedArgType>(Wrapper);


This is performed in both host and device, isn't it?
Why is that needed?

@romanovvlad made these changes, but the root of it is that we're required to have the kernel instantiated on both sides now, since otherwise we cant tell whether a lambda is used to name a kernel on the host-side.

s-kanaev · 2021-06-03T11:13:14Z

sycl/include/CL/sycl/handler.hpp

      (void)NumWorkItems;
      kernel_parallel_for_wrapper<NameT, TransformedArgType>(KernelFunc);


Same question here and in similar places

premanandrao

LGTM

bader · 2021-06-03T18:17:24Z

@erichkeane, this patch breaks "clang-12 + shared libraries" post-commit configuration - https://github.com/intel/llvm/runs/2738419633?check_suite_focus=true.
Could you take a look, please?

erichkeane · 2021-06-03T18:20:57Z

@erichkeane, this patch breaks "clang-12 + shared libraries" post-commit configuration - https://github.com/intel/llvm/runs/2738419633?check_suite_focus=true.
Could you take a look, please?

It appears the problem is in the library changes that @romanovvlad committed. The 'simple' fix is to just push a commit that casts the parameters in each of these functions to 'void'. I can prepare a patch to do that @bader, though I'm not sure how to submit it?

bader · 2021-06-03T18:26:50Z

Could you open a pull request with the patch you proposed, please? If possible, please, check locally that "self-build + shared lib" configuration passes compilation + lit tests. Unfortunately, this configuration is covered by post-commit only, so pre-commit doesn't validate that it's fixed.

erichkeane · 2021-06-03T18:28:21Z

Submitted a PR here: #3874

I don't have a self-build + shared-lib config available on my test environment, but the errors/changes seem quite simple.

bader · 2021-06-03T18:36:48Z

Interesting... I also see that Clang :: Driver/sycl-device-lib.cpp test from check-clang fails. Do you think it's related to the same problem and will be fixed by #3874?
+@mdtoguchi

erichkeane · 2021-06-03T18:38:56Z

I have seen that failure in my local repo for a few months now, I don't think it is related to this at all.

bader · 2021-06-03T18:42:38Z

I think it might be related to #3630, but I thought that it's fixed by #3639. @jinge90, could you check, please?

erichkeane · 2021-06-03T18:47:18Z

So my mistake, that one no longer fails for me (but did for quite a while). Even so, I don't see anything in this patch or the #3874 having the ability to change the behavior of that test at all.

FWIW, my environment's failures are now:
Clang :: CodeGenCUDA/host-used-device-var.cu
Clang :: CodeGenCUDA/unused-global-var.cu
Clang :: Driver/amdgpu-openmp-toolchain.c
Clang :: Driver/hip-include-path.hip

bader · 2021-06-03T18:55:22Z

I suspect that it might be that issue #3630 is not fixed, but your fix #3874 will hide it for CI. It's reproducible if you run clang tests w/o building sycl project. CI always builds sycl project before running clang lit tests. In this case it fails to build sycl project and run check-clang.

* sycl: (320 commits) [SYCL] Silence a "local variable is initialized but not referenced" warning; NFC (intel#3870) [SYCL] Improve SYCL_DEVICE_ALLOWLIST (intel#3826) [SPIR-V] Change return value of mapType function (intel#3871) [SYCL] Fix post-commit failure in handler.hpp from unused-parameters. (intel#3874) [Driver][SYCL] Do not imply defaultlib msvcrt for Linux based driver on Windows (intel#3827) [SYCL] Unique stable name rebase (intel#3835) [SYCL] Align behavior of empty command groups with SYCL2020 (intel#3822) [SYCL][ESIMD] Make typenames and constants consistent with SYCL API style. (intel#3850) [SYCL] Allow __failed_assertion to support libstdc++-11 (intel#3774) [SYCL] Refactor stream class handing implementation (intel#3646) [SYCL] Fix syntax error introduced in intel#3401 (intel#3861) [SYCL] SYCL 2020 sub_group algorithms (intel#3786) [Buildbot][NFC] Add option to use LLD as linker (intel#3866) Revert "Emit correct location lists with basic block sections." [SPIRITTAnnotations] Fix debug info for ITT calls. (intel#3829) [SYCL][Doc] Fix build of Sphinx docs (intel#3863) [SYCL][FPGA][NFC] Tidy up intel_fpga_reg codegen test (intel#3810) [CODEOWNERS] Fix SPIRITTAnnnotations tests ownership (intel#3859) [SYCL][ESIMD] Host-compile simd.cpp test, fix errors & warnings. (intel#3846) [SYCL] Store pointers to memory allocations instead of iterators (intel#3860) ...

Resolved by reverting: eba69b5 Reimplement __builtin_unique_stable_name- And applying: 70281ea [SYCL] Unique stable name rebase (intel#3835) CONFLICT (content): Merge conflict in clang/lib/Sema/SemaTemplateInstantiateDecl.cpp CONFLICT (content): Merge conflict in clang/lib/Sema/SemaSYCL.cpp CONFLICT (content): Merge conflict in clang/lib/Sema/SemaExpr.cpp CONFLICT (content): Merge conflict in clang/lib/Parse/ParseExpr.cpp CONFLICT (content): Merge conflict in clang/lib/Basic/IdentifierTable.cpp CONFLICT (content): Merge conflict in clang/lib/AST/ItaniumMangle.cpp CONFLICT (content): Merge conflict in clang/include/clang/Sema/Sema.h CONFLICT (content): Merge conflict in clang/include/clang/Parse/Parser.h CONFLICT (content): Merge conflict in clang/include/clang/Basic/TokenKinds.def CONFLICT (content): Merge conflict in clang/include/clang/AST/Mangle.h CONFLICT (content): Merge conflict in clang/docs/LanguageExtensions.rst

Erich Keane added 4 commits May 27, 2021 07:25

Revert old implementation sans tests

bd97889

Correct comment that referred to old name

6672484

Update a bunch of library uses of unique-stable-name to have the corr…

47791c4

…ect spelling

erichkeane requested review from AaronBallman, elizabethandrews, premanandrao and a team as code owners May 27, 2021 15:18

erichkeane requested a review from s-kanaev May 27, 2021 15:18

Fix the SYCL-Sema uses

8bdd49a

s-kanaev previously approved these changes May 27, 2021

View reviewed changes

Erich Keane added 3 commits May 27, 2021 09:01

Fix the sema-sycl tests, which just needed check values updated

eb0b1c3

update unique-stable-name codegen test

a723f9a

Fix most of the CodeGenSYCL tests, int_header_sycl2020_spec_const.cpp…

fa740a5

… crashes

erichkeane dismissed s-kanaev’s stale review via fa740a5 May 27, 2021 17:49

Erich Keane added 3 commits May 27, 2021 13:47

Add Wno-sycl-2020-compat to test so that it passes with no diagnostics

2c29e59

Update buffer-location-codegen test to match the new mangling

09a07c6

romanovvlad and others added 5 commits June 1, 2021 06:42

[SYCL] Instantiate sycl_kernel functions in host compilation

3a21d47

Add workaround for library-host mode

72738e0

[SYCL] Use if constexpr instead of enable_if to avoid msvc bug

5f82ac6

Fix clang-format

1ccdb17

Another clang-format ifx

255a450

AaronBallman previously approved these changes Jun 1, 2021

View reviewed changes

Disabling issues in 2 tests to hopefully make CI happy

45d489e

erichkeane dismissed AaronBallman’s stale review via 45d489e June 1, 2021 14:50

Erich Keane added 3 commits June 2, 2021 13:17

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

6c947a9

…me_rebase

Update windows-diff test to force emit of unique-stable-name

47a130a

bader requested review from s-kanaev, AaronBallman and romanovvlad June 3, 2021 09:23

bader changed the title ~~[SYCL]Unique stable name rebase~~ [SYCL] Unique stable name rebase Jun 3, 2021

AaronBallman reviewed Jun 3, 2021

View reviewed changes

s-kanaev approved these changes Jun 3, 2021

View reviewed changes

premanandrao approved these changes Jun 3, 2021

View reviewed changes

bader merged commit 70281ea into intel:sycl Jun 3, 2021

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL] Unique stable name rebase #3835

[SYCL] Unique stable name rebase #3835

erichkeane commented May 27, 2021

erichkeane commented May 27, 2021

s-kanaev left a comment

erichkeane commented May 27, 2021

AlexeySachkov commented May 31, 2021

erichkeane commented May 31, 2021

AaronBallman left a comment

erichkeane commented Jun 1, 2021

romanovvlad commented Jun 1, 2021

AaronBallman left a comment

s-kanaev left a comment

s-kanaev Jun 3, 2021

erichkeane Jun 3, 2021

s-kanaev Jun 3, 2021 •

edited

Loading

premanandrao left a comment

bader commented Jun 3, 2021

erichkeane commented Jun 3, 2021

bader commented Jun 3, 2021

erichkeane commented Jun 3, 2021

bader commented Jun 3, 2021

erichkeane commented Jun 3, 2021

bader commented Jun 3, 2021

erichkeane commented Jun 3, 2021

bader commented Jun 3, 2021

		(void)NumWorkItems;
		kernel_parallel_for_wrapper<NameT, TransformedArgType>(KernelFunc);

[SYCL] Unique stable name rebase #3835

[SYCL] Unique stable name rebase #3835

Conversation

erichkeane commented May 27, 2021

erichkeane commented May 27, 2021

s-kanaev left a comment

Choose a reason for hiding this comment

erichkeane commented May 27, 2021

AlexeySachkov commented May 31, 2021

erichkeane commented May 31, 2021

AaronBallman left a comment

Choose a reason for hiding this comment

erichkeane commented Jun 1, 2021

romanovvlad commented Jun 1, 2021

AaronBallman left a comment

Choose a reason for hiding this comment

s-kanaev left a comment

Choose a reason for hiding this comment

s-kanaev Jun 3, 2021

Choose a reason for hiding this comment

erichkeane Jun 3, 2021

Choose a reason for hiding this comment

s-kanaev Jun 3, 2021 • edited Loading

Choose a reason for hiding this comment

premanandrao left a comment

Choose a reason for hiding this comment

bader commented Jun 3, 2021

erichkeane commented Jun 3, 2021

bader commented Jun 3, 2021

erichkeane commented Jun 3, 2021

bader commented Jun 3, 2021

erichkeane commented Jun 3, 2021

bader commented Jun 3, 2021

erichkeane commented Jun 3, 2021

bader commented Jun 3, 2021

s-kanaev Jun 3, 2021 •

edited

Loading