KhronosGroup / KhronosGroup/OpenCL-CTS

test_half has overlapping memory read/write among work items when vector size is 3

Open
#1,149 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

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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.