sycl/test-e2e/GroupAlgorithm out of bounds access
- 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
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