[SYCL][Doc] non_uniform_groups extension spec fixes #13702
Conversation
Changes: - fixed descriptions of `get_local_linear_id` methods so they are not recursive anymore - fixed descriptions of `get_[group|local]_linear_range` methods to use correct return type and references - added clarification that `get_group_id` is relative to a parent group
| Ballot-groups are always created in a range of two: the first ballot-group | ||
| contains all work-items where the predicate is true, and the second | ||
| ballot-group contains all work-items where the predicate is false. | ||
| Ballot-groups are always created in a range of two: the one ballot-group |
There was a problem hiding this comment.
Note for reviewers: I removed mentions of "first" and "second", because below we explicitly say that work-items where the predicate is true are in ballot-group with local id = 1 and the rest of work-items are in ballot-group with local id = 0.
Which meant that 1 is "first" and 0 is "second". I think here we shouldn't use any ordering-related terms to avoid contradiction with the detailed description of methods below.
| @@ -335,7 +335,7 @@ than only sub-groups. | |||
| ---- | |||
| id_type get_group_id() const; | |||
| ---- | |||
| _Returns_: An `id` representing the index of the ballot-group. | |||
| _Returns_: An `id` representing the index of the ballot-group in a parent group. | |||
There was a problem hiding this comment.
Note for reviewers: I'm not entirely sure if this clarification is needed, perhaps that piece of information is implied and should be obvious, but I though that it would be better to state it explicitly.
If it is decided that it makes sense to add such clarification, then I should have another pass over the doc, because I haven't been consistent with adding it.
There was a problem hiding this comment.
I am not even sure the clarification is that clear. What does it mean to have an "index in a parent group". You can create multiple ballot-groups from the same sub-group, and they may have the same indices.
There was a problem hiding this comment.
The clearest thing to do here might be to say something really precise (and verbose), like:
| _Returns_: An `id` representing the index of the ballot-group in a parent group. | |
| _Returns_: An `id` representing the index of the ballot-group to which the | |
| calling work-item belongs, as determined by the work-item's associated | |
| `predicate` value. |
...but I appreciate this isn't perfect either. One of the things we've been exploring is to replace the concept of a ballot-group with "the result of a logical partitioning", which might make this easier to explain. This would then read as "the index of the logical partition to which..."
| work-item is the leader of the fixed-size-group, and `false` for all other | ||
| work-items in the fixed-size-group. The leader of the fixed-size-group is guaranteed | ||
| to be the work-item for which `get_local_id()` returns 0. | ||
| _Returns_: `true` for exactly one work-item in the fixed-size-group, if the |
There was a problem hiding this comment.
Note for reviewers: no real changes here, just a re-formatting to fit into 80 symbol limit per line
AlexeySachkov
changed the title
| Ballot-groups are always created in a range of two: the first ballot-group | ||
| contains all work-items where the predicate is true, and the second | ||
| ballot-group contains all work-items where the predicate is false. | ||
| Ballot-groups are always created in a range of two: the one ballot-group |
There was a problem hiding this comment.
I would also remove the "the", though I think both are correct.
| Ballot-groups are always created in a range of two: the one ballot-group | |
| Ballot-groups are always created in a range of two: one ballot-group |
There was a problem hiding this comment.
I agree that "the" should be removed here.
| @@ -335,7 +335,7 @@ than only sub-groups. | |||
| ---- | |||
| id_type get_group_id() const; | |||
| ---- | |||
| _Returns_: An `id` representing the index of the ballot-group. | |||
| _Returns_: An `id` representing the index of the ballot-group in a parent group. | |||
There was a problem hiding this comment.
I am not even sure the clarification is that clear. What does it mean to have an "index in a parent group". You can create multiple ballot-groups from the same sub-group, and they may have the same indices.
| Ballot-groups are always created in a range of two: the first ballot-group | ||
| contains all work-items where the predicate is true, and the second | ||
| ballot-group contains all work-items where the predicate is false. | ||
| Ballot-groups are always created in a range of two: the one ballot-group |
There was a problem hiding this comment.
I agree that "the" should be removed here.
| @@ -335,7 +335,7 @@ than only sub-groups. | |||
| ---- | |||
| id_type get_group_id() const; | |||
| ---- | |||
| _Returns_: An `id` representing the index of the ballot-group. | |||
| _Returns_: An `id` representing the index of the ballot-group in a parent group. | |||
There was a problem hiding this comment.
The clearest thing to do here might be to say something really precise (and verbose), like:
| _Returns_: An `id` representing the index of the ballot-group in a parent group. | |
| _Returns_: An `id` representing the index of the ballot-group to which the | |
| calling work-item belongs, as determined by the work-item's associated | |
| `predicate` value. |
...but I appreciate this isn't perfect either. One of the things we've been exploring is to replace the concept of a ballot-group with "the result of a logical partitioning", which might make this easier to explain. This would then read as "the index of the logical partition to which..."
| _Returns_: A `range` representing the number of ballot-groups in a parent group. | ||
|
|
||
| NOTE: This will always return a `range` of 2, as there will always be two groups; | ||
| one representing the group of work-items where the predicate was true and |
There was a problem hiding this comment.
Maybe we could simplify this to:
| _Returns_: A `range` representing the number of ballot-groups in a parent group. | |
| NOTE: This will always return a `range` of 2, as there will always be two groups; | |
| one representing the group of work-items where the predicate was true and | |
| _Returns_: The result of `range<1>(2)`. | |
| NOTE: This always returns a `range` of 2, as there will always be two groups; | |
| one representing the group of work-items where the predicate was true and |
| @@ -376,19 +376,19 @@ _Returns_: A linearized version of the `id` returned by `get_group_id()`. | |||
| ---- | |||
| id_type get_local_linear_id() const; | |||
| ---- | |||
| _Returns_: A linearized version of the `id` returned by `get_local_linear_id()`. | |||
| _Returns_: A linearized version of the `id` returned by `get_local_id()`. | |||
| _Returns_: A `range` representing the number of fixed-size-groups in a parent | ||
| group. |
There was a problem hiding this comment.
Maybe something like:
| _Returns_: A `range` representing the number of fixed-size-groups in a parent | |
| group. | |
| _Returns_: A `range` representing the number of fixed-size-groups created by | |
| partitioning the parent group with `Size`. |
I'm being pedantic, but the number of fixed-size-groups in a parent group might be a concept that some people struggle with. If somebody partitions a group into fixed-size-groups of size 1 and again into fixed-size-groups of size 2, those two partitionings don't know anything about one another.
| @@ -668,7 +669,7 @@ public: | |||
| ---- | |||
| id_type get_group_id() const; | |||
| ---- | |||
| _Returns_: An `id` representing the index of the tangle-group. | |||
| _Returns_: An `id` representing the index of the tangle-group in a parent group. | |||
There was a problem hiding this comment.
| _Returns_: An `id` representing the index of the tangle-group in a parent group. | |
| _Returns_: The result of `id<1>(0)`. |
| _Returns_: An `id` representing the index of the opportunistic-group in a parent | ||
| group. |
There was a problem hiding this comment.
| _Returns_: An `id` representing the index of the opportunistic-group in a parent | |
| group. | |
| _Returns_: The result of `id<1>(0)`. |
|
|
||
| . Inconsistencies in `fixed_size_group` | ||
| + | ||
| -- | ||
| `fixed_size_group` accepts template argument `PartitionSize` of type | ||
| `std::size_t`, which can be an unsigned 64-bit integer. However, | ||
| `linear_id_type` is defined as `uint32_t` for this group. Even though with the | ||
| current restrictions it is impossible to encounter such situation, there is | ||
| still a mismatch allowing to create a group so big that linear IDs won't be | ||
| correct for some of its work-items. | ||
| -- |
There was a problem hiding this comment.
I think this is just a bug, and we should update linear_id_type with size_t.
My initial prototype only works with sub-groups (which are limited to uint32_t) and that's probably how this crept in. If/when we allow partitioning of work-groups (and other groups) into fixed-size groups, we may need size_t.
|
The spec is being updated in #14604, I've moved all my comments there |
Signed-off-by: John Pennycook <john.pennycook@intel.com>
Changes:
get_local_linear_idmethods so they are not recursive anymoreget_[group|local]_linear_rangemethods to use correct return type and referencesget_group_idis relative to a parent group