KhronosGroup / KhronosGroup/OpenCL-CTS
test_half has overlapping memory read/write among work items when vector size is 3
- Dominant language
- C++
- Stars
- 232
- Forks
- 235
- Avg merge
- 8d 7h
- Merged PRs (30d)
- 18
Description
CTS test_half test has following kernel in Test_vStoreHalf.cpp file:
```
const char *source_v3[] = {
"__kernel void test( __global float *p, __global half *f,\n"
" uint extra_last_thread)\n"
"{\n"
" size_t i = get_global_id(0);\n"
" size_t last_i = get_global_size(0)-1;\n"
" size_t adjust = 0;\n"
" if(last_i == i && extra_last_thread != 0) {\n"
" adjust = 3-extra_last_thread;\n"
" } "
" vstore_half3",roundName,"( vload3(i, p-adjust), i, f-adjust );\n"
"}\n"
};
```
When the input parameter "extra_last_thread" is 2, the "adjust" variable will be 1. Then the "f-adjust" for vstore_half3 will cause 2 byte overlapping memory write between the last thread and the (last-1) thread. This overlapping memory write will cause a failure in CTS half/test_half vstore_half_rtz test on VeriSilicon GPUs.
```
vstore_half_rtz...
...............
Failure at [524285] with -0x1.7ffffap+126: *0xfbff vs 0xfbfe, vector_size = 3, address_space = global
vstore_half_rtz FAILED
vstorea_half_rtz...
................vstorea_half_rtz passed
```
Contributor guide
No contributing guide indexed for this repository
Research direction
Inspect the source_v3 kernel in Test_vStoreHalf.cpp, focusing on the extra_last_thread adjustment and the vstore_half3 destination indexing. Reproduce the vector_size=3 vstore_half_rtz failure described in test_half, then verify that the last work items no longer overlap and the test passes on the affected case.
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