[HEADER] Overloadable function candidates for half/double types

[sidorovd]

Summary

Overloadable function candidate should not be viable if it uses extensions (Half or Double type) that are not enabled or supported.

Patch by Egor Churaev

[Anastasia]

lib/Sema/SemaOverload.cpp

6063

I would imagine if extension isn't supported the candidate function with data type defined by extension shouldn't be available at all during compilation?

Also is there any good way to generalize this for all types and extensions including vendor ones that are added with the pragmas?
https://clang.llvm.org/docs/UsersManual.html#opencl-extensions

test/SemaOpenCL/half-double-overload.cl

8

I think it will be clearer if candidates with non-half parameters moved out of extension enabled block.

19

Wondering if better message could be:

candidate unavailable as it requires OpenCL extension to be disabled

Could it also print the extension name?

[Anastasia]

I think this is related to: https://reviews.llvm.org/D46501

[sidorovd]

lib/Sema/SemaOverload.cpp

6063

I would imagine if extension isn't supported the candidate function with data type defined by extension shouldn't be available at all during compilation?

Yes, you are right.

Also is there any good way to generalize this for all types and extensions including vendor ones that are added with the pragmas?

https://clang.llvm.org/docs/UsersManual.html#opencl-extensions

There might be a way but I can't properly answer this question right now. By default types and extensions aren't associated to each other.

test/SemaOpenCL/half-double-overload.cl

19

Sounds good. And I'd prefer to split it into another patch, because it would affect unrelated to the change tests and also require changes in Sema to allow extension name printing.

[Anastasia]

lib/Sema/SemaOverload.cpp

6063

We have a map of types and associated extensions with them in OpenCLTypeExtMap in Sema.h... not sure what would be the cost of such generic check though.

test/SemaOpenCL/half-double-overload.cl

19

Ok, makes sense for this to be a separate change!

[sidorovd]

For now this patch is on hold since it adds no new functionality. Its future will be decided after @asavonic finished https://reviews.llvm.org/D51544

[sidorovd]

@Anastasia @yaxunl
Hi, I am working on generalizing this patch and several questions have raised during this, so I want to discuss them with you:

1. Should #pragma OPENCL EXTENSION ext_name : begin enables the extension as well? For now I see it's not, as an example:

#pragma OPENCL EXTENSION cl_khr_fp16 : enable
half __attribute__((overloadable)) goo(half in1, half in2); // all ok
#pragma OPENCL EXTENSION cl_khr_fp16 : disable

#pragma OPENCL EXTENSION cl_khr_fp16 : begin
half __attribute__((overloadable)) goo(half in1, half in2); // declaring function parameter of type 'half' is not allowed; did you forget * ?
#pragma OPENCL EXTENSION cl_khr_fp16 : end

2. As far as I understand, when declaring an extension we shall have 1 #pragma begin and 1 #pragma end. But here is a test called extension-begin and in its header one can see this construction:

#pragma OPENCL EXTENSION all : begin
#pragma OPENCL EXTENSION all : end 

#pragma OPENCL EXTENSION my_ext : begin
///some code
#pragma OPENCL EXTENSION my_ext : end 
#pragma OPENCL EXTENSION my_ext : end // why?
}

so here my_ext has double ending. And in this way the test passes, but if I remove second ending (which is redundant from my perspective), I see following diagnostics: " OpenCL extension end directive mismatches begin directive - ignoring". Is it a bug or it's supposed to work that way?

[yaxunl]

In D51341#1301047, @sidorovd wrote:

@Anastasia @yaxunl
Hi, I am working on generalizing this patch and several questions have raised during this, so I want to discuss them with you:

1. Should #pragma OPENCL EXTENSION ext_name : begin enables the extension as well? For now I see it's not, as an example:

#pragma OPENCL EXTENSION cl_khr_fp16 : enable
half __attribute__((overloadable)) goo(half in1, half in2); // all ok
#pragma OPENCL EXTENSION cl_khr_fp16 : disable

#pragma OPENCL EXTENSION cl_khr_fp16 : begin
half __attribute__((overloadable)) goo(half in1, half in2); // declaring function parameter of type 'half' is not allowed; did you forget * ?
#pragma OPENCL EXTENSION cl_khr_fp16 : end

#pragma OPENCL EXTENSION ext_name : begin should not enable the extension. There are cases that you want to declare functions and types associated with an extension but do not want to enable the extension.

2. As far as I understand, when declaring an extension we shall have 1 #pragma begin and 1 #pragma end. But here is a test called extension-begin and in its header one can see this construction:

#pragma OPENCL EXTENSION all : begin
#pragma OPENCL EXTENSION all : end 

#pragma OPENCL EXTENSION my_ext : begin
///some code
#pragma OPENCL EXTENSION my_ext : end 
#pragma OPENCL EXTENSION my_ext : end // why?
}

so here my_ext has double ending. And in this way the test passes, but if I remove second ending (which is redundant from my perspective), I see following diagnostics: " OpenCL extension end directive mismatches begin directive - ignoring". Is it a bug or it's supposed to work that way?

It seems to be a bug.

include/clang/AST/Type.h

public:
   bool isFloat128Type() const;
   bool isRealType() const;         // C99 6.2.5p17 (real floating + integer)
   bool isArithmeticType() const;   // C99 6.2.5p18 (integer + floating)
   bool isVoidType() const;         // C99 6.2.5p19
   bool isScalarType() const;       // C99 6.2.5p21 (arithmetic + pointers)
   bool isAggregateType() const;
   bool isFundamentalType() const;
   bool isCompoundType() const;
+  bool isDoubleType() const;       // (double + long double)
 
   // Type Predicates: Check to see if this type is structurally the specified
   // type, ignoring typedefs and qualifiers.
   bool isFunctionType() const;
   bool isFunctionNoProtoType() const { return getAs<FunctionNoProtoType>(); }
   bool isFunctionProtoType() const { return getAs<FunctionProtoType>(); }
   bool isPointerType() const;
   bool isAnyPointerType() const;   // Any C pointer or ObjC object pointer

lib/AST/Type.cpp

bool Type::isFloatingType() const {
   if (const auto *BT = dyn_cast<BuiltinType>(CanonicalType))
     return BT->getKind() >= BuiltinType::Half &&
            BT->getKind() <= BuiltinType::Float128;
   if (const auto *CT = dyn_cast<ComplexType>(CanonicalType))
     return CT->getElementType()->isFloatingType();
   return false;
 }
 
+bool Type::isDoubleType() const {
+  if (const BuiltinType *BT = dyn_cast<BuiltinType>(CanonicalType))
+    return BT->getKind() >= BuiltinType::Double &&
+      BT->getKind() <= BuiltinType::LongDouble;
+  return false;
+}
+
 bool Type::hasFloatingRepresentation() const {
   if (const auto *VT = dyn_cast<VectorType>(CanonicalType))
     return VT->getElementType()->isFloatingType();
   else
     return isFloatingType();
 }
 
 bool Type::isRealFloatingType() const {

lib/Sema/SemaOverload.cpp

Sema::AddOverloadCandidate(FunctionDecl *Function,
   unsigned MinRequiredArgs = Function->getMinRequiredArguments();
   if (Args.size() < MinRequiredArgs && !PartialOverloading) {
     // Not enough arguments.
     Candidate.Viable = false;
     Candidate.FailureKind = ovl_fail_too_few_arguments;
     return;
   }
 
+  // OpenCL: A candidate function that uses extentions that are not enabled or
+  // supported is not viable.
+  if (getLangOpts().OpenCL) {
+    bool HasHalf =
+      getOpenCLOptions().isSupported("cl_khr_fp16", getLangOpts().OpenCLVersion)
+      && getOpenCLOptions().isEnabled("cl_khr_fp16");
+    bool HasDouble =
+      getOpenCLOptions().isSupported("cl_khr_fp64", getLangOpts().OpenCLVersion)
+      && getOpenCLOptions().isEnabled("cl_khr_fp64");
+
+    if ((!HasHalf && Function->getReturnType()->isHalfType()) ||
+        (!HasDouble && Function->getReturnType()->isDoubleType())) {
+      Candidate.Viable = false;
+      Candidate.FailureKind = ovl_fail_ext_disabled;
+      return;
+    }
+    for (const auto *PI : Function->parameters()) {
+      QualType PTy = PI->getType();
+      if ((!HasHalf && PTy->isHalfType()) ||
+          (!HasDouble && PTy->isDoubleType())) {
+        Candidate.Viable = false;
+        Candidate.FailureKind = ovl_fail_ext_disabled;
+        return;
+      }
+    }
+  }
+
   // (CUDA B.1): Check for invalid calls between targets.
   if (getLangOpts().CUDA)
     if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
       // Skip the check for callers that are implicit members, because in this
       // case we may not yet know what the member's target is; the target is
       // inferred for the member automatically, based on the bases and fields of
       // the class.
       if (!Caller->isImplicit() && !IsAllowedCUDACall(Caller, Function)) {

test/SemaOpenCL/half-double-overload.cl

+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL1.2
+
+float __attribute__((overloadable)) foo(float in1, float in2);
+int __attribute__((overloadable)) goo(float in1, float in2);
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+float __attribute__((overloadable)) foo(half in1, half in2);
+half __attribute__((overloadable)) goo(double in1, double in2);
+#pragma OPENCL EXTENSION cl_khr_fp16 : disable
+
+__kernel void vi(int x, int y) {
+  foo(x, y);
+  goo(x, y);
+}
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+float __attribute__((overloadable)) foo_err(half in1, half in2);
+// expected-note@-1 {{candidate disabled due to OpenCL extension}}
+float __attribute__((overloadable)) foo_err(half in1, int in2);
+// expected-note@-1 {{candidate disabled due to OpenCL extension}}
+#pragma OPENCL EXTENSION cl_khr_fp16 : disable
+
+__kernel void vi_err(int x, int y) {
+  foo_err(x, y);
+  // expected-error@-1 {{no matching function for call to 'foo_err'}}
+}