[SYCL] Rework MarkDevice and children #3475

erichkeane · 2021-04-02T15:05:58Z

We identified during a previous review that collectKernelAttributes and
the MarkDeviceFunction object are in a situation that is very difficult
to maintain, so this attempts to fix that by making the ownership model
clear.

We identified during a previous review that collectKernelAttributes and the MarkDeviceFunction object are in a situation that is very difficult to maintain, so this attempts to fix that by making the ownership model clear.

erichkeane · 2021-04-02T15:08:05Z

FYI @jinge90, @smanna12, @premanandrao, and @AaronBallman. I think I fixed a couple of bugs I found as I went through here, but want to see if we break anything later in the process. Note there are a bunch of "TODOs" as I discovered things that we could likely do different/better, but I'd like to see the results of the pre-commit first to see if I break any of this.

erichkeane · 2021-04-02T15:16:37Z

@AaronBallman and @keryell : Note the last patch that adds a test. Seemingly the constexpr-if example already works!

erichkeane · 2021-04-02T15:20:15Z

clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp

@@ -10,15 +10,15 @@
 }

 #else
-[[cl::reqd_work_group_size(2, 2, 2)]] void not_direct_two() {} // expected-note {{conflicting attribute is here}}


Note that the changes here are because we are doing a DFS instead of a BFS. This obviously does not change the validity of any program.

IMO, the diagnostics are better now, since you are more likely to have the first diagnostic be 'child of the thing that added it' rather than a peer, which seems easier to debug as an end user to me. AND as added benefit, we only end up walking the call graph 1x instead of 2x.

erichkeane · 2021-04-02T15:23:43Z

clang/test/SemaSYCL/allow-constexpr-recursion.cpp

@@ -6,13 +6,11 @@

 sycl::queue q;

-// expected-note@+1{{function implemented using recursion declared here}}


So we were only able to figure this out by pure happenstance before, there is a bug in SemaSYCL.cpp currently where we decide that the caller of a recursive function is ALSO recursive. This is not a sound assumption, and even in the cases it is true, it doesn't catch all cases.

Fixing that ends up meaning we cannot catch the following example, but this does not change the validity of the program (since we get 1 diagnostic here instead of 2 like previously):

int func1(); int func2() { func1(); } int func1() { func2(); } kernel<class Foo>([](){ func1(); });

Because we have to give up on running through diagnostics, we catch that func1 is recursive, but cannot catch func2, because we don't continue running down the list.

erichkeane · 2021-04-02T15:24:45Z

clang/lib/Sema/SemaSYCL.cpp

@@ -328,6 +328,8 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD,
               SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
  });

+  // TODO: ERICH: We should probably warn on the bottom3 as well and turn this


This is one I'll probably visit once I get the build-bots to agree I'm OK with my refactor so far. It seems that we should still be warning instead of silently dropping here, but even if that is a problem, I still want to make it so we have much less copy/paste here.

Ugh, we actually have a number of lit tests that depend on us not diagnosing or removing the bottom 3 (now 4 after merging @zahiraam s patch). I simplified this now, but we can't combine all 4. That is perhaps a future "TODO".

erichkeane · 2021-04-02T15:24:58Z

clang/lib/Sema/SemaSYCL.cpp

-        }
-      }
-    }
+// TODO: ERICH: This likely needs a better name and documentation.


Suggestions welcome :)

erichkeane · 2021-04-02T15:25:20Z

clang/lib/Sema/SemaSYCL.cpp

-      WorkList.pop_back();
-      if (!Visited.insert(FD).second)
-        continue; // We've already seen this Decl
+  // TODO: is this too clever?  Should this just be a called function?


Let me know what you think here. See MarkDevices for how this is used.

clang/lib/Sema/SemaSYCL.cpp

erichkeane · 2021-04-02T15:27:30Z

clang/lib/Sema/SemaSYCL.cpp

+static void PropagateAndDiagnoseDeviceAttr(Sema &S, Attr *A,
+                                           FunctionDecl *SYCLKernel,
+                                           FunctionDecl *KernelBody) {
+  switch (A->getKind()) {


I pulled the attribute propagation into its own function here, MarkDevice/MarkDevices was just plain unwieldy as a result. Most of this is a copy/paste, except it is no longer a member function and has a different tab depth.

erichkeane · 2021-04-02T15:28:24Z

clang/lib/Sema/SemaSYCL.cpp

+}
+
+void Sema::MarkDevices(void) {
+  // This Tracker object ensures that the SyclDeviceDecls collection includes


So this is the entry-point here. The idea is that Tracker keeps track of the info for the entire TU, ensures that all SYCL_EXTERNAL functions are here, and ensures diagnostics happen on destruction.

clang/lib/Sema/SemaSYCL.cpp

erichkeane · 2021-04-02T15:30:06Z

clang/lib/Sema/SemaSYCL.cpp

+    // This type does the actual analysis on a per-kernel basis. It does this to
+    // make sure that we're only ever dealing with the context of a single
+    // kernel at a time.
+    SingleDeviceFunctionTracker T{Tracker, SYCLKernel};


I found we were constantly doing a set/reset of the collections owned by the parent when we were at a point where we only cared about the current kernel. This object does the analysis work, and just rejoins its information into the parent collection when it is done.

erichkeane · 2021-04-06T13:41:06Z

Alright, I think this is ready for review now (as am I). Thanks in advance all!

premanandrao · 2021-04-06T15:44:45Z

clang/lib/Sema/SemaSYCL.cpp

+                 SYCLIntelFPGADisableLoopPipeliningAttr,
+                 SYCLIntelFPGAInitiationIntervalAttr>(A);
+    });
+  }


Erich, can you explain what is happening with this copy_if? How are these attributes getting excluded?

So this is copying these ONLY IF this is a 'directly-called'. These attributes are all the set from the 'old' line 575, which only added them if they were 'directly called' from the kernel. They get copied if, for instance, they are the attribute on the lambda.

clang/lib/Sema/SemaSYCL.cpp

premanandrao · 2021-04-07T15:34:39Z

clang/lib/Sema/SemaSYCL.cpp

+  DeviceFunctionTracker Tracker(*this);
+
+  for (Decl *D : syclDeviceDecls()) {
+    FunctionDecl *SYCLKernel = cast<FunctionDecl>(D);


Make this 'auto'?

premanandrao · 2021-04-07T15:44:38Z

clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp


-[[intel::max_work_group_size(1, 1, 1)]] // expected-note {{conflicting attribute is here}}
+[[intel::max_work_group_size(1, 1, 1)]] // expected-note 3 {{conflicting attribute is here}}


What are the three conflicts? If I squint, I can see may be two :-)

Here is the test output:

/export/iusers/ekeane1/workspaces/sycl/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp:52:34: error: conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function parallel_for<class KernelName>([]() {}); // expected-error 3 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} ^ /export/iusers/ekeane1/workspaces/sycl/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp:15:3: note: conflicting attribute is here [[intel::max_work_group_size(1, 1, 1)]] // expected-note 3 {{conflicting attribute is here}} ^ /export/iusers/ekeane1/workspaces/sycl/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp:13:3: note: conflicting attribute is here [[cl::reqd_work_group_size(2, 2, 2)]] void not_direct_two() {} // expected-note 2 {{conflicting attribute is here}} ^ /export/iusers/ekeane1/workspaces/sycl/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp:52:34: error: conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function parallel_for<class KernelName>([]() {}); // expected-error 3 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} ^ /export/iusers/ekeane1/workspaces/sycl/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp:15:3: note: conflicting attribute is here [[intel::max_work_group_size(1, 1, 1)]] // expected-note 3 {{conflicting attribute is here}} ^ /export/iusers/ekeane1/workspaces/sycl/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp:21:3: note: conflicting attribute is here [[cl::reqd_work_group_size(4, 4, 4)]] // expected-note 1 {{conflicting attribute is here}} ^ /export/iusers/ekeane1/workspaces/sycl/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp:52:34: error: conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function parallel_for<class KernelName>([]() {}); // expected-error 3 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} ^ /export/iusers/ekeane1/workspaces/sycl/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp:15:3: note: conflicting attribute is here [[intel::max_work_group_size(1, 1, 1)]] // expected-note 3 {{conflicting attribute is here}} ^ /export/iusers/ekeane1/workspaces/sycl/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp:13:3: note: conflicting attribute is here [[cl::reqd_work_group_size(2, 2, 2)]] void not_direct_two() {} // expected-note 2 {{conflicting attribute is here}}

It conflicts with 'not_direct_two', func_three, then not-direct-2 does again. This is a side-effect of us not removing these attributes and leaving them in place. This is consistent with the existing implementation, except in this case we hit it 2x (instead of the others conflicting with eachother) due to the Depth-first nature.

AaronBallman

How awful would it be to also make this more const correct given that you're reworking much of it anyway?

clang/lib/Sema/SemaSYCL.cpp

erichkeane · 2021-04-07T20:33:14Z

How awful would it be to also make this more const correct given that you're reworking much of it anyway?

Can you point out some examples of where you mean? I tried to add const in a few places, but some of the dependencies make that difficult/impossible. If you have some offensive examples, I can at least try to work harder on those. For example, RecursiveASTVisitor's traverse functions for some reason are non-const, which made little sense to me. Additionally, we end up 'modifying' the function decls in a few places (like by adding or removing attributes!) so it tends to be difficult in other places too.

FYI: just building a patch to fix the comments you mentioned, so I didn't just resolve them for no reason :)

AaronBallman · 2021-04-08T13:21:07Z

clang/lib/Sema/SemaSYCL.cpp

-  MarkDeviceFunction(Sema &S)
-      : RecursiveASTVisitor<MarkDeviceFunction>(), SemaRef(S) {}
+  DiagDeviceFunction(
+      Sema &S, const llvm::SmallPtrSetImpl<FunctionDecl *> &RecursiveFuncs)


I was hoping we could make RecursiveFuncs hold const FunctionDecl * because nothing mutates the functions once they get passed in here. However, I see why that's really annoying (because we take the container rather than using a more general interface with iterators), so I suppose this is as const correct as we're likely to get.

Actually, DeviceFunctions cannot be const I think, but the RecursiveFuncs I can make const. I'll do that, thanks!

AaronBallman · 2021-04-08T13:51:24Z

clang/lib/Sema/SemaSYCL.cpp

+      if (Existing->getXDimVal(Ctx) != RWGSA->getXDimVal(Ctx) ||
+          Existing->getYDimVal(Ctx) != RWGSA->getYDimVal(Ctx) ||
+          Existing->getZDimVal(Ctx) != RWGSA->getZDimVal(Ctx)) {


This isn't specific to your review, but I think there's a bug here -- we're not checking to see how this attribute is spelled and so we may be comparing things we don't expect to compare. e.g., the OpenCL spelling for this attribute has argument order X, Y, Z and the SYCL spelling has argument order Z, Y, X.

lol... that is depressing. I just moved this code though.

Are we sure that the SemaDeclAttr part doesn't correct the order of these?

Yeah, that's why this is an existing bug. I'm pretttttty sure that SemaDeclAttr isn't correcting the order of those because we need to keep the order sensible for AST dumping, pretty printing, etc. IIRC, we tried that route but had to roll it back. @smanna12 may know more (or be able to look into whether this is a real issue or just me being paranoid).

Yes, this is a bug with the existing implementation. I plan to refactor this attribute (separate SYCL vs OpenCL spelling).

AaronBallman · 2021-04-08T13:51:58Z

clang/lib/Sema/SemaSYCL.cpp

+  }
+  case attr::Kind::SYCLIntelMaxWorkGroupSize: {
+    auto *SIMWGSA = cast<SYCLIntelMaxWorkGroupSizeAttr>(A);
+    if (auto *Existing = SYCLKernel->getAttr<ReqdWorkGroupSizeAttr>()) {


Similar concern here about not caring about the way the attribute was spelled.

AaronBallman · 2021-04-08T13:54:23Z

clang/lib/Sema/SemaSYCL.cpp

+  case attr::Kind::SYCLIntelFPGADisableLoopPipelining:
+  case attr::Kind::SYCLIntelFPGAInitiationInterval:
+  case attr::Kind::SYCLSimd: {
+    if ((A->getKind() == attr::Kind::SYCLSimd) && KernelBody &&


This is a bit messy -- put this logic in the attr::Kind::SYCLSimd case, then use LLVM_FALLTHROUGH?

Oof... i didn't notice that when moving this. THAT is simple enough and offensive enough I'll fix it :)

AaronBallman

Thanks! I think this generally LGTM. We can address the open question about reqd_work_group_size behavior in follow-ups.

Last time I use the web app for this :)

erichkeane · 2021-04-08T18:04:46Z

Only change with the merge was the conflict with the recursion becoming an error (so a different message). This needs re-review, but the change is quite minimal.

premanandrao

I am good with the changes.

AaronBallman

LGTM, thank you for this cleanup!

bader · 2021-04-09T12:04:30Z

/summary:run

Still need to do the code gen, but that is broken at least until the rebase on intel#3475, which is awaiting merge.

* upstream/sycl: (39 commits) [CI] Switch to default clang-format version. (intel#3540) [Driver][NFC] Cleanup some option setting for SYCL offload (intel#3542) [GitHub Actions] Update main branch sync schedule [SYCL][NFC] Fix potential namespace conflicts with PSTL in tuple.hpp (intel#3541) [SYCL] Bump sycl library minor version (intel#3538) [SYCL][CUDA] Implemented cuda_piextUSMEnqueueMemAdvise (intel#3365) [SYCL][FPGA] Add mutual diagnostic of max_concurrency attribute in conjunction of disable_loop_pipelining attribute (intel#3512) [SYCL] [MATRIX] Enable joint_matrix_load, joint_matrix_store, and joint_matrix_mad for AMX (intel#3503) [ESIMD] Skip rewriting functions used through function pointers (intel#3527) [SYCL] Fix address space for spec constants buffer (intel#3521) [SYCL] Correct the tablegen for checking mutually exclusive stmt attrs (intel#3519) [SYCL][PI][L0][NFC] Refactor setting of LastCommandEvent (intel#3528) [SYCL] Fix group local memory sharing issue (intel#3489) [SYCL][NFC] Fix post-commit failure (intel#3532) [SYCL][Doc] Remove extension mechanism (intel#3526) [SYCL] Move sycl.hpp in install directory and adjust driver to match (intel#3523) [SYCL][ESIMD] Update ESIMD docs to address recent user comments: (intel#3516) [NFCI][SYCL] Correct -fdeclare-spirv-builtins to use marshalling (intel#3515) [SYCL] Rework MarkDevice and children (intel#3475) [SYCL] Fix StringLiteral Ctor issue from intel#3504. (intel#3520) ...

[SYCL] Rework MarkDevice and children

1880471

We identified during a previous review that collectKernelAttributes and the MarkDeviceFunction object are in a situation that is very difficult to maintain, so this attempts to fix that by making the ownership model clear.

Erich Keane added 2 commits April 2, 2021 08:08

Clang-format fixes

dc10304

Add test to ensure constexpr-if inactive branches work

c51a88c

Remove copy paste comment error

aa44fd9

erichkeane commented Apr 2, 2021

View reviewed changes

clang/lib/Sema/SemaSYCL.cpp Show resolved Hide resolved

erichkeane commented Apr 2, 2021

View reviewed changes

Erich Keane added 3 commits April 2, 2021 10:25

Remove some 'proven' todos, combine attribute collection checks

9efd1f4

Merge branch 'sycl' into mark_device_rewrite

91eabf3

document the new types

efb756e

erichkeane marked this pull request as ready for review April 6, 2021 13:40

erichkeane requested review from AaronBallman, elizabethandrews and premanandrao as code owners April 6, 2021 13:40

erichkeane mentioned this pull request Apr 7, 2021

[SYCL] Report Compiling error for recursion in SYCL kernel #3390

Merged

premanandrao reviewed Apr 7, 2021

View reviewed changes

Apply Prem's suggestions

d2e3531

AaronBallman reviewed Apr 7, 2021

View reviewed changes

Erich Keane added 2 commits April 7, 2021 13:36

A bunch of fixups as suggested by Aaron

569250b

Clang-format fix

99747fd

AaronBallman reviewed Apr 8, 2021

View reviewed changes

Add more const, mess with sycl-simd in switch

bda527d

AaronBallman previously approved these changes Apr 8, 2021

View reviewed changes

Merge branch 'sycl' into mark_device_rewrite

ca07344

erichkeane dismissed AaronBallman’s stale review via ca07344 April 8, 2021 18:01

Remove newline inadvertently added during the 'merge'

83444aa

Last time I use the web app for this :)

premanandrao approved these changes Apr 8, 2021

View reviewed changes

AaronBallman approved these changes Apr 8, 2021

View reviewed changes

erichkeane pushed a commit to erichkeane/llvm that referenced this pull request Apr 9, 2021

Allow string parameters, finish up the sema tests.

5036d3a

Still need to do the code gen, but that is broken at least until the rebase on intel#3475, which is awaiting merge.

bader merged commit 551f528 into intel:sycl Apr 9, 2021

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL] Rework MarkDevice and children #3475

[SYCL] Rework MarkDevice and children #3475

erichkeane commented Apr 2, 2021

erichkeane commented Apr 2, 2021

erichkeane commented Apr 2, 2021

erichkeane Apr 2, 2021

erichkeane Apr 2, 2021

erichkeane Apr 2, 2021

erichkeane Apr 2, 2021

erichkeane Apr 2, 2021

erichkeane Apr 2, 2021

erichkeane Apr 2, 2021

erichkeane Apr 2, 2021

erichkeane Apr 2, 2021

erichkeane commented Apr 6, 2021

premanandrao Apr 6, 2021

erichkeane Apr 7, 2021

premanandrao Apr 7, 2021

premanandrao Apr 7, 2021

erichkeane Apr 7, 2021

AaronBallman left a comment

erichkeane commented Apr 7, 2021 •

edited

Loading

AaronBallman Apr 8, 2021

erichkeane Apr 8, 2021

AaronBallman Apr 8, 2021

erichkeane Apr 8, 2021

AaronBallman Apr 8, 2021

smanna12 Apr 8, 2021 •

edited

Loading

AaronBallman Apr 8, 2021

AaronBallman Apr 8, 2021

erichkeane Apr 8, 2021

AaronBallman left a comment

erichkeane commented Apr 8, 2021

premanandrao left a comment

AaronBallman left a comment

bader commented Apr 9, 2021

		@@ -6,13 +6,11 @@

		sycl::queue q;

		// expected-note@+1{{function implemented using recursion declared here}}


		[[intel::max_work_group_size(1, 1, 1)]] // expected-note {{conflicting attribute is here}}
		[[intel::max_work_group_size(1, 1, 1)]] // expected-note 3 {{conflicting attribute is here}}

[SYCL] Rework MarkDevice and children #3475

[SYCL] Rework MarkDevice and children #3475

Conversation

erichkeane commented Apr 2, 2021

erichkeane commented Apr 2, 2021

erichkeane commented Apr 2, 2021

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

erichkeane commented Apr 6, 2021

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

AaronBallman left a comment

Choose a reason for hiding this comment

erichkeane commented Apr 7, 2021 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

smanna12 Apr 8, 2021 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

AaronBallman left a comment

Choose a reason for hiding this comment

erichkeane commented Apr 8, 2021

premanandrao left a comment

Choose a reason for hiding this comment

AaronBallman left a comment

Choose a reason for hiding this comment

bader commented Apr 9, 2021

erichkeane commented Apr 7, 2021 •

edited

Loading

smanna12 Apr 8, 2021 •

edited

Loading