intel / intel/llvm

sycl/test-e2e/GroupAlgorithm out of bounds access

Open
#13,939 0 comments 1 reaction 0 assignees View on GitHub
bug confirmed
Dominant language
LLVM
Stars
1.5k
Forks
854
Avg merge
3d 17h
Merged PRs (30d)
137

Description

### Describe the bug

Please see https://github.com/intel/llvm/blob/371bd980c5ef8c2899f540f10096615b6a1ed6e3/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp:

```
template
void test(const InputContainer &input, BinaryOperation binary_op,
typename InputContainer::value_type identity) {
...
constexpr size_t G = 16;
...
buffer in_buf(input.data(), input.size());
buffer out_buf(output.data(), output.size());
q.submit([&](handler &cgh) {
accessor in{in_buf, cgh, sycl::read_only};
accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
cgh.parallel_for(nd_range<1>(G, G), [=](nd_item<1> it) {
group<1> g = it.get_group();
int lid = it.get_local_id(0);
out[lid] = inclusive_scan_over_group(g, in[lid], binary_op);
});
...
}
...
// Smaller size as the multiplication test
// will result in computing of a factorial
// 12! fits in a 32 bits integer.
constexpr int M = 12;
std::array input_small;
...
test(input, sycl::minimum(),
std::numeric_limits::max());
...
```
This defines a buffer of 12 integers, invokes a SYCL kernel with a global and local size of 16, and uses that to write beyond the end of `out_buf`.

The same problem appears to exist in `exclusive_scan_sycl2020.cpp` as well.

### To reproduce

N/A for now because current OpenCL implementations reject this test for other reasons. Once they get updated to actually run, this out of bounds access will come up at some point when just running SYCL E2E tests normally.

### Environment

_No response_

### Additional context

_No response_

Contributor guide

Open the contributing guide

Research direction

Start with sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp and compare the 12-element input/output buffers with the nd_range global and local size of 16. Inspect exclusive_scan_sycl2020.cpp for the same pattern. The work is done when both tests avoid out-of-bounds buffer access and pass when the relevant SYCL E2E tests run.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
testing-qa
Issue type
Bug
Difficulty
2/5
Estimated time
1-3 hours
Activity status
Stale
Clarity
Clearly specified
Newbie friendliness
45/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.