Clarify default constructed accessor and "require" #434

gmlueck · 2023-06-23T18:23:13Z

The previous wording caused some confusion about whether a default constructed accessor is a placeholder. In general, the accessor constructors that do not take a handler parameter construct placeholder accessors, which could give the impression that the default constructed accessor is a placeholder. Clarify this by specifically stating that the default constructed accessor is not a placeholder and that it may be passed as a kernel parameter without calling handler::require.

The default constructed accessor and the default constructed local_accessor are similar because they can both be passed as kernel parameters and neither allocates any memory for use in the kernel. Adjust the wording for these two constructors to be more similar.

Remove wording from handler::require saying that it is an error to call this function for a zero sized accessor. We allow non-placeholder accessor that are zero sized, so it makes no sense to disallow this for placeholder accessors.

Remove wording from the default accessor constructor saying that it is an error to pass such an accessor to handler::require. We allow other non-placeholder accessors to be passed to handler::require, so it makes sense to allow this for the default constructed accessor too.

Clarify the wording for accessor::is_placeholder. The previous wording "if the accessor was constructed as a placeholder" caused confusion. What if an accessor was constructed as a placeholder and then later assigned to a non-placeholder accessor (or swapped to a non-placeholder accessor)? Presumably, we want is_placeholder to return false in such a case. The new wording removes this confusing phrasing.

Closes #408.

The previous wording caused some confusion about whether a default constructed `accessor` is a placeholder. In general, the accessor constructors that do not take a `handler` parameter construct placeholder accessors, which could give the impression that the default constructed accessor is a placeholder. Clarify this by specifically stating that the default constructed accessor is not a placeholder and that it may be passed as a kernel parameter without calling `handler::require`. The default constructed `accessor` and the default constructed `local_accessor` are similar because they can both be passed as kernel parameters and neither allocates any memory for use in the kernel. Adjust the wording for these two constructors to be more similar. Remove wording from `handler::require` saying that it is an error to call this function for a zero sized accessor. We allow non-placeholder accessor that are zero sized, so it makes no sense to disallow this for placeholder accessors. Remove wording from the default `accessor` constructor saying that it is an error to pass such an accessor to `handler::require`. We allow other non-placeholder accessors to be passed to `handler::require`, so it makes sense to allow this for the default constructed accessor too. Clarify the wording for `accessor::is_placeholder`. The previous wording "if the accessor was constructed as a placeholder" caused confusion. What if an accessor was constructed as a placeholder and then later assigned to a non-placeholder accessor (or swapped to a non-placeholder accessor)? Presumably, we want `is_placeholder` to return `false` in such a case. The new wording removes this confusing phrasing. Closes KhronosGroup#408.

maarquitos14 · 2023-06-26T11:24:37Z

Looks good to me. Only question is, requiring an empty accessor should create a dependency? Say we have an empty buffer, and we create a write accessor and a read accessor to it, should that create a write/read dependency or should we ignore it because it's empty?

gmlueck · 2023-06-26T14:30:45Z

Looks good to me. Only question is, requiring an empty accessor should create a dependency?

There are several ways to create an empty accessor, and I think the answer to your question depends on the way the accessor was created:

If the accessor was created from a ranged accessor with a zero-sized range, I think the spec is clear that there is a dependency on the entire buffer. Section 4.7.6.8 states that a ranged accessor creates a dependency on the entire buffer, and there is no exception for a zero-sized range. Therefore, this applies even to zero-size ranged accessors.
If the accessor was created from a zero-sized buffer, I think it still creates a dependency. Section 3.8.1 specifies how accessors create dependencies on buffers, and there is no exception here for zero-sized buffers. Therefore, I think there are dependencies even when the buffer has zero size.
If the accessor was created from the default constructor, there is no buffer at all, so it is impossible to create a dependency. This PR clarifies that default constructed accessors do not create requisites (i.e. dependencies). See the change in this PR to the section titled "SYCL functions for adding requirements".

keryell

Great subtle clarifications.

adoc/chapters/programming_interface.adoc

Avoid using the phrase "the underlying memory" for a default constructed accessor because this implies that there is some underlying memory.

keryell

Thanks, this is clearer.

gmlueck · 2023-07-13T11:39:38Z

Ping to the group .... looking for another reviewer on this.

…sed to a command (#10110) Current mechanism for checking whether an exception should be thrown if a placeholder accessor was passed to a command was not working properly under certain conditions, either by not throwing when required, or throwing when unrequired. Additionally, following clarifications in KhronosGroup/SYCL-Docs#434, this PR makes default-constructed accessors to stop being placeholders. --------- Signed-off-by: Maronas, Marcos <marcos.maronas@intel.com>

AerialMantis

A few suggestions, but overall these changes look good to me.

One comment I have is I think we should also add wording to https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:data.access.and.storage to establish that zero-sized accessors are permitted, and perhaps add more general wording to the accessor section to describe the concept of a zero-sized accessor and the ways in which it can be created (i.e. zero-sized buffer, ranged accessor with size of zero or default constructor).

AerialMantis · 2023-07-18T07:01:53Z

adoc/chapters/programming_interface.adoc

@@ -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


Suggested change

* A default constructed local accessor can be passed to a

* A default constructed [code]#local_accessor# can be passed to a

adoc/chapters/programming_interface.adoc

AerialMantis · 2023-07-18T09:38:39Z

adoc/chapters/programming_interface.adoc

-
-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>>,


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>>,

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.

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?

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... :-)

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.

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

AerialMantis · 2023-07-18T09:50:00Z

adoc/chapters/programming_interface.adoc

-  * Trying to access the underlying memory is undefined behavior.
+  * A default constructed local accessor can be passed to a
+    <<sycl-kernel-function>>, but attempting to access data elements from it
+    produces undefined behavior.


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).

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.

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.

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.

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.

adoc/chapters/programming_interface.adoc

gmlueck · 2023-07-18T12:53:30Z

One comment I have is I think we should also add wording to https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:data.access.and.storage to establish that zero-sized accessors are permitted, and perhaps add more general wording to the accessor section to describe the concept of a zero-sized accessor and the ways in which it can be created (i.e. zero-sized buffer, ranged accessor with size of zero or default constructor).

How about expanding this existing paragraph in the introduction:

Zero-sized buffers are permitted. In this case, the buffer still handles ownership as normal, but does not need to store the zero-sized data itself.

The updated wording could be like this:

Zero-sized buffers and accessors are permitted, but attempting to access data within them produces undefined behavior, similar to dereferencing a null pointer in C++. Note that zero-sized accessors can be created in several ways: by creating an accessor from a zero-sized buffer, by creating an accessor with a zero-sized buffer sub-range, or by creating an accessor with its default constructor.

I intentionally dropped the last sentence here ("In this case, the buffer still handles ownership as normal, but does not need to store the zero-sized data itself.") because I don't think this specifies anything useful. What does it mean to say that the buffer "owns" zero-sized data? What do we mean by saying the ownership is "normal"? Is it illegal for an implementation to have a special case when the buffer has zero-size? In my opinion, this sentence does not add any clarity.

Use glossary term "sycl-kernel-function" and clarify that dereferencing a zero-sized accessor is UB.

AerialMantis · 2023-07-20T14:06:27Z

One comment I have is I think we should also add wording to https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:data.access.and.storage to establish that zero-sized accessors are permitted, and perhaps add more general wording to the accessor section to describe the concept of a zero-sized accessor and the ways in which it can be created (i.e. zero-sized buffer, ranged accessor with size of zero or default constructor).

How about expanding this existing paragraph in the introduction:

Zero-sized buffers are permitted. In this case, the buffer still handles ownership as normal, but does not need to store the zero-sized data itself.

The updated wording could be like this:

Zero-sized buffers and accessors are permitted, but attempting to access data within them produces undefined behavior, similar to dereferencing a null pointer in C++. Note that zero-sized accessors can be created in several ways: by creating an accessor from a zero-sized buffer, by creating an accessor with a zero-sized buffer sub-range, or by creating an accessor with its default constructor.

I intentionally dropped the last sentence here ("In this case, the buffer still handles ownership as normal, but does not need to store the zero-sized data itself.") because I don't think this specifies anything useful. What does it mean to say that the buffer "owns" zero-sized data? What do we mean by saying the ownership is "normal"? Is it illegal for an implementation to have a special case when the buffer has zero-size? In my opinion, this sentence does not add any clarity.

This wording sounds good to me, and I agree with the removal of the last sentence, if there's no data for the buffer to manage, then there's no ownership relationship to describe. Thanks.

Add a paragraph to the introduction of the section on buffers and accessors stating that zero-sized accessors are legal.

gmlueck · 2023-07-20T14:49:33Z

This wording sounds good to me, and I agree with the removal of the last sentence, if there's no data for the buffer to manage, then there's no ownership relationship to describe. Thanks.

Done in d6654c0

fraggamuffin · 2023-07-20T15:31:51Z

LGTM

Following the removal of the require() exception condition in KhronosGroup/SYCL-Docs#434 this commit removes the corresponding test. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Following the resolution of KhronosGroup/SYCL-Docs#408 with KhronosGroup/SYCL-Docs#434 DPC++ has been changed to not throw for empty accessors. As such, the test cases using check_zero_length_buffer_constructor can now be enabled. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Clarify default constructed accessor and "require". The previous wording caused some confusion about whether a default constructed accessor is a placeholder. In general, the accessor constructors that do not take a handler parameter construct placeholder accessors, which could give the impression that the default constructed accessor is a placeholder. Clarify this by specifically stating that the default constructed accessor is not a placeholder and that it may be passed as a kernel parameter without calling handler::require. The default constructed accessor and the default constructed local_accessor are similar because they can both be passed as kernel parameters and neither allocates any memory for use in the kernel. Adjust the wording for these two constructors to be more similar. Remove wording from handler::require saying that it is an error to call this function for a zero sized accessor. We allow non-placeholder accessor that are zero sized, so it makes no sense to disallow this for placeholder accessors. Remove wording from the default accessor constructor saying that it is an error to pass such an accessor to handler::require. We allow other non-placeholder accessors to be passed to handler::require, so it makes sense to allow this for the default constructed accessor too. Clarify the wording for accessor::is_placeholder. The previous wording "if the accessor was constructed as a placeholder" caused confusion. What if an accessor was constructed as a placeholder and then later assigned to a non-placeholder accessor (or swapped to a non-placeholder accessor)? Presumably, we want is_placeholder to return false in such a case. The new wording removes this confusing phrasing. Closes #408.

gmlueck mentioned this pull request Jun 23, 2023

Should accessors created from a command group handler throw if passed empty buffers? #408

Closed

maarquitos14 mentioned this pull request Jun 28, 2023

[SYCL] Fix mechanism to throw exception when placeholder accessor passed to a command intel/llvm#10110

Merged

keryell approved these changes Jun 28, 2023

View reviewed changes

adoc/chapters/programming_interface.adoc Outdated Show resolved Hide resolved

adoc/chapters/programming_interface.adoc Show resolved Hide resolved

Clarify wording of default-constructed accessors

a9f7355

Avoid using the phrase "the underlying memory" for a default constructed accessor because this implies that there is some underlying memory.

keryell approved these changes Jul 1, 2023

View reviewed changes

AerialMantis reviewed Jul 18, 2023

View reviewed changes

Review feedback from Gordon

bc8e2b6

Use glossary term "sycl-kernel-function" and clarify that dereferencing a zero-sized accessor is UB.

Further clarify that zero-sized accessors legal

d6654c0

Add a paragraph to the introduction of the section on buffers and accessors stating that zero-sized accessors are legal.

AerialMantis approved these changes Jul 20, 2023

View reviewed changes

keryell merged commit 418c24a into KhronosGroup:SYCL-2020/master Jul 20, 2023
1 check passed

gmlueck deleted the gmlueck/require-accessor branch July 20, 2023 16:33

steffenlarsen mentioned this pull request Aug 1, 2023

Remove require() exception test KhronosGroup/SYCL-CTS#770

Merged

steffenlarsen mentioned this pull request Aug 2, 2023

Enable check_zero_length_buffer_constructor for DPCPP KhronosGroup/SYCL-CTS#775

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Clarify default constructed accessor and "require" #434

Clarify default constructed accessor and "require" #434

gmlueck commented Jun 23, 2023

maarquitos14 commented Jun 26, 2023

gmlueck commented Jun 26, 2023

keryell left a comment

keryell left a comment

gmlueck commented Jul 13, 2023

AerialMantis left a comment

AerialMantis Jul 18, 2023

AerialMantis Jul 18, 2023

gmlueck Jul 18, 2023

AerialMantis Jul 20, 2023

keryell Jul 20, 2023

gmlueck Jul 20, 2023

AerialMantis Jul 20, 2023

AerialMantis Jul 18, 2023

gmlueck Jul 18, 2023

AerialMantis Jul 20, 2023

gmlueck Jul 20, 2023

AerialMantis Jul 20, 2023

gmlueck commented Jul 18, 2023

AerialMantis commented Jul 20, 2023

gmlueck commented Jul 20, 2023

fraggamuffin commented Jul 20, 2023

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

	* 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>>,

Clarify default constructed accessor and "require" #434

Clarify default constructed accessor and "require" #434

Conversation

gmlueck commented Jun 23, 2023

maarquitos14 commented Jun 26, 2023

gmlueck commented Jun 26, 2023

keryell left a comment

Choose a reason for hiding this comment

keryell left a comment

Choose a reason for hiding this comment

gmlueck commented Jul 13, 2023

AerialMantis left a 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

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

gmlueck commented Jul 18, 2023

AerialMantis commented Jul 20, 2023

gmlueck commented Jul 20, 2023

fraggamuffin commented Jul 20, 2023