search

llvm / llvm-project

The LLVM Project is a collection of modular and reusable compiler and toolchain technologies.
http://llvm.org
Other
28.82k stars 11.91k forks source link

[NVPTX] Clang errors with non-trivial scalar-to-vector conversion when using opencl vector of two 8-bit integer as input #107219

Open rashedmyt opened 1 month ago

rashedmyt commented 1 month ago

This issue started occurring with LLVM 18. Below is the reproduction code and the command used to reproduce the issue.

typedef char cchar __attribute__((ext_vector_type(2)));

__global__ void kernel(cchar* out, cchar in) {
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    out[tid] = in;
}
$ clang --cuda-path=<cuda_root> --cuda-gpu-arch=sm_60 --cuda-device-only -O3 -S repro.cu -o repro.ptx
error: non-trivial scalar-to-vector conversion
1 error generated when compiling for sm_60.

Here is a link of compiler explorer where issue is reproduced - https://cuda.godbolt.org/z/c8jscoach

More Observations:

  1. The issue only occurs when using a scalar input of the 8-bit integer vector of size 2. For other sizes and when a pointer is used, PTX is generated.
  2. The issue doesn't occur when using other lower bit integers such as short for 16-bit integers
Artem-B commented 1 month ago

It may be related to https://github.com/llvm/llvm-project/issues/104864