KhronosGroup / KhronosGroup/OpenCL-CTS
Missing 'and' instruction in atomic_and_local_{int,uint}.bc of spir test
- Dominant language
- C++
- Stars
- 232
- Forks
- 235
- Avg merge
- 8d 7h
- Merged PRs (30d)
- 18
Description
test_atomic_fn.atomic_and_local_int.bc64 in spir/atomics.zip is wrong.
The OpenCL C part of the test is as follows
```opencl
// test_atomic_fn.atomic_and_local_int.cl in spir/atomics.zip
kernel void test_atomic_fn(global int *Src, global int *Dst,
volatile local int *Local){
int tid = get_global_id(0), lid = get_local_id(0), gid = get_group_id(0);
size_t numBits;
// Copy from global source to local buffer.
Local[lid] = Src[tid];
barrier(CLK_LOCAL_MEM_FENCE);
numBits = sizeof(Local[0]) * 8;
int bitIndex = tid - (gid * numBits);
atomic_and(&Local[0], ~(1U << bitIndex));
barrier(CLK_LOCAL_MEM_FENCE);
Dst[gid] = Local[0];
}
```
The `bitIndex` may be negative in some cases, and `1U << bitIndex` seems to be undefined behavior. But [OpenCL C spec rules it to be valid](https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#operators-shift). To conform the spec, an `bitIndex & 31` should be generated by compiler frontend to avoid over-width bit shifting, and clang actually does it (this can be verified by executing `clang -cc1 -triple spir64 -xcl -cl-std=CL2.0 -include opencl-c.h test_atomic_fn.atomic_and_local_int.cl -S -emit-llvm -disable-llvm-passes`).
However, the corresponding IR part test_atomic_fn.atomic_and_local_int.bc64 is wrong:
```llvm
%11 = sub i64 %4, %10 ; bitIndex = tid - (gid * numBits)
%12 = trunc i64 %11 to i32
%13 = shl i32 1, %12 ; 1U << bitIndex
```
So, an `and` is missing between the `sub` and the `shl` instructions. The correct code should be
``` llvm
%11 = sub i64 %4, %10 ; bitIndex = tid - (gid * numBits)
%12 = trunc i64 %11 to i32
%shl.mask = and i32 %12, 31
%13 = shl i32 1, %shl.mask ; 1U << bitIndex
```
The missing `and` in the bc make the semantic inconsistent between the .cl and the .bc file.
atomic_and_local_uint is also affected.
Contributor guide
No contributing guide indexed for this repository
Research direction
Compare the atomic_and_local_int and atomic_and_local_uint entries in spir/atomics.zip with their OpenCL C sources and generated IR. Use the provided clang -cc1 SPIR64 command to verify the frontend output, then confirm both .bc files consistently represent the masked shift instruction.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- c
- Domain
- testing-qa
- Issue type
- Bug
- Difficulty
- 2/5
- Estimated time
- 1-3 hours
- Activity status
- Stale
- Clarity
- Clearly specified
- Newbie friendliness
- 50/100