KhronosGroup / KhronosGroup/SYCL-Docs

[Proposal] Deprecate SYCL Offset Accessors

Open
#569 8 comments 0 reactions 0 assignees View on GitHub
enhancement
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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.