Naming conflict between built-in functions of CUDA and OCL

[MathisonBao]

The void sincosf(float, float *, float *) is a CUDA built-in function, the float sincos(float, float *) is a OCL built-in function but not CUDA. So the user could define a function as same as the OCL one.

In the implementation of sincosf function, sincos will be called. (in order to use the SPIR-V Instruction: OpExtInst sincos), but if a user defines a function, named sincos, the sincosf will return the wrong result (sincosf will call the function that the user defined).

For example:

__device__
float sincos(float x, float &y) {
    *y = x + 1.0;
    return x + 2.0;
}

__global__
void test(float *IN, float *OUT0, float *OUT1, int num) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < num) {
        float x, y;
        x = sincos(IN[i], &y);
        float a, b;
        sincosf(IN[i], &a, &b);
        OUT0[i] = x + y;
        OUT1[i] = a + b;
    }
}

The data in OUT0 and the data in OUT1 will be the same. Is this a bug? Is there any solution? Thanks a lot.

[pjaaskel]

A bug. We might need to rename the conflicting user functions using macros.

[MathisonBao]

Hi, pjaaskel:

Thanks for your reply. Could you explain in detail how to solve it using macro? I'm not sure how to use macros to differentiate between device or host functions with the same function name.

Thanks, Mathison

[pjaaskel]

I meant that hipcc could (force) include a header that has lines such as ``#define sincos _user_sincos` that avoid the name clash. But after a bit of thinking, that's not an ideal solution as it renames user facing symbols, so we should do something similar internally for the sincos() builtin calls.

[linehill]

A proper way to fix this issue is to use SPIR-V builtins instead of OpenCL built-ins in the headers. E.g. HIP's sincosf would be mapped to __spirv_ocl_sincos.

[MathisonBao]

Okay, thanks pjaaskel and linehill!