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: 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: 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: 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
+
+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'}}
+}