KhronosGroup / KhronosGroup/OpenCL-CTS

test_non_uniform_work_group/non_uniform_3d_barriers has segfault if global buffer size is adjusted

Open
#1,142 0 comments 0 reactions 0 assignees View on GitHub
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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.