OpenCL compiler doesn't implicitly convert size_t to unsigned

[Joose_S_]

At least on mad24(get_global_id(1), get_global_size(0), get_global_id(1)) compiler gives this so I have to explicitly conver all of the size_t to unsigned.  At least on AMD GPUs the code compiles fine without explicit casting.

    line 347: error: more
          than one instance of overloaded function "mad24" matches the
          argument list:
            function "mad24(int, int, int) C++"
            function "mad24(int2, int2, int2) C++"
            function "mad24(int3, int3, int3) C++"
            function "mad24(int4, int4, int4) C++"
            function "mad24(int8, int8, int8) C++"
            function "mad24(int16, int16, int16) C++"
            function "mad24(uint, uint, uint) C++"
            function "mad24(uint2, uint2, uint2) C++"
            function "mad24(uint3, uint3, uint3) C++"
            function "mad24(uint4, uint4, uint4) C++"
            function "mad24(uint8, uint8, uint8) C++"
            function "mad24(uint16, uint16, uint16) C++"
            argument types are: (size_t, size_t, size_t)
    out[mad24(get_global_id(1), get_global_size(0), get_global_id(1))] = value;
        ^

[Michal_M_Intel]

Hello,

In case binary is compiled as 64-bit , then size_t type is 8 bytes long.

Therefore implicit cast to unsigned int would cause overflow in some cases.

Additionally basing on the error message, compiler doesn't know how to handle input parameters of size_t.

https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/mad24.html

mad24 doesn't accept size_t as gentype.

By doing explicit cast compiler knows what is the gentype and can find proper function to handle mad24 operation.

[Joose_S_]

Yes but mad24 expects the values to be 24 bits anyways so it doesn't matter wether the value is 64 bit unsigned or 32 bit unsigned

[Ben_A_Intel]

I suspect on the devices where your kernel works size_t is a 32-bit type.  I see errors consistently across other OpenCL implementations when mad24() is used with 64-bit types.  Could you try the following kernel as an experiment?

__kernel void foo( __global ulong* dst )
{
    uint  ui32 = (uint)get_global_id(0);
    ulong ui64 = get_global_id(0);
    dst[0] = mad24( ui32, ui32, ui32 );    // ok
    dst[1] = mad24( ui64, ui64, ui64 );    // error
}

Because mad24() is defined only for 32-bit types, and there is no implicit conversion between 64-bit types and 32-bit types in C99 (see section 6.5 in the C99 spec) all OpenCL compilers should report an error for this kernel.

[Joose_S_]

You are actually correct. When I compile a 64-bit build on AMD it also fails. Sorry about the trouble.