as_* function: expected src and dst have the same size

[MaxDZ8]

An user running 14.12 (apparently on Win8?) sent me a screenshot looking like

>oclcckvck.exe
Found 1 OpenCL platform for processing.
Platform 0 counts 1 device to test.
- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - 
OpenCL error -11 for Luffa_1W.cl.Luffa_1way, attempted compile with "-D LUFFA_HE...
ERROR LOG:
"<redacted>\temp\tmpfile.cl", line 174: error: bad
        argument type to opencl as_* function: expected src and dst have the
        same size
            M = (uint8)(wuData[16], wuData[17], wuData[18], as_uint(as_uchar4(...

"<redacted>\temp\tmpfile.cl", line 174: error: bad
        argument type to opencl as_* function: expected src and dst have the
        same size
            M = (uint8)(wuData[16], wuData[17], wuData[18], as_uint(as_uchar4(...
                                                                    ^
"<redacted>\temp\tmpfile.cl", line 175: error: not
        enough initializer values
                             0x80000000u, 0, 0, 0);

3 errors detected in the compilation of "<redacted>
Frontend phase failed compilation.

>

[MaxDZ8]

Some elaborations.

Leaving aside the behavior of as_uint when dealing with multi-component is implementation-defined... This is extremely odd.

By reading the CL 1.2 spec, section 6.2 I can tell for sure this isn't going to happen with casts.

It obviously cannot be convert_* (it's basically a packed per-component cast) either as page 210 reads

The number of elements in the source and destination vectors must match.

Section 6.2.4 hints "between the lines" reinterpreting data using as_* is the preferred, most performant way:

Several methods to achieve this (non-) conversion

And even goes in details with note 30:

some other extensions to the C language designed to support particular vector ISA (...) use such conversions in conjunction with swizzle operators to achieve type unconversion. So as to support legacy code of this type, as_typen() allows conversions between vectors of the same size but different numbers of elements, even though the behavior of this sort of conversion is not likely to be portable except to other OpenCL implementations for the same hardware architecture.

Interesting: section 6.2.4.1 - reinterpreting types using unions.

6.2.4.2 Reinterpreting Types Using as_type() and as_typen()

(page 214)

All data types described in tables 6.1 and 6.2 (except bool, half³² and void) may be also reinterpreted as another data type of the same size using the as_type() operator for scalar data types and the as_typen() operator³³ for vector data types.

...

³³ ... as_type() and as_typen() constructs are intended to reflect the organization of data in register. The as_type() and as_typen() constructs are intended to _compile to _no* instructions* on devices that use a shared register file designed to operate on both the operand and result types ... an architecture may load a char into a 32-bit register, or a char vector into a SIMD vector register with fixed 32-bit element size ...

It is an error to use as_type() or as_typen() operator to reinterpret data to a type of a different number of bytes.

Which pretty much says the opposite of above.

Also:

int i;
// Legal. Result is implementation-defined.
short2 j = as_short2(i);
int4 i;
// Legal. Result is implementation-defined.
short8 j = as_short8(i);

Which seems to be pretty much my case here.

I'd sure need to get more data on this error frequency. Quite sure I'm on the good side but that cannot be let go.

[MaxDZ8]

Managed to replicate. Somehow. Omega drivers. I'm pretty sure this used to work... Full report is (taken from M8M however):

OpenCL error -11 for Luffa_1W.cl.Luffa_1way, attempted compile with "-D LUFFA_HEAD"
ERROR LOG:
"<tempDir>\OCL4232T4.cl", line 174: error: bad
          argument type to opencl as_* function: expected src and dst have the
          same size
              M = (uint8)(wuData[16], wuData[17], wuData[18], as_uint(as_uchar4(get_global_id(0)).wzyx),
                                                                                ^

"<tempDir>\OCL4232T4.cl", line 174: error: bad
          argument type to opencl as_* function: expected src and dst have the
          same size
              M = (uint8)(wuData[16], wuData[17], wuData[18], as_uint(as_uchar4(get_global_id(0)).wzyx),
                                                                      ^

"<tempDir>\OCL4232T4.cl", line 175: error: not
          enough initializer values
                          0x80000000u, 0, 0, 0);

[MaxDZ8]

It turns out that means the opposite of what's written. Fixed in internal source by truncating to uint. To be published after M8M update.