atom_add with int* and uint arguments

Hi Anastasia/Alexey/Xiuli,

When we compile some OpenCL conformance test for atom_add with clang trunk we got overloading resolution issue:

__kernel void test_atomic_fn(volatile __global uint *destMemory, __global uint *oldValues)
{
    int tid = get_global_id(0);

oldValues[tid] = atom_add( &destMemory[0], tid + 3 );
atom_add( &destMemory[0], tid + 3 );
   atom_add( &destMemory[0], tid + 3 );
   atom_add( &destMemory[0], tid + 3 );

}

/tmp/AMD_17652_49/t_17652_51.cl:6:19: error: call to 'atom_add' is ambiguous
        oldValues[tid] = atom_add( &destMemory[0], tid + 3 );
                         ^~~~~~~~
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14336:12: note: candidate function
int __ovld atom_add(volatile __global int *p, int val);
           ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14337:21: note: candidate function
unsigned int __ovld atom_add(volatile __global unsigned int *p, unsigned int val);
                    ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14345:13: note: candidate function
long __ovld atom_add(volatile __global long *p, long val);
            ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14346:22: note: candidate function
unsigned long __ovld atom_add(volatile __global unsigned long *p, unsigned long val);
                     ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14347:13: note: candidate function
long __ovld atom_add(volatile __local long *p, long val);
            ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14348:22: note: candidate function
unsigned long __ovld atom_add(volatile __local unsigned long *p, unsigned long val);
                     ^
/tmp/AMD_17652_49/t_17652_51.cl:7:2: error: call to 'atom_add' is ambiguous
        atom_add( &destMemory[0], tid + 3 );
        ^~~~~~~~

The reason is that the call expects atom_add(volatile global uint*, int) but cannot find a better one from existing declarations.

I suggest to add the missing declarations for mixed-signedness arguments to opencl-c.h. What's your opioin?

Thanks.

Sam

Hi Anastasia/Alexey/Xiuli,

When we compile some OpenCL conformance test for atom_add with clang trunk we got overloading resolution issue:

__kernel void test_atomic_fn(volatile __global uint *destMemory, __global uint *oldValues)
{
    int tid = get_global_id(0);

oldValues[tid] = atom_add( &destMemory[0], tid + 3 );
atom_add( &destMemory[0], tid + 3 );
   atom_add( &destMemory[0], tid + 3 );
   atom_add( &destMemory[0], tid + 3 );

}

/tmp/AMD_17652_49/t_17652_51.cl:6:19: error: call to 'atom_add' is ambiguous
        oldValues[tid] = atom_add( &destMemory[0], tid + 3 );
                         ^~~~~~~~
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14336:12: note: candidate function
int __ovld atom_add(volatile __global int *p, int val);
           ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14337:21: note: candidate function
unsigned int __ovld atom_add(volatile __global unsigned int *p, unsigned int val);
                    ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14345:13: note: candidate function
long __ovld atom_add(volatile __global long *p, long val);
            ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14346:22: note: candidate function
unsigned long __ovld atom_add(volatile __global unsigned long *p, unsigned long val);
                     ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14347:13: note: candidate function
long __ovld atom_add(volatile __local long *p, long val);
            ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14348:22: note: candidate function
unsigned long __ovld atom_add(volatile __local unsigned long *p, unsigned long val);
                     ^
/tmp/AMD_17652_49/t_17652_51.cl:7:2: error: call to 'atom_add' is ambiguous
        atom_add( &destMemory[0], tid + 3 );
        ^~~~~~~~

The reason is that the call expects atom_add(volatile global uint*, int) but cannot find a better one from existing declarations.

I suggest to add the missing declarations for mixed-signedness arguments to opencl-c.h. What's your opioin?

Hi,

This was a regression caused by: ⚙ D24113 Allow implicit conversions between incompatible pointer types in overload resolution in C.. I'm
not sure if this was intended or not.

-Tom

Thanks Tom. I will ask the author.

Sam