[SYCL] Add indeterminate constructor to work group memory interface #16003

lbushi25 · 2024-11-06T17:08:45Z

Implement the changes made to the work group memory spec in this PR: #15933.
Essentially, the default constructor is removed from the spec and instead a constructor that takes an indeterminate argument is made available in order to emphasize that any operation on such a work group memory object is undefined behavior except for assigning to it another work group memory object. Because the frontend needs special types to have a default constructor, we make it private instead and when its time for the frontend to call it, simply override the access specifier temporarily.

Pennycook · 2024-11-07T14:42:41Z

sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp

+// Frontend requires special types to have a default constructor in device
+// compilation mode in order to have a uniform way of initializing an object of
+// special type to then call the __init method on it. This is purely an
+// implementation detail and not part of the spec.
+#ifdef __SYCL_DEVICE_ONLY__
  work_group_memory() = default;
+#endif


I don't think this is legal. Doesn't it violate ODR?

This is observable to the user, so isn't just an implementation detail. Exposing a public default constructor allows a user to call work_group_memory() in device code, and also causes std::is_default_constructible_v<work_group_memory> to produce different values on host and device (which could break all sorts of things).

One fix would be to make this a protected function instead, and then create a factory function in the detail:: namespace that creates an uninitialized work_group_memory and calls __init on it.

This compiles fine. Exposing it to the user in device code is a problem, but unfortunately the way our frontend is set-up it requires special types to have a default constructor so that it will initialize them using that constructor in device code and then call __init on it. I did explore other avenues but the frontend team, @elizabethandrews , really don't like to have different handlings for different special types. Presumably, this sets a bad precedent for the future where other special types will also have their own hacks like this and the code will be non-uniform and harder to reason about and in my opinion its a totally valid concern.
Surprisingly, another class which does not have a default constructor as per the spec but seems to add one in device compilation mode is sycl::stream.( see line 862 in stream.hpp)

I didn't mean to imply that the code wouldn't compile. I think it's invalid because it results in multiple definitions of the same class (one for the host, one for the device), because it's not what the specification says, and because it could break user code in surprising ways.

I didn't realize we were already doing this for sycl::stream. I don't think we should be doing that, either.

Since the function call is only required by the front-end, can we just mark these default constructors private, and then ask the front-end to ignore the access specifier and call it anyway?

Tagging @gmlueck for awareness.

We could probably use some knowledge input from @intel/dpcpp-cfe-reviewers.

Can we do what @Pennycook suggested and have the Semantic analyzer ignore access specifiers just for the duration of this default initialization for special types?

This would allow us to mark these default constructors private for classes that do not have them in the spec so that users dont accidentally call them. This would be something that would apply to all special types, so we would preserve uniformity.

I didn't mean to imply that the code wouldn't compile. I think it's invalid because it results in multiple definitions of the same class (one for the host, one for the device), because it's not what the specification says, and because it could break user code in surprising ways.

I didn't realize we were already doing this for sycl::stream. I don't think we should be doing that, either.

Since the function call is only required by the front-end, can we just mark these default constructors private, and then ask the front-end to ignore the access specifier and call it anyway?

Tagging @gmlueck for awareness.

I made the requested changes.
Let's see if this looks ok to @intel/dpcpp-cfe-reviewers .

Sorry, our comments crossed. I'm not opposed to your option (2) if that seems less hacky. If we do that, can we put the factory function in the detail namespace?

I think so? I'm just not sure the effort required for this off the top of my head. As Tom mentioned, it may just not be worth it at this point. We could just go with option 1 adding a comment along the lines of "This is a temporary hack" and focus on designing whatever we upstream.

Either (1) or (2) is fine with me. I'm happy as long as we don't expose public APIs to the user which are not official public APIs.

Personally, I don't see (1) as any more of a hack than the use of the __init function as a whole. It is the compiler making the decision to use some member function(s) of the type, by instruction of the directives we give it. Luckily it isn't a user-facing feature, so we only have ourselves to blame if we mess it up.

Note also that the __init() member function can be - and in most classes is - declared private. @lbushi25 - Could we also move that to private while we are at it?

@tahonermann & @elizabethandrews - SYCLSpecialClassDocs in clang/include/clang/Basic/AttrDocs.td does not mention the default ctor requirement, though it has been a requirement for as long as I can remember. Should we amend it and, if we go with (1) here, note the access specifier relaxation?

I don't think we need to update the documentation. Comments in the code should suffice since it is an incorrect implementation assumption which we need to change eventually.

I agree with @elizabethandrews. We'll need to revisit the docs with whatever design we publish upstream anyway.

…tor access specifier

Pennycook

The changes in e531b05 address my concerns about the implementation of work_group_memory following the specification. I defer to @intel/dpcpp-cfe-reviewers for an assessment of whether a factory function or ignoring the access specifiers is an acceptable solution.

Fznamznon · 2024-11-12T13:52:57Z

clang/lib/Sema/SemaSYCL.cpp

    addFieldInit(FD, Ty, std::nullopt,
                 InitializationKind::CreateDefault(KernelCallerSrcLoc));
-
+    DefaultConstructor->setAccess(DefaultConstructorAccess);


Does it actually complain during kernel construction if you don't set access?
If yes, does it complain when it should after this change?

Does it actually complain during kernel construction if you don't set access? If yes, does it complain when it should after this change?

During the construction of the SYCL Kernel Body, it complains that the constructor has an incorrect number of arguments because it doesn't see a public default constructor. After this change, it doesn't complain anymore during kernel construction but it does complain when using the default constructor in any other place for example if the user uses the default constructor, they will get a compilation error because its not public to them. Its also an approach that works for any special type so we don't have to change our handling based on the special type.

Does it actually complain during kernel construction if you don't set access? If yes, does it complain when it should after this change?

This line in particular that you have reviewed sets the access back to its original value so that we don't accidentally leave the default constructor visible when it shouldn't be. Frankly, even without this line it seemed like the user could still not call the default constructor if it was defined private, but I chose to err on the safe side.

I see. I suppose I would also prefer the safe side.
Shouldn't you be doing the same for each variant of handleSpecialType?

I see. I suppose I would also prefer the safe side. Shouldn't you be doing the same for each variant of handleSpecialType?

Good question. This change is only relevant for the semantic action of the SyclKernelBodyCreator class which has 3 members named handleSpecialType. One is for kernel handler which has its own thing going on so I don't think it needs these changes, the other is the one I've changed here and the remaining one is the one taking a CXXBaseSpecifier. I suppose I should add it to this one too but I forgot, is that what you were asking about?

I suppose I should add it to this one too but I forgot, is that what you were asking about?

Right. Passing precommit also suggests lack of corresponding tests. I suppose a test should contain work group memory as a base to exercise the code.

Ok. I will add such a test.

I do not have another solution but this just feels very hacky and I really doubt community clang will accept this. This is out of the scope of the PR but the design of how we handle initialization of special types probably needs to change if we cannot use a default constructor. @tahonermann please weigh in here.

lbushi25 · 2024-11-12T22:19:04Z

The changes in e531b05 address my concerns about the implementation of work_group_memory following the specification. I defer to @intel/dpcpp-cfe-reviewers for an assessment of whether a factory function or ignoring the access specifiers is an acceptable solution.

@Pennycook @gmlueck
Based on @tahonermann's feedback, I've reverted back to following the precedent set by sycl::stream, that is, making the default constructor public. I've also opened an issue about this and put a link in a comment. I realize this is deviating from the spec but based on the discussion here, this seems to be the consensus solution at least in the short term. We either have the option to accept this until the frontend handling changes or we can revert the substitution of the default constructor with the one taking an indeterminate as an argument.

This reverts commit 21c27fa.

This reverts commit 8c94bfc.

…structor" This reverts commit eb31ca0.

This reverts commit 54abcd0.

… for inheriting from work group memory

lbushi25 · 2024-11-15T02:29:16Z

@intel/llvm-gatekeepers This is ready to merge.

Add indeterminate constructor to work group memory interface

e574aa7

lbushi25 requested a review from a team as a code owner November 6, 2024 17:08

lbushi25 requested a review from bso-intel November 6, 2024 17:08

lbushi25 had a problem deploying to WindowsCILock November 6, 2024 17:10 — with GitHub Actions Error

Update tests to use indeterminate constructor instead of default one

8164ff2

lbushi25 had a problem deploying to WindowsCILock November 6, 2024 17:12 — with GitHub Actions Error

lbushi25 added 3 commits November 6, 2024 09:13

Add indeterminate constructor to work group memory interface

543aa6f

Fix typo

a34b463

Fix compiler error

65a8e23

lbushi25 had a problem deploying to WindowsCILock November 6, 2024 17:18 — with GitHub Actions Error

Formatting changes

aac82e2

lbushi25 had a problem deploying to WindowsCILock November 6, 2024 17:28 — with GitHub Actions Error

Fix unused variable warning

d6709ed

lbushi25 temporarily deployed to WindowsCILock November 6, 2024 17:41 — with GitHub Actions Inactive

lbushi25 temporarily deployed to WindowsCILock November 6, 2024 18:27 — with GitHub Actions Inactive

lbushi25 requested a review from steffenlarsen November 6, 2024 19:15

lbushi25 mentioned this pull request Nov 7, 2024

[SYCL] Write tests following the test plan for the work group memory extension #15928

Merged

Pennycook requested changes Nov 7, 2024

View reviewed changes

bso-intel approved these changes Nov 7, 2024

View reviewed changes

Modify handling of SYCL special types to account for default construc…

847163d

…tor access specifier

lbushi25 requested a review from a team as a code owner November 8, 2024 22:40

lbushi25 had a problem deploying to WindowsCILock November 8, 2024 22:40 — with GitHub Actions Error

Modify handling of SYCL special types to account for default construc…

f551ba4

…tor access specifier

lbushi25 had a problem deploying to WindowsCILock November 8, 2024 22:42 — with GitHub Actions Error

Modify handling of SYCL special types to account for default construc…

172d717

…tor access specifier

lbushi25 had a problem deploying to WindowsCILock November 8, 2024 22:43 — with GitHub Actions Error

Make default constructor private

e531b05

lbushi25 had a problem deploying to WindowsCILock November 8, 2024 23:21 — with GitHub Actions Error

Formatting changes

52bf8a4

lbushi25 temporarily deployed to WindowsCILock November 8, 2024 23:27 — with GitHub Actions Inactive

lbushi25 had a problem deploying to WindowsCILock November 11, 2024 17:58 — with GitHub Actions Error

Use std::find_if instead of std::for_each to find default constructor

eb31ca0

lbushi25 temporarily deployed to WindowsCILock November 11, 2024 18:58 — with GitHub Actions Inactive

lbushi25 temporarily deployed to WindowsCILock November 11, 2024 19:49 — with GitHub Actions Inactive

Pennycook approved these changes Nov 12, 2024

View reviewed changes

Fznamznon reviewed Nov 12, 2024

View reviewed changes

Revert back changes to default constructor

54abcd0

lbushi25 had a problem deploying to WindowsCILock November 12, 2024 22:09 — with GitHub Actions Error

Fix formatting errors

21c27fa

lbushi25 closed this Nov 12, 2024

lbushi25 reopened this Nov 12, 2024

lbushi25 temporarily deployed to WindowsCILock November 12, 2024 22:19 — with GitHub Actions Inactive

lbushi25 temporarily deployed to WindowsCILock November 12, 2024 22:46 — with GitHub Actions Inactive

lbushi25 added 6 commits November 14, 2024 06:45

Revert "Fix formatting errors"

8c94bfc

This reverts commit 21c27fa.

Revert "Revert "Fix formatting errors""

d5c79ec

This reverts commit 8c94bfc.

Revert "Use std::find_if instead of std::for_each to find default con…

0e3ae8a

…structor" This reverts commit eb31ca0.

Revert "Revert back changes to default constructor"

be386bc

This reverts commit 54abcd0.

Bypass default constructor access specifier in frontend and add tests…

17639ae

… for inheriting from work group memory

Add TODO to revisit the handling of special types

6cc99d3

lbushi25 had a problem deploying to WindowsCILock November 14, 2024 17:16 — with GitHub Actions Error

Formatting changes

90bceb0

lbushi25 had a problem deploying to WindowsCILock November 14, 2024 17:32 — with GitHub Actions Failure

elizabethandrews approved these changes Nov 14, 2024

View reviewed changes

lbushi25 temporarily deployed to WindowsCILock November 14, 2024 19:38 — with GitHub Actions Inactive

Fix regex in filecheck pattern

3ea9103

lbushi25 temporarily deployed to WindowsCILock November 14, 2024 21:52 — with GitHub Actions Inactive

lbushi25 temporarily deployed to WindowsCILock November 14, 2024 22:45 — with GitHub Actions Inactive

martygrant merged commit d76b552 into intel:sycl Nov 15, 2024
13 checks passed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL] Add indeterminate constructor to work group memory interface #16003

[SYCL] Add indeterminate constructor to work group memory interface #16003

lbushi25 commented Nov 6, 2024 •

edited

Loading

Pennycook Nov 7, 2024

lbushi25 Nov 7, 2024 •

edited

Loading

Pennycook Nov 7, 2024

lbushi25 Nov 8, 2024 •

edited

Loading

lbushi25 Nov 8, 2024 •

edited

Loading

elizabethandrews Nov 13, 2024

gmlueck Nov 13, 2024

steffenlarsen Nov 14, 2024 •

edited

Loading

elizabethandrews Nov 14, 2024

tahonermann Nov 14, 2024

Pennycook left a comment

Fznamznon Nov 12, 2024

lbushi25 Nov 12, 2024 •

edited

Loading

lbushi25 Nov 12, 2024

Fznamznon Nov 12, 2024

lbushi25 Nov 12, 2024 •

edited

Loading

Fznamznon Nov 12, 2024

lbushi25 Nov 12, 2024

elizabethandrews Nov 12, 2024

lbushi25 commented Nov 12, 2024 •

edited

Loading

lbushi25 commented Nov 15, 2024

[SYCL] Add indeterminate constructor to work group memory interface #16003

[SYCL] Add indeterminate constructor to work group memory interface #16003

Conversation

lbushi25 commented Nov 6, 2024 • edited Loading

Choose a reason for hiding this comment

lbushi25 Nov 7, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

lbushi25 Nov 8, 2024 • edited Loading

Choose a reason for hiding this comment

lbushi25 Nov 8, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

steffenlarsen Nov 14, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Pennycook left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

lbushi25 Nov 12, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

lbushi25 Nov 12, 2024 • 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

lbushi25 commented Nov 12, 2024 • edited Loading

lbushi25 commented Nov 15, 2024

lbushi25 commented Nov 6, 2024 •

edited

Loading

lbushi25 Nov 7, 2024 •

edited

Loading

lbushi25 Nov 8, 2024 •

edited

Loading

lbushi25 Nov 8, 2024 •

edited

Loading

steffenlarsen Nov 14, 2024 •

edited

Loading

lbushi25 Nov 12, 2024 •

edited

Loading

lbushi25 Nov 12, 2024 •

edited

Loading

lbushi25 commented Nov 12, 2024 •

edited

Loading