[OpenCL] Ambiguous function call.

[echuraev]

Summary

Added warning about potential ambiguity error with built-in overloading.

Patch by Alexey Bader

[Anastasia]

This change seems to modify normal C behavior again. Is there any strong motivation for doing this and if yes could it be done generically with C?

lib/Sema/SemaChecking.cpp

2479

How does it check that this is a built-in function?

[bader]

In D27334#611703, @Anastasia wrote:

This change seems to modify normal C behavior again. Is there any strong motivation for doing this and if yes could it be done generically with C?

Motivation:

// Non-portable OpenCL 1.2 code 
__kernel void foo(global float* out) {
  out[get_global_id(0)] = sin(get_global_id(0));
}

This program compiles fine on OpenCL platform w/o doubles support and fails otherwise.
If OpenCL driver supports doubles it provides at least two versions of 'sin' built-in math function and compiler will not be able to choose the right one for 'size_t' argument.
The goal of this warning is to let OpenCL developer know about potential issues with code portability.

This might be not so serious issue for C, since C doesn't support function overloading, whereas OpenCL built-in functions library must overload most of the functions.

[Anastasia]

In D27334#612858, @bader wrote:

In D27334#611703, @Anastasia wrote:

This change seems to modify normal C behavior again. Is there any strong motivation for doing this and if yes could it be done generically with C?

Motivation:

// Non-portable OpenCL 1.2 code 
__kernel void foo(global float* out) {
  out[get_global_id(0)] = sin(get_global_id(0));
}

This program compiles fine on OpenCL platform w/o doubles support and fails otherwise.
If OpenCL driver supports doubles it provides at least two versions of 'sin' built-in math function and compiler will not be able to choose the right one for 'size_t' argument.
The goal of this warning is to let OpenCL developer know about potential issues with code portability.

I would argue this improves the portability much as it can also be misleading in some situations (because it refers to a potentially hypothetical problem). For example there can be builtin functions that only have a float parameter (without a double version of it). This is for example the case with read_image functions that take a float coordinate value between 0 and 1. Unfortunately this warning won't be triggered on read_image functions because there is an overload candidate with an int type of the same parameter too. But we can't exclude this situations to appear in the future or from some vendor extensions or even custom OpenCL code.

[bader]

In D27334#613504, @Anastasia wrote:

In D27334#612858, @bader wrote:

In D27334#611703, @Anastasia wrote:

This change seems to modify normal C behavior again. Is there any strong motivation for doing this and if yes could it be done generically with C?

Motivation:

// Non-portable OpenCL 1.2 code 
__kernel void foo(global float* out) {
  out[get_global_id(0)] = sin(get_global_id(0));
}

This program compiles fine on OpenCL platform w/o doubles support and fails otherwise.
If OpenCL driver supports doubles it provides at least two versions of 'sin' built-in math function and compiler will not be able to choose the right one for 'size_t' argument.
The goal of this warning is to let OpenCL developer know about potential issues with code portability.

I would argue this improves the portability much as it can also be misleading in some situations (because it refers to a potentially hypothetical problem). For example there can be builtin functions that only have a float parameter (without a double version of it). This is for example the case with read_image functions that take a float coordinate value between 0 and 1. Unfortunately this warning won't be triggered on read_image functions because there is an overload candidate with an int type of the same parameter too. But we can't exclude this situations to appear in the future or from some vendor extensions or even custom OpenCL code.

As much as any other warning it's not always means that there is an error in the code. It just means that developer should inspect the construction triggering a warning.
Passing integer value to a function with floating point parameters is not always an error, but some times it might be so.
Do you suggest dropping the diagnostics at all or changing the diagnostics message?

[Anastasia]

In D27334#614389, @bader wrote:

In D27334#613504, @Anastasia wrote:

In D27334#612858, @bader wrote:

In D27334#611703, @Anastasia wrote:

This change seems to modify normal C behavior again. Is there any strong motivation for doing this and if yes could it be done generically with C?

Motivation:

// Non-portable OpenCL 1.2 code 
__kernel void foo(global float* out) {
  out[get_global_id(0)] = sin(get_global_id(0));
}

This program compiles fine on OpenCL platform w/o doubles support and fails otherwise.
If OpenCL driver supports doubles it provides at least two versions of 'sin' built-in math function and compiler will not be able to choose the right one for 'size_t' argument.
The goal of this warning is to let OpenCL developer know about potential issues with code portability.

I would argue this improves the portability much as it can also be misleading in some situations (because it refers to a potentially hypothetical problem). For example there can be builtin functions that only have a float parameter (without a double version of it). This is for example the case with read_image functions that take a float coordinate value between 0 and 1. Unfortunately this warning won't be triggered on read_image functions because there is an overload candidate with an int type of the same parameter too. But we can't exclude this situations to appear in the future or from some vendor extensions or even custom OpenCL code.

As much as any other warning it's not always means that there is an error in the code. It just means that developer should inspect the construction triggering a warning.
Passing integer value to a function with floating point parameters is not always an error, but some times it might be so.
Do you suggest dropping the diagnostics at all or changing the diagnostics message?

I agree warnings don't always signal a definite issue (even thought it's good to make them as precise as we can). We could try to reword the diagnostic message. However, the biggest issue I have here is that the message can be given in the situations that are unrelated to the problem (i.e. the overload candidates that don't have anything to do with the parameter being diagnosed or don't overload with the double precision). Therefore, it feels like the diagnostic can be confusing in some cases even though they are not very probable ones.

[echuraev]

In D27334#614826, @Anastasia wrote:

In D27334#614389, @bader wrote:

In D27334#613504, @Anastasia wrote:

In D27334#612858, @bader wrote:

In D27334#611703, @Anastasia wrote:

This change seems to modify normal C behavior again. Is there any strong motivation for doing this and if yes could it be done generically with C?

Motivation:

// Non-portable OpenCL 1.2 code 
__kernel void foo(global float* out) {
  out[get_global_id(0)] = sin(get_global_id(0));
}

This program compiles fine on OpenCL platform w/o doubles support and fails otherwise.
If OpenCL driver supports doubles it provides at least two versions of 'sin' built-in math function and compiler will not be able to choose the right one for 'size_t' argument.
The goal of this warning is to let OpenCL developer know about potential issues with code portability.

I would argue this improves the portability much as it can also be misleading in some situations (because it refers to a potentially hypothetical problem). For example there can be builtin functions that only have a float parameter (without a double version of it). This is for example the case with read_image functions that take a float coordinate value between 0 and 1. Unfortunately this warning won't be triggered on read_image functions because there is an overload candidate with an int type of the same parameter too. But we can't exclude this situations to appear in the future or from some vendor extensions or even custom OpenCL code.

As much as any other warning it's not always means that there is an error in the code. It just means that developer should inspect the construction triggering a warning.
Passing integer value to a function with floating point parameters is not always an error, but some times it might be so.
Do you suggest dropping the diagnostics at all or changing the diagnostics message?

I agree warnings don't always signal a definite issue (even thought it's good to make them as precise as we can). We could try to reword the diagnostic message. However, the biggest issue I have here is that the message can be given in the situations that are unrelated to the problem (i.e. the overload candidates that don't have anything to do with the parameter being diagnosed or don't overload with the double precision). Therefore, it feels like the diagnostic can be confusing in some cases even though they are not very probable ones.

@Anastasia, do you have any suggestions how is it better to reword the diagnostic message? Yes, this message can be given in some situations that are unrelated to the problem but in this case it will be a notification for developer that this function call can be potential ambiguous.

lib/Sema/SemaChecking.cpp

2479

In OpenCL we can overload only built-in functions. So, I think that we can recognize that the function is built-in by checking that the language is OpenCL and the function has Overloadable attribute.

[Anastasia]

I don't actually. But remembering the follow up discussion:
http://lists.llvm.org/pipermail/cfe-commits/Week-of-Mon-20161205/178846.html
and since we have to deviate from the standard C/C++ implementation anyways I am wondering whether modifying overloading resolution is a better way to go?

[echuraev]

In D27334#700521, @Anastasia wrote:

I don't actually. But remembering the follow up discussion:
http://lists.llvm.org/pipermail/cfe-commits/Week-of-Mon-20161205/178846.html
and since we have to deviate from the standard C/C++ implementation anyways I am wondering whether modifying overloading resolution is a better way to go?

Yes, I remember this discussion. But we saw this problem in the real OpenCL code and this warning is able to help developer to understand where his code could be ambiguous. Actually, I don't know the modifying overloading resolution is a better way to go or not. Do you have any suggestions how could we notify developer about this potential problem in better way?

[echuraev]

So, I think that we have to do some decision about this patch. @Anastasia, What do you think about it? Please see my comment above. What should we do with this patch?

[Anastasia]

In D27334#751921, @echuraev wrote:

So, I think that we have to do some decision about this patch. @Anastasia, What do you think about it? Please see my comment above. What should we do with this patch?

I am still not convinced adding the diagnostics in this shape is a good way to go....

include/clang/Basic/DiagnosticSemaKinds.td

 def error_opencl_cast_non_zero_to_event_t : Error<
   "cannot cast non-zero value '%0' to 'event_t'">;
 def err_opencl_global_invalid_addr_space : Error<
   "%select{program scope|static local|extern}0 variable must reside in %1 address space">;
 def err_missing_actual_pipe_type : Error<
   "missing actual type specifier for pipe">;
 def err_reference_pipe_type : Error <
   "pipes packet types cannot be of reference type">;
+def warn_ocl_bultin_potential_ambiguity : Warning<
+    "implicit conversion from integral type to floating point type for"
+    " overloadable function might lead to ambiguity">,
+  InGroup<DiagGroup<"conversion-might-lead-to-ambiguity">>, DefaultIgnore;
 def err_opencl_no_main : Error<"%select{function|kernel}0 cannot be called 'main'">;
 def err_opencl_kernel_attr :
   Error<"attribute %0 can only be applied to a kernel function">;
 def err_opencl_return_value_with_address_space : Error<
   "return value cannot be qualified with address space">;
 def err_opencl_constant_no_init : Error<
   "variable in constant address space must be initialized">;
 def err_atomic_init_constant : Error<

lib/Sema/SemaChecking.cpp

void Sema::checkCall(NamedDecl *FDecl, const FunctionProtoType *Proto,
 
   if (FDecl || Proto) {
     CheckNonNullArguments(*this, FDecl, Proto, Args, Loc);
 
     // Type safety checking.
     if (FDecl) {
       for (const auto *I : FDecl->specific_attrs<ArgumentWithTypeTagAttr>())
         CheckArgumentWithTypeTag(I, Args.data());
+
+      if (getLangOpts().OpenCL) {
+        // Check if overloadble built-in function with floating point arguments takes
+        // integer values.
+        if (FDecl->hasAttr<OverloadableAttr>()) {
+          for (const auto* Arg : Args) {
+            const ImplicitCastExpr *ICE = dyn_cast<ImplicitCastExpr>(Arg);
+            if (!ICE || ICE->getCastKind() != CK_IntegralToFloating)
+              continue;
+            Diag(Loc, diag::warn_ocl_bultin_potential_ambiguity) << Range;
+          }
+        }
+      }
     }
   }
 }
 
 /// CheckConstructorCall - Check a constructor call for correctness and safety
 /// properties not enforced by the C type system.
 void Sema::CheckConstructorCall(FunctionDecl *FDecl,
                                 ArrayRef<const Expr *> Args,

test/SemaOpenCL/warn-potential-abiguity.cl

+// RUN: %clang_cc1 -fsyntax-only -verify -pedantic -Wconversion-might-lead-to-ambiguity %s
+
+float __attribute__((overloadable)) ocl_builtin_func(float f) { return f; }
+float __attribute__((overloadable)) ocl_builtin_func_2args(float arg1, float arg2) { return arg1 + arg2; }
+
+__kernel void test() {
+  int p = ocl_builtin_func(3); // expected-warning {{implicit conversion from integral type to floating point type for overloadable function might lead to ambiguity}}
+  int q = ocl_builtin_func_2args(3.f, 3); // expected-warning {{implicit conversion from integral type to floating point type for overloadable function might lead to ambiguity}}
+}