Binary Instruction seen with illegal int type. Legalization support missing when compiling kernel using unsigned long long

[ukasz]

When compiling the following simple kernel I get an error "Binary instruction seen with illegal int type". This happens mostly on unsigned long long, but I believe I've seen similar issue on other types.
Here is the kernel:

#define fuzz_t unsigned long long
__kernel void foo(__global fuzz_t * mem, unsigned int n) {
        fuzz_t tmp[] = { 0x5856ee2a6d198413, 0x0b81b3a5685ed7c8, 0x61dc5f842ce859cb, 0x685ec65f3cdb7e77, };

        tmp[0] = mem[206619138 % n];
        tmp[3] *= 0x10dcf06f60227df6;
        tmp[2] += tmp[0];
        tmp[0] += tmp[1];
        tmp[1] <<= (tmp[2] % 8);

        fuzz_t all = 0;
        for(unsigned int i = 0; i < 4; i++) all += tmp[i];
        for(unsigned int i = 0; i < n; i++) all += mem[i];
        mem[0] = all;
}

I attached tar.gz package to reproduce the issue, steps to reproduce:

tar xzf cd 1-u64-5-66.tar.gz
cd 1-u64-5-66
./main-opencl
Device selected: GPU
Binary Instruction seen with illegal int type. Legalization support missing. Inst opcode:25Error building program: -11
Build log:
IGC: Internal Compiler Error: Segmentation violation

[ukasz]

1-u64-5-66.tar.gz

[karolzwolak]

The long long type in OpenCL C is 128 bits wide and it is a reserved data type, unlike in C. A reserved datatype means it's recognized by the compiler, but you should not use it, and so our backend doesn't have support for it, hence the crash.
Our frontend incorrectly accepts this reserved datatype without any error or warning. This has been corrected with in llvm/llvm-project#196002, which introduces a warning.

In conclusion, this is expected behavior.