KhronosGroup / KhronosGroup/OpenCL-CTS

Missing 'and' instruction in atomic_and_local_{int,uint}.bc of spir test

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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.