An odd bug regarding as_* nonconversions: shall they be the same size?

Some time ago, an user sent me an odd bug report where the OpenCL compiler complained about as_* arguments being the same size.
The code looks like

01kernel void Luffa_1way(global uint *wuData, ...) {02
03
04    // snip
05
06
07    uint8 M = (uint8)(wuData[0], wuData[1], wuData[2], wuData[3],
08                      wuData[4], wuData[5], wuData[6], wuData[7]);
09    for(uint i = 0; i < 5; i++)
10    {
11
12
13        // snip
14
15
16        if(i == 0) {
17            M = (uint8)(wuData[ 8], wuData[ 9], wuData[10], wuData[11],
18                        wuData[12], wuData[13], wuData[14], wuData[15]);
19        } else if(i == 1) {
20            M = (uint8)(wuData[16], wuData[17], wuData[18], as_uint(as_uchar4(get_global_id(0)).wzyx),
21                        0x80000000u, 0, 0, 0);
22        } else if(i == 2) {
23            M = (uint8)(0);
24        } else if(i == 3) {
25            // snip
26        }
27    }
28    // snip
29}

and the error reported, for LN20, twice: “argument type to opencl as_* function: expected src and dst have the same size”. This makes no sense to me. It seems as_* functions are non-conversions meant to be used exactly this way.

I’m pretty sure everything compiled fine last week and I don’t recall updating anything. I’m using AMD Omega drivers.

Can you check my reasoning please? I must be missing something.

I will most likely signal this in AMD forums as well in the meanwhile, how do you suggest me to workaround?

If you use the CPU device and your app is compiled for x64, get_global_id() returns a size_t value with is 64-bit wide.
In this case, as_uchar4(get_global_id(0)) is not legal.

You should first check CL_DEVICE_ADDRESS_BITS to ensure that size_t is 32-bit.

So that error message really means the parameters are of different sizes? I got it the other way around!

Thank you very much. I checked my device selection methods and I found out I forgot a negation in condition evaluation.

Given this specific context I will just truncate to uint.

I will have to investigate how this happened to the original user but I assume Intel iGPUs might export 64 bits as well.