KhronosGroup / KhronosGroup/OpenCL-CTS
test_non_uniform_work_group/non_uniform_3d_barriers has segfault if global buffer size is adjusted
- Dominant language
- C++
- Stars
- 232
- Forks
- 235
- Avg merge
- 8d 7h
- Merged PRs (30d)
- 18
Description
In test test_non_uniform_work_group/non_uniform_3d_barriers, initial global buffer size the product of _globalSize.
Function TestNonUniformWorkGroup::adjustGlobalBufferSize adjusts global buffer size according to query result of CL_DEVICE_MAX_MEM_ALLOC_SIZE.
If global buffer size is adjusted, there is segfault during kernel execution, because _globalSize is not adjusted accordingly. testBarriers function in the kernel accesses global buffer with globalIndex. globalIndex exceeds global buffer bound for some workitems.
```
void testBarriers(__global unsigned int *errorCounterBuffer, __local unsigned int *testLocalBuffer, __global unsigned int *testGlobalBuffer) {
uint gid0 = get_global_id(0);
uint gid1 = get_global_id(1);
uint gid2 = get_global_id(2);
...
uint globalIndex = getGlobalIndex(gid2-get_global_offset(2), gid1-get_global_offset(1), gid0-get_global_offset(0));
...
testGlobalBuffer[globalIndex] = 0;
}
__kernel void testKernel(__global DataContainerAttrib *results, ...
{
testBarriers(errorCounterBuffer, testLocalBuffer, testGlobalBuffer);
}
```
Contributor guide
No contributing guide indexed for this repository
Research direction
Start with TestNonUniformWorkGroup::adjustGlobalBufferSize and the test_non_uniform_work_group/non_uniform_3d_barriers test, then inspect the testBarriers kernel and its globalIndex calculation. Run the test with a reduced CL_DEVICE_MAX_MEM_ALLOC_SIZE and verify that kernel accesses remain within the allocated global buffer without a segfault.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- testing-qa
- Issue type
- Bug
- Difficulty
- 3/5
- Estimated time
- 1-2 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 45/100