Community
cancel
Showing results for 
Search instead for 
Did you mean: 
Highlighted
Beginner
66 Views

OpenCL compiler doesn't implicitly convert size_t to unsigned

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;
        ^

 

0 Kudos
4 Replies
Highlighted
Employee
66 Views

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.

0 Kudos
Highlighted
Beginner
66 Views

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

0 Kudos
Highlighted
Employee
66 Views

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.

0 Kudos
Highlighted
Beginner
66 Views

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

0 Kudos