KhronosGroup / KhronosGroup/SYCL-Docs
[Proposal] Deprecate SYCL Offset Accessors
- Dominant language
- JavaScript
- Stars
- 158
- Forks
- 80
- Avg merge
- 7d 6h
- Merged PRs (30d)
- 5
Description
The SYCL spec allows for an extra offset argument to be used when constructing an accessor.
```c++
/* Available only when: (Dimensions > 0) */
template
accessor(buffer& bufferRef,
handler& commandGroupHandlerRef, range accessRange,
id accessOffset, const property_list& propList = {});
/* Available only when: (Dimensions > 0) */
template
accessor(buffer& bufferRef,
handler& commandGroupHandlerRef, range accessRange,
id accessOffset, TagT tag,
const property_list& propList = {});
```
This seems like a good idea to allow the programmer to offset into a mem arg, such that `acc[0]` refers not to the base of some native allocation.
However, this means that when calculating `acc[arbitrary_idx]`, internally the SYCL implementation must add `arbitrary_idx + accessOffset` in order to generate the required index.
In 99% (guess) of cases, `accessOffset` is zero. However most compilers (DPC++, for one) lack the host-device optimizations necessary in order to propagate the zero value to the device compilation pass. This means that the access offset becomes a kernel argument, and the zero values are loaded at runtime and then added to each accessor base pointer, at least once at accessor initialization. This results in more kernel args, clock cycles wasted adding zero to ptrs, and increased register usage within the kernel.
As a quick demonstration:
A simple SYCL kernel just doing `acc[0] = 1`:
```asm
.weak .entry _ZTS6kernel(
.param .u64 _ZTS6kernel_param_0,
.param .align 8 .b8 _ZTS6kernel_param_1[8] // THIS IS THE OFFSET ARG WHICH IS USUALLY ZERO
)
{
.reg .b32 %r<2>;
.reg .b64 %rd<8>;
// %bb.0: // %entry
ld.param.u32 %rd1, [_ZTS6kernel_param_0+4];
shl.b64 %rd2, %rd1, 32;
ld.param.u32 %rd3, [_ZTS6kernel_param_0];
or.b64 %rd4, %rd2, %rd3;
ld.param.u64 %rd5, [_ZTS6kernel_param_1]; // LOAD OFFSET ARG
shl.b64 %rd6, %rd5, 2;
add.s64 %rd7, %rd4, %rd6; // COMBINE OFFSET ARG WITH PTR
mov.b32 %r1, 1;
st.global.u32 [%rd7], %r1;
ret;
// -- End function
}
```
vs a simple SYCL kernel doing `acc.get_multi_ptr()[0]`
```asm
.weak .entry _ZTS6kernel(
.param .u64 _ZTS6kernel_param_0
)
{
.reg .b32 %r<2>;
.reg .b64 %rd<5>;
// %bb.0: // %entry
ld.param.u32 %rd1, [_ZTS6kernel_param_0+4];
shl.b64 %rd2, %rd1, 32;
ld.param.u32 %rd3, [_ZTS6kernel_param_0];
or.b64 %rd4, %rd2, %rd3;
mov.b32 %r1, 1;
st.global.u32 [%rd4], %r1;
ret;
// -- End function
}
```
Note that when `get_multi_ptr()` is called, the compiler is able to see that the `accessOffset` member is not used in the kernel and it can remove a kernel argument.
This can be solved in two ways in the SYCL specification:
1. Add an extra template arg to the `accessor` class, which would describe whether the accessor was offset or not. This would allow the use of `if constexpr` in the accessor setup which would use the accessOffset or not.
2. Deprecating and eventually removing support for offset accessors, which presumably are not being used much anyway.
Contributor guide
No contributing guide indexed for this repository
Research direction
Start by reviewing the SYCL accessor constructor specification and the two alternatives described in this issue. This is done when the community agrees on a resolution for offset accessors and the corresponding specification changes are identified and documented.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- documentation
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100