sidorovd updated this revision to Diff 162812. https://reviews.llvm.org/D51341
Files: include/clang/AST/Type.h lib/AST/Type.cpp lib/Sema/SemaOverload.cpp test/SemaOpenCL/half-double-overload.cl Index: test/SemaOpenCL/half-double-overload.cl =================================================================== --- /dev/null +++ test/SemaOpenCL/half-double-overload.cl @@ -0,0 +1,26 @@ +// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL1.2 + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +float __attribute__((overloadable)) foo(float in1, float in2); +float __attribute__((overloadable)) foo(half in1, half in2); + +int __attribute__((overloadable)) goo(float in1, float 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'}} +} Index: lib/Sema/SemaOverload.cpp =================================================================== --- lib/Sema/SemaOverload.cpp +++ lib/Sema/SemaOverload.cpp @@ -6059,6 +6059,33 @@ 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)) Index: lib/AST/Type.cpp =================================================================== --- lib/AST/Type.cpp +++ lib/AST/Type.cpp @@ -1919,6 +1919,13 @@ 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(); Index: include/clang/AST/Type.h =================================================================== --- include/clang/AST/Type.h +++ include/clang/AST/Type.h @@ -1924,6 +1924,7 @@ 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.
Index: test/SemaOpenCL/half-double-overload.cl =================================================================== --- /dev/null +++ test/SemaOpenCL/half-double-overload.cl @@ -0,0 +1,26 @@ +// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL1.2 + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +float __attribute__((overloadable)) foo(float in1, float in2); +float __attribute__((overloadable)) foo(half in1, half in2); + +int __attribute__((overloadable)) goo(float in1, float 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'}} +} Index: lib/Sema/SemaOverload.cpp =================================================================== --- lib/Sema/SemaOverload.cpp +++ lib/Sema/SemaOverload.cpp @@ -6059,6 +6059,33 @@ 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)) Index: lib/AST/Type.cpp =================================================================== --- lib/AST/Type.cpp +++ lib/AST/Type.cpp @@ -1919,6 +1919,13 @@ 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(); Index: include/clang/AST/Type.h =================================================================== --- include/clang/AST/Type.h +++ include/clang/AST/Type.h @@ -1924,6 +1924,7 @@ 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.
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits