When compiling CUDA benchmark that defines vector operators itself (specifically, boxfilter-cuda in HeCBench that uses "helper_math.h" header file provided by NVIDIA cuda-samples), the following "use of overloaded operator ambiguous" compilation errors are produced.
warning: cucc is a work-in-progress. It is incomplete and may behave incorrectly.
Please, report issues at https://github.com/CHIP-SPV/chipStar/issues.
warning: Ignored option '-arch'
In file included from main.cu:17:
./helper_math.h:1455:30: error: use of overloaded operator '/' is ambiguous (with operand types 'float2' (aka 'HIP_vector_type<float, 2>') and 'float2')
1455 | float2 y = clamp((x - a) / (b - a), 0.0f, 1.0f);
| ~~~~~~~ ^ ~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:498:5: note: candidate function
498 | operator/(HIP_vector_type x, const HIP_vector_type &y) noexcept {
| ^
./helper_math.h:980:35: note: candidate function
980 | inline __host__ __device__ float2 operator/(float2 a, float2 b)
| ^
./helper_math.h:1456:14: error: use of overloaded operator '*' is ambiguous (with operand types 'float2' (aka 'HIP_vector_type<float, 2>') and 'float2')
1456 | return (y*y*(make_float2(3.0f) - (make_float2(2.0f)*y)));
| ~^~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:485:61: note: candidate function
485 | friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator*(
| ^
./helper_math.h:751:35: note: candidate function
751 | inline __host__ __device__ float2 operator*(float2 a, float2 b)
| ^
./helper_math.h:1456:56: error: use of overloaded operator '*' is ambiguous (with operand types 'float2' (aka 'HIP_vector_type<float, 2>') and 'float2')
1456 | return (y*y*(make_float2(3.0f) - (make_float2(2.0f)*y)));
| ~~~~~~~~~~~~~~~~~^~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:485:61: note: candidate function
485 | friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator*(
| ^
./helper_math.h:751:35: note: candidate function
751 | inline __host__ __device__ float2 operator*(float2 a, float2 b)
| ^
./helper_math.h:1460:30: error: use of overloaded operator '/' is ambiguous (with operand types 'float3' (aka 'HIP_vector_type<float, 3>') and 'float3')
1460 | float3 y = clamp((x - a) / (b - a), 0.0f, 1.0f);
| ~~~~~~~ ^ ~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:498:5: note: candidate function
498 | operator/(HIP_vector_type x, const HIP_vector_type &y) noexcept {
| ^
./helper_math.h:1003:35: note: candidate function
1003 | inline __host__ __device__ float3 operator/(float3 a, float3 b)
| ^
./helper_math.h:1461:14: error: use of overloaded operator '*' is ambiguous (with operand types 'float3' (aka 'HIP_vector_type<float, 3>') and 'float3')
1461 | return (y*y*(make_float3(3.0f) - (make_float3(2.0f)*y)));
| ~^~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:485:61: note: candidate function
485 | friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator*(
| ^
./helper_math.h:820:35: note: candidate function
820 | inline __host__ __device__ float3 operator*(float3 a, float3 b)
| ^
./helper_math.h:1461:56: error: use of overloaded operator '*' is ambiguous (with operand types 'float3' (aka 'HIP_vector_type<float, 3>') and 'float3')
1461 | return (y*y*(make_float3(3.0f) - (make_float3(2.0f)*y)));
| ~~~~~~~~~~~~~~~~~^~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:485:61: note: candidate function
485 | friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator*(
| ^
./helper_math.h:820:35: note: candidate function
820 | inline __host__ __device__ float3 operator*(float3 a, float3 b)
| ^
./helper_math.h:1465:30: error: use of overloaded operator '/' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
1465 | float4 y = clamp((x - a) / (b - a), 0.0f, 1.0f);
| ~~~~~~~ ^ ~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:498:5: note: candidate function
498 | operator/(HIP_vector_type x, const HIP_vector_type &y) noexcept {
| ^
./helper_math.h:1028:35: note: candidate function
1028 | inline __host__ __device__ float4 operator/(float4 a, float4 b)
| ^
./helper_math.h:1466:14: error: use of overloaded operator '*' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
1466 | return (y*y*(make_float4(3.0f) - (make_float4(2.0f)*y)));
| ~^~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:485:61: note: candidate function
485 | friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator*(
| ^
./helper_math.h:895:35: note: candidate function
895 | inline __host__ __device__ float4 operator*(float4 a, float4 b)
| ^
./helper_math.h:1466:56: error: use of overloaded operator '*' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
1466 | return (y*y*(make_float4(3.0f) - (make_float4(2.0f)*y)));
| ~~~~~~~~~~~~~~~~~^~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:485:61: note: candidate function
485 | friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator*(
| ^
./helper_math.h:895:35: note: candidate function
895 | inline __host__ __device__ float4 operator*(float4 a, float4 b)
| ^
main.cu:124:21: error: use of overloaded operator '*' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
124 | f4Sum = top_color *
| ~~~~~~~~~ ^
125 | make_float4((float)iRadius, (float)iRadius, (float)iRadius, (float)iRadius);
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:485:61: note: candidate function
485 | friend __HOST_DEVICE__ inline constexpr HIP_vector_type operator*(
| ^
./helper_math.h:895:35: note: candidate function
895 | inline __host__ __device__ float4 operator*(float4 a, float4 b)
| ^
main.cu:128:11: error: use of overloaded operator '+=' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
128 | f4Sum += rgbaUintToFloat4(uiInputImage[y * uiWidth]);
| ~~~~~ ^ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:457:22: note: candidate function
457 | HIP_vector_type &operator+=(const HIP_vector_type &x) noexcept {
| ^
./helper_math.h:445:33: note: candidate function
445 | inline __host__ __device__ void operator+=(float4 &a, float4 b)
| ^
main.cu:134:11: error: use of overloaded operator '+=' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
134 | f4Sum += rgbaUintToFloat4(uiInputImage[(y + iRadius) * uiWidth]);
| ~~~~~ ^ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:457:22: note: candidate function
457 | HIP_vector_type &operator+=(const HIP_vector_type &x) noexcept {
| ^
./helper_math.h:445:33: note: candidate function
445 | inline __host__ __device__ void operator+=(float4 &a, float4 b)
| ^
main.cu:135:11: error: use of overloaded operator '-=' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
135 | f4Sum -= top_color;
| ~~~~~ ^ ~~~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:468:22: note: candidate function
468 | HIP_vector_type &operator-=(const HIP_vector_type &x) noexcept {
| ^
./helper_math.h:674:33: note: candidate function
674 | inline __host__ __device__ void operator-=(float4 &a, float4 b)
| ^
main.cu:141:11: error: use of overloaded operator '+=' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
141 | f4Sum += rgbaUintToFloat4(uiInputImage[(y + iRadius) * uiWidth]);
| ~~~~~ ^ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:457:22: note: candidate function
457 | HIP_vector_type &operator+=(const HIP_vector_type &x) noexcept {
| ^
./helper_math.h:445:33: note: candidate function
445 | inline __host__ __device__ void operator+=(float4 &a, float4 b)
| ^
main.cu:142:11: error: use of overloaded operator '-=' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
142 | f4Sum -= rgbaUintToFloat4(uiInputImage[((y - iRadius) * uiWidth) - uiWidth]);
| ~~~~~ ^ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:468:22: note: candidate function
468 | HIP_vector_type &operator-=(const HIP_vector_type &x) noexcept {
| ^
./helper_math.h:674:33: note: candidate function
674 | inline __host__ __device__ void operator-=(float4 &a, float4 b)
| ^
main.cu:148:11: error: use of overloaded operator '+=' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
148 | f4Sum += bot_color;
| ~~~~~ ^ ~~~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:457:22: note: candidate function
457 | HIP_vector_type &operator+=(const HIP_vector_type &x) noexcept {
| ^
./helper_math.h:445:33: note: candidate function
445 | inline __host__ __device__ void operator+=(float4 &a, float4 b)
| ^
main.cu:149:11: error: use of overloaded operator '-=' is ambiguous (with operand types 'float4' (aka 'HIP_vector_type<float, 4>') and 'float4')
149 | f4Sum -= rgbaUintToFloat4(uiInputImage[((y - iRadius) * uiWidth) - uiWidth]);
| ~~~~~ ^ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/tsaini.chen/install/chipStar/chipStar_06052024/include/hip/spirv_hip_vector_types.h:468:22: note: candidate function
468 | HIP_vector_type &operator-=(const HIP_vector_type &x) noexcept {
| ^
./helper_math.h:674:33: note: candidate function
674 | inline __host__ __device__ void operator-=(float4 &a, float4 b)
| ^
17 errors generated when compiling for .
Upon investigations, it appears that CUDA doesn't have built-in operator overloadings for vectors, and the CUDA programs need to define them (or include the helper_math.h that can be found in NVIDIA cuda-samples repo).
However, chipStar does have built-in vector operators here, from amd_detail/hip_vector_types.h. Therefore, when CUDA programs overload the vector operators, two definitions, one from CUDA program, one from spirv_hip_vector_types.h, are present, leading to compilation errors.
Here is a proposed solution, which wraps the operator overloadings in spirv_hip_vector_types.h with preprocessor directives (__CUDACC__) to prevent double defining them.
When compiling CUDA benchmark that defines vector operators itself (specifically, boxfilter-cuda in HeCBench that uses "helper_math.h" header file provided by NVIDIA cuda-samples), the following "use of overloaded operator ambiguous" compilation errors are produced.
Upon investigations, it appears that CUDA doesn't have built-in operator overloadings for vectors, and the CUDA programs need to define them (or include the helper_math.h that can be found in NVIDIA cuda-samples repo). However, chipStar does have built-in vector operators here, from amd_detail/hip_vector_types.h. Therefore, when CUDA programs overload the vector operators, two definitions, one from CUDA program, one from spirv_hip_vector_types.h, are present, leading to compilation errors.
Here is a proposed solution, which wraps the operator overloadings in spirv_hip_vector_types.h with preprocessor directives (
__CUDACC__) to prevent double defining them.