Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Clarify default constructed accessor and "require" #434

Merged
merged 4 commits into from
Jul 20, 2023
Merged

Clarify default constructed accessor and "require" #434

Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
14a3729
Clarify default constructed accessor and "require"
gmlueck Jun 23, 2023
a9f7355
Clarify wording of default-constructed accessors
gmlueck Jun 29, 2023
bc8e2b6
Review feedback from Gordon
gmlueck Jul 18, 2023
d6654c0
Further clarify that zero-sized accessors legal
gmlueck Jul 20, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
34 changes: 18 additions & 16 deletions adoc/chapters/programming_interface.adoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6671,10 +6671,8 @@ accessor()
* All size queries return [code]#0#.
* The return values of [code]#get_pointer()# and [code]#get_multi_ptr()# are
unspecified.
* Trying to access the underlying memory is undefined behavior.

A default constructed accessor can be passed to a <<sycl-kernel-function>>
but it is not valid to register it with the command group handler.
* A default constructed accessor can be passed to a <<sycl-kernel-function>>,
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
* A default constructed accessor can be passed to a <<sycl-kernel-function>>,
* A default constructed [code]#accessor# can be passed to a <<sycl-kernel-function>>,

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If you feel strongly about this, I don't mind adding the code-font. However, we do not seem to use code font for every occurrence of class names, especially the common ones like "queue", "buffer", and "accessor". For example, see the first sentence of this function description:

Constructs an empty accessor which fulfills the following post-conditions

(There is no code-font for "accessor" here.)

I've actually been trying to tone-down the amount of code font in cases like this because I think all the font changes are a little distracting. I don't really have a clear rule for when to use code font, but using it for every occurrence of a class name seems a bit too much.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't feel strongly about this, so I'm happy leaving this as-is, maybe it's worth having a separate discussion about general conventions for formatting?

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In the past I think that I made some choice depending on whether it was either very important that it is about the sycl:: object or more vaguely the concept. Then we have the discussion about whether we should point to the glossary or not... :-)

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In the past I think that I made some choice depending on whether it was either very important that it is about the sycl:: object or more vaguely the concept.

Yes, this make sense.

Then we have the discussion about whether we should point to the glossary or not... :-)

Agreed that this is another point to consider. When writing new text, I've been trying to avoid excessive links to the glossary entries too. For example, the first occurrence of a glossary entry in a paragraph (or group of related paragraphs) can link to the glossary entry, but other occurrences of the glossary term in that same paragraph (or group of related paragraphs) does not need to be a link.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree, I think this is generally a good convention to have. Maybe we should open a new issue for discussing this.

but attempting to access the underlying memory produces undefined behavior.
gmlueck marked this conversation as resolved.
Show resolved Hide resolved
--

a@
Expand Down Expand Up @@ -6940,8 +6938,8 @@ a@
----
bool is_placeholder() const
----
a@ Returns [code]#true# if the accessor was constructed as a placeholder.
Otherwise returns [code]#false#.
a@ Returns [code]#true# if the accessor is a placeholder. Otherwise returns
[code]#false#.

a@
[source]
Expand Down Expand Up @@ -8195,7 +8193,9 @@ local_accessor()
* All size queries return [code]#0#.
* The return values of [code]#get_pointer()# and [code]#get_multi_ptr()# are
unspecified.
* Trying to access the underlying memory is undefined behavior.
* A default constructed local accessor can be passed to a
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
* A default constructed local accessor can be passed to a
* A default constructed [code]#local_accessor# can be passed to a

<<sycl-kernel-function>>, but attempting to access the underlying memory
produces undefined behavior.
gmlueck marked this conversation as resolved.
Show resolved Hide resolved
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This does somewhat create a foot gun, I think it's an acceptable side effect of allowing default constructed accessors to be passed to a kernel, and I don't think we can provide any stronger wording than saying it's UB without impacting implementations, however, we could potentially require that the value of the pointer for a zero-sized accessor is nullptr, to allow users to check for this case within kernel code (without having to query the size).

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We discussed the question about nullptr already in #232, see this comment.

My opinion is pretty much the same as before. I think we should treat all zero-sized accessors the same in this regard (either say they all return nullptr or say they all return an unspecified value). It is not hard to guarantee the nullptr, but it does require extra logic in the implementation. For example, when constructing a ranged accessor, the implementation would need a special check to see if the range was zero. Guaranteeing a nullptr does not seem that useful to me. Applications still need to check that the index is within range of the accessor. If applications follow this rule, then they should never attempt to dereference the pointer for a zero-sized accessor. If applications do want to check for this special case, it seems easy enough to call accessor::empty.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry, I missed that discussion, I think that's a fair point, though I think there is still a use case for this, though admittedly pretty niche. If you had a kernel where an accessor is conditionally used, being left as default constructed if it's not, and when it is used each work-item only accesses its own index, and it's guaranteed to be the same size as the nd-range, then you may want to check if the accessor is zero-sized, without having to check the size (and potentially removing the implementation's ability to optimize away that kernel argument).

Though saying this I think while size() does require an argument to be passed to the kernel to provide this value, empty could be implemented based on whether the pointer is nullptr. So I think this is a problem that can be handled by the implementation, and so doesn't need to be in the specification.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Rather than holding up this PR, lets create an issue to discuss whether a zero-sized accessor should be required to return nullptr from get_pointer. I think this is an orthogonal discussion from the clarification I'm making in this PR.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree, I don't think we should hold up the PR for this, I'm comfortable that this is implementable with the current wording, so I think we can leave this as-is, also happy to discuss this further in a separate ticket.

--

a@
Expand Down Expand Up @@ -13162,6 +13162,11 @@ this does not happen when creating a [keyword]#placeholder# accessor. In order
to create a *requirement* for a [keyword]#placeholder# accessor, code
must call the [code]#handler::require()# member function.

Note that the default constructed [code]#accessor# is not a placeholder, so it
may be passed to a kernel without calling [code]#handler::require()#. However,
gmlueck marked this conversation as resolved.
Show resolved Hide resolved
this accessor also has no underlying memory object, so such an accessor does
not create any *requirement* for the command group.
gmlueck marked this conversation as resolved.
Show resolved Hide resolved

SYCL events may also be used to create requirements for a <<command-group>>.
Such requirements state that the actions represented by the events must
complete before the <<command-group>> may execute. Such requirements
Expand All @@ -13180,15 +13185,12 @@ template <typename DataT, int Dimensions, access_mode AccessMode,
void require(
accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder> acc)
----
a@ Requires access to the memory object associated with the accessor.

The <<command-group>> now has a *requirement* to gain access
to the given memory object before executing the kernel.
If the accessor has already been registered with the <<command-group>>,
calling this function has no effect.

Throws [code]#exception# with the [code]#errc::invalid# error code
if [code]#(acc.empty() == true)#.
a@ Calling this function has no effect unless [code]#acc# is a placeholder
accessor. When [code]#acc# is a placeholder accessor, this function
adds a *requirement* to the handler's <<command-group>> for the memory
object represented by [code]#acc#. If the accessor has already been
registered with the <<command-group>>, calling this function has no
effect.

a@
[source]
Expand Down