[SYCL] Deprecate old sycl_ext_oneapi_group_load_store
#15405
Conversation
Newer version is under `proposed/`, I'm going to move it to `experimental/` in a separate PR for a cleaner `git diff`.
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
| the documentation in its current state does not reflect the intended long-term | ||
| direction of the extension. | ||
| NOTE: This extension has been replaced with a new version under the same name | ||
| that completely changed the interfaces. |
There was a problem hiding this comment.
In the next PR, can you add a link here to the new spec?
aelovikov-intel
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
aelovikov-intel
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
aelovikov-intel
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
aelovikov-intel
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
| LocOffsetA, | ||
| slmVecA); | ||
| vec<TOperand, 2> slmVecA; | ||
| group_load(sg, SrcA, slmVecA, properties(data_placement_striped)); |
There was a problem hiding this comment.
why did you choose here data_placement_striped instead of data_placement_blocked. Is the striped one the one that guarantees vectorization inside the SG?
There was a problem hiding this comment.
I chose it because (IIUC) that matches the behavior of the old API. I didn't use the property for scalar overloads because both data layouts are essentially the same for scalars.
There was a problem hiding this comment.
Sounds good if this is old API behavior.
It is weird then that by default, they made default placement is the blocked one.
There was a problem hiding this comment.
Sounds good if this is old API behavior.
Old:
Returns: Returns N elements per work-item, corresponding to the N memory locations at src + i * get_max_local_range() + get_local_id() for i between 0 and N.
New:
Effects: Loads N elements from in_iter to out using the g group object. Properties may specify data placement. Default data placement is a blocked one: out[i] = in_iter[g.get_local_linear_id() * N + i]; in striped case: out[i] = in_iter[g.get_local_linear_id() + g.get_local_linear_range() * i]; for i between 0 and N.
Newer version is under
proposed/, I'm going to move it toexperimental/in a separate PR for a cleanergit diff.