[cfe-dev] atom_add with int* and uint arguments

[Liu, Yaxun (Sam) via cfe-dev]

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
-------------- next part --------------
A non-text attachment was scrubbed...
Name: winmail.dat
Type: application/ms-tnef
Size: 18290 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20160907/bb7d73a1/attachment.bin>