[SYCL][Doc] Add draft forward progress extension #7598

Pennycook · 2022-11-30T21:05:32Z

This extension strengthens the forward progress guarantees provided by a SYCL implementation, and exposes the guarantees of particular implementations/backends/devices via new device queries and kernel properties.

Signed-off-by: John Pennycook [email protected]

This extension strengthens the forward progress guarantees provided by a SYCL implementation, and exposes the guarantees of particular implementations/backends/devices via new device queries and kernel properties. Signed-off-by: John Pennycook <[email protected]>

Pennycook · 2022-12-09T15:26:35Z

Note to self and future reviewers: the launch configuration queries as currently written are insufficient, because there may be dependencies between query results. For example, the maximum number of work-groups that can execute concurrently on an NVIDIA GPU depends on the work-group size (see the documentation for cudaOccupancyMaxActiveBlocksPerMultiprocessor).

In general, and looking across implementations, I think:

The maximum number of work-groups may depend on a user-supplied work-group size
The maximum work-group size may depend on a user-supplied number of sub-groups
The maximum number of sub-groups depends on the device
The maximum sub-group size depends on the kernel

Pennycook · 2023-03-13T18:12:00Z

@intel/dpcpp-specification-reviewers - I think this is ready for review, now.

Pennycook · 2023-06-20T15:18:02Z

@intel/dpcpp-specification-reviewers: Do you think this is ready to be merged? If I recall correctly, we reviewed the draft internally ahead of the IWOCL presentation. There's been no activity here for a few months.

The ability of a device to simultaneously execute work-groups may depend on the amount of dynamically allocated work-group local memory (e.g. from accessors). Signed-off-by: John Pennycook <[email protected]>

Pennycook · 2023-07-19T16:22:49Z

@bashbaug: I've just tried to add the local memory query that we talked about. Please let me know if this isn't the solution you expected.

AerialMantis

Looks great! Just a few suggestions.

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

AerialMantis · 2023-07-21T15:49:38Z

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

+----
+namespace ext::oneapi::experimental::sycl {
+
+enum class forward_progress_guarantee {


Since the values of this enum represent semantics within a fixed hierarchy, should we specify that they have set values, either explicitly in the synopsis or as an implementation requirement?

I don't think we should do this. P2300 also defines a forward_progress_guarantee enum, and does not specify values. I think we should stay aligned with that, at least for now.

I'm also worried that we might have to add additional progress guarantees at some point (either because SYCL needs them, or because C++ introduces them), and updating previously specified values might break user code.

That's a good point, having fixed values could complicate adding additional guarantees in the future, especially if they fit in between existing guarantees or if it's no longer a simpler linear hierarchy.

AerialMantis · 2023-07-21T15:50:07Z

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

+----
+namespace ext::oneapi::experimental::sycl {
+
+enum class execution_scope {


Similar to the comment above, since the values of this enum represent semantics within a fixed hierarchy, should we specify that they have set values?

I don't think we should view this as a fixed hierarchy. There may be implementations that want to add additional execution scopes in the middle of the existing hierarchy. For example, if it's possible to provide stronger forward progress guarantees for all work-groups executing on the same cluster of SMs on an NVIDIA GPU, it would make sense to add something like a cluster execution scope between work_group and root_group.

Aside: These sorts of concerns are why we split things into a coordination and target scope. Initially, we considered implicitly specifying a forward progress guarantee relative to the immediate parent (i.e. work-item forward progress guarantees would always be relative to the sub-group, and so on), but this would break user code unexpectedly if/when new layers were added to the hierarchy. By making both scopes explicit, we can add new layers to the hierarchy safely.

Yeah that's a good point, I agree it makes sense then to leave these without set values.

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

bashbaug

Thanks for writing this up! This is a complex topic and I think you've done a great job making it as approachable as possible.

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

- They must respect all other properties - They must respect usage of work-group local memory

- Require a minimum of 1 if a kernel is supported - Require a value of 0 if a kernel is not supported - Clarify that values respect all properties

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

Co-authored-by: Greg Lueck <[email protected]>

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

Co-authored-by: Greg Lueck <[email protected]>

Co-authored-by: Ben Ashbaugh <[email protected]>

This is in the ISO C++ definition, but replicating it here in a non-normative note may help readers understand the implications of the execution model.

Queries and properties relate to the intermediate threads of execution, and so they must also be part of the associated scope(s).

This hopefully clarifies what the return value of this query really represents, and how it might be used.

gmlueck

Thanks for the updates, @Pennycook. I think the meaning of the forward-progress properties are now more formally defined according to the execution model. I added a few more comments below, though.

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc

Each work-group must be the same size, and the sub-group size has already been determined (because the query is relative to a specific work-group size). The number of sub-groups must therefore be known.

sycl/doc/extensions/proposed/sycl_ext_oneapi_launch_queries.asciidoc

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

gmlueck

If we want the Param::return_type and Param::argument_types to be implementation details, I think it's better to leave them out of the spec entirely. The comments below show one way to do this, which aligns with the proposed wording in the core spec in KhronosGroup/SYCL-Docs#440.

sycl/doc/extensions/proposed/sycl_ext_oneapi_launch_queries.asciidoc

Co-authored-by: Greg Lueck <[email protected]>

Pennycook · 2023-08-01T17:40:19Z

If we want the Param::return_type and Param::argument_types to be implementation details, I think it's better to leave them out of the spec entirely. The comments below show one way to do this, which aligns with the proposed wording in the core spec in KhronosGroup/SYCL-Docs#440.

I prefer this too. I've made these changes in 36433dd and 1cdabef.

Note that Param::return_type is currently how SYCL 2020 defines its queries, and it's not very well defined there. It might be a good idea to swap it out for /* return-type */ at some point.

gmlueck · 2023-08-01T19:16:53Z

Note that Param::return_type is currently how SYCL 2020 defines its queries, and it's not very well defined there. It might be a good idea to swap it out for /* return-type */ at some point.

I created internal Khronos issue 663 to discuss this.

AerialMantis

Thanks for the updates @Pennycook, looks great! I've just added a couple of additonal comments.

sycl/doc/extensions/proposed/sycl_ext_oneapi_launch_queries.asciidoc

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

Pennycook · 2023-08-25T14:16:10Z

@intel/llvm-gatekeepers: This can be merged! 🥳

AlexeySachkov · 2024-02-28T09:51:56Z

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

+
+[source,c++]
+----
+namespace ext::oneapi::experimental::sycl {


Should this (and other) namespace definition be sycl::ext::oneapi::experimental instead? (sycl:: in the beginning and not at the end)

...Yes. I have no idea how this happened! Probably a bad vim search and replace. Good catch!

Pennycook added the spec extension All issues/PRs related to extensions specifications label Nov 30, 2022

Pennycook marked this pull request as ready for review March 13, 2023 18:11

Pennycook requested a review from a team as a code owner March 13, 2023 18:11

Allow kernel info descriptors to require arguments

6ffcacf

Pennycook mentioned this pull request May 10, 2023

[SYCL] Implement sycl_ext_oneapi_root_group #9396

Merged

Fix table formatting

da438f8

Pennycook requested a review from gmlueck May 17, 2023 13:55

Limit max_num_work_groups based on local memory

ccbdcf0

The ability of a device to simultaneously execute work-groups may depend on the amount of dynamically allocated work-group local memory (e.g. from accessors). Signed-off-by: John Pennycook <[email protected]>

Pennycook requested review from AerialMantis and bashbaug July 19, 2023 16:21

AerialMantis reviewed Jul 21, 2023

View reviewed changes

Pennycook added 2 commits July 21, 2023 14:36

Clarify that coordination scope > target scope

1c16d2f

Prevent invalid forward_progress_guarantees

cb8c546

bashbaug reviewed Jul 24, 2023

View reviewed changes

Pennycook added 5 commits July 24, 2023 09:03

Clarify that work-group local memory is dynamic

7319e89

Leave door open to broader scopes than root-group

5f382df

Clarify how launch queries are calculated

b09639d

- They must respect all other properties - They must respect usage of work-group local memory

Clarify the return values of launch queries

da6d477

- Require a minimum of 1 if a kernel is supported - Require a value of 0 if a kernel is not supported - Clarify that values respect all properties

Add some advice for choosing coordination scopes

8836cfb

gmlueck reviewed Jul 24, 2023

View reviewed changes

Pennycook and others added 2 commits July 24, 2023 10:00

Bump copyright to 2023

4ce7593

Co-authored-by: Greg Lueck <[email protected]>

Fix link to kernel properties

5178d69

bashbaug reviewed Jul 24, 2023

View reviewed changes

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc Outdated Show resolved Hide resolved

Pennycook and others added 2 commits July 24, 2023 10:04

Fix formatting of ...

b2678ae

Co-authored-by: Greg Lueck <[email protected]>

Weaken "will" to "may"

7eb6c9e

Co-authored-by: Ben Ashbaugh <[email protected]>

Pennycook added 5 commits July 28, 2023 11:56

Clarify that strengthening is not permanent

03e7e12

This is in the ISO C++ definition, but replicating it here in a non-normative note may help readers understand the implications of the execution model.

Extend execution scopes to include all threads

300f7f6

Queries and properties relate to the intermediate threads of execution, and so they must also be part of the associated scope(s).

Tie device queries to kernel properties

c42ae01

This hopefully clarifies what the return value of this query really represents, and how it might be used.

Define how coordination scope works precisely

d317e7b

Fix typo: - => _

d7dbf7d

Pennycook force-pushed the forward_progress branch from 5db4836 to d7dbf7d Compare July 28, 2023 21:27

gmlueck reviewed Jul 31, 2023

View reviewed changes

Pennycook added 5 commits July 31, 2023 07:56

Fix namespaces for device queries

0a85dd4

Split launch queries into a separate extension

ec5f998

Split launch queries into two tables

14abbed

Remove "max" from num_sub_groups

2693865

Each work-group must be the same size, and the sub-group size has already been determined (because the query is relative to a specific work-group size). The number of sub-groups must therefore be known.

Add overview for launch queries extension

9adb3cb

gmlueck reviewed Aug 1, 2023

View reviewed changes

Pennycook added 2 commits August 1, 2023 09:41

Fix broken links to other extensions

652cc22

Clarify Param::argument_types

f9570bb

gmlueck reviewed Aug 1, 2023

View reviewed changes

Pennycook and others added 2 commits August 1, 2023 10:36

Remove Param::return_type and argument_types

36433dd

Co-authored-by: Greg Lueck <[email protected]>

Remove two more references to Param:: members

1cdabef

gmlueck approved these changes Aug 1, 2023

View reviewed changes

Pennycook requested a review from AerialMantis August 3, 2023 18:17

AerialMantis reviewed Aug 18, 2023

View reviewed changes

Pennycook mentioned this pull request Aug 21, 2023

[SYCL] Added ext_oneapi_non_uniform_groups aspect #10902

Merged

AerialMantis approved these changes Aug 25, 2023

View reviewed changes

steffenlarsen merged commit 6746114 into intel:sycl Aug 25, 2023

Pennycook deleted the forward_progress branch August 25, 2023 17:40

AlexeySachkov reviewed Feb 28, 2024

View reviewed changes

AlexeySachkov mentioned this pull request Sep 6, 2024

[SYCL] Implement max_num_work_groups from the launch queries extension #14333

Merged

[SYCL][Doc] Add draft forward progress extension #7598

[SYCL][Doc] Add draft forward progress extension #7598

Uh oh!

Conversation

Pennycook commented Nov 30, 2022

Uh oh!

Pennycook commented Dec 9, 2022

Uh oh!

Pennycook commented Mar 13, 2023

Uh oh!

Pennycook commented Jun 20, 2023

Uh oh!

Pennycook commented Jul 19, 2023

Uh oh!

AerialMantis left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

AerialMantis Jul 21, 2023

Choose a reason for hiding this comment

Uh oh!

Pennycook Jul 21, 2023

Choose a reason for hiding this comment

Uh oh!

AerialMantis Aug 18, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

AerialMantis Jul 21, 2023

Choose a reason for hiding this comment

Uh oh!

Pennycook Jul 21, 2023

Choose a reason for hiding this comment

Uh oh!

AerialMantis Aug 18, 2023

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

bashbaug left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

gmlueck left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

gmlueck left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Pennycook commented Aug 1, 2023

Uh oh!

gmlueck commented Aug 1, 2023

Uh oh!

AerialMantis left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

AerialMantis Aug 18, 2023 •

edited

Loading