[OpenCL] Implement OpenCL math builtins: fp -> fp

[asavonic]

Summary

This patch adds OpenCL builtin functions that take any built-in
floating point type (float, double, half) and return a value of the
same floating point type.

include/clang/Basic/Builtins.def

 // Win64-compatible va_list functions
 BUILTIN(__builtin_ms_va_start, "vc*&.", "nt")
 BUILTIN(__builtin_ms_va_end, "vc*&", "n")
 BUILTIN(__builtin_ms_va_copy, "vc*&c*&", "n")
 
 // OpenCL v1.1/1.2/2.0 s6.2.3 - Explicit conversions
 LANGBUILTIN(__builtin_opencl_convert, "v*", "nt", ALL_OCLC_LANGUAGES)
 
+// OpenCL v1.1 s6.11.2, v1.2 s6.12.2, v2.0 s6.13.2 - Math functions
+LANGBUILTIN(__builtin_opencl_acos,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_acosh,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_acospi,   "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_asin,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_asinh,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_asinpi,   "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_atan,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_atanh,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_atanpi,   "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_cbrt,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_ceil,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_copysign, "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_cos,      "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_cosh,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_cospi,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_erfc,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_erf,      "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_exp,      "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_exp2,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_exp10,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_expm1,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_fabs,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_floor,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_ilogb,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_lgamma,   "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_log,      "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_log2,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_log10,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_log1p,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_logb,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_rint,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_round,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_rsqrt,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_sin,      "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_sinh,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_sinpi,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_sqrt,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_tan,      "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_tanh,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_tanpi,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_tgamma,   "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_trunk,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_degrees,  "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_radians,  "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_sign,     "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_isfinite, "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_isinf,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_isnan,    "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_isnormal, "v*", "nt", ALL_OCLC_LANGUAGES)
+LANGBUILTIN(__builtin_opencl_signbit,  "v*", "nt", ALL_OCLC_LANGUAGES)
+
 #undef BUILTIN
 #undef LIBBUILTIN
 #undef LANGBUILTIN

lib/CodeGen/CGBuiltin.cpp

Attribute::AttrKind Attrs[] = {
     Attribute::Convergent,
     Attribute::NoUnwind,
     Attribute::ReadNone
   };
 
   return emitOpenCLBuiltin(CGF, MangledName, E->getType(), Args, Attrs);
 }
 
+static StringRef getOpenCLMathBuiltinName(unsigned BuiltinID) {
+  switch (BuiltinID) {
+  case Builtin::BI__builtin_opencl_acos:
+    return "acos";
+  case Builtin::BI__builtin_opencl_acosh:
+    return "acosh";
+  case Builtin::BI__builtin_opencl_acospi:
+    return "acospi";
+  case Builtin::BI__builtin_opencl_asin:
+    return "asin";
+  case Builtin::BI__builtin_opencl_asinh:
+    return "asinh";
+  case Builtin::BI__builtin_opencl_asinpi:
+    return "asinpi";
+  case Builtin::BI__builtin_opencl_atan:
+    return "atan";
+  case Builtin::BI__builtin_opencl_atanh:
+    return "atanh";
+  case Builtin::BI__builtin_opencl_atanpi:
+    return "atanpi";
+  case Builtin::BI__builtin_opencl_cbrt:
+    return "cbrt";
+  case Builtin::BI__builtin_opencl_ceil:
+    return "ceil";
+  case Builtin::BI__builtin_opencl_copysign:
+    return "copysign";
+  case Builtin::BI__builtin_opencl_cos:
+    return "cos";
+  case Builtin::BI__builtin_opencl_cosh:
+    return "cosh";
+  case Builtin::BI__builtin_opencl_cospi:
+    return "cospi";
+  case Builtin::BI__builtin_opencl_erfc:
+    return "erfc";
+  case Builtin::BI__builtin_opencl_erf:
+    return "erf";
+  case Builtin::BI__builtin_opencl_exp:
+    return "exp";
+  case Builtin::BI__builtin_opencl_exp2:
+    return "exp2";
+  case Builtin::BI__builtin_opencl_exp10:
+    return "exp10";
+  case Builtin::BI__builtin_opencl_expm1:
+    return "expm1";
+  case Builtin::BI__builtin_opencl_fabs:
+    return "fabs";
+  case Builtin::BI__builtin_opencl_floor:
+    return "floor";
+  case Builtin::BI__builtin_opencl_ilogb:
+    return "ilogb";
+  case Builtin::BI__builtin_opencl_lgamma:
+    return "lgamma";
+  case Builtin::BI__builtin_opencl_log:
+    return "log";
+  case Builtin::BI__builtin_opencl_log2:
+    return "log2";
+  case Builtin::BI__builtin_opencl_log10:
+    return "log10";
+  case Builtin::BI__builtin_opencl_log1p:
+    return "log1p";
+  case Builtin::BI__builtin_opencl_logb:
+    return "logb";
+  case Builtin::BI__builtin_opencl_rint:
+    return "rint";
+  case Builtin::BI__builtin_opencl_round:
+    return "round";
+  case Builtin::BI__builtin_opencl_rsqrt:
+    return "rsqrt";
+  case Builtin::BI__builtin_opencl_sin:
+    return "sin";
+  case Builtin::BI__builtin_opencl_sinh:
+    return "sinh";
+  case Builtin::BI__builtin_opencl_sinpi:
+    return "sinpi";
+  case Builtin::BI__builtin_opencl_sqrt:
+    return "sqrt";
+  case Builtin::BI__builtin_opencl_tan:
+    return "tan";
+  case Builtin::BI__builtin_opencl_tanh:
+    return "tanh";
+  case Builtin::BI__builtin_opencl_tanpi:
+    return "tanpi";
+  case Builtin::BI__builtin_opencl_tgamma:
+    return "tgamma";
+  case Builtin::BI__builtin_opencl_trunk:
+    return "trunk";
+  case Builtin::BI__builtin_opencl_degrees:
+    return "degrees";
+  case Builtin::BI__builtin_opencl_radians:
+    return "radians";
+  case Builtin::BI__builtin_opencl_sign:
+    return "sign";
+  case Builtin::BI__builtin_opencl_isfinite:
+    return "isfinite";
+  case Builtin::BI__builtin_opencl_isinf:
+    return "isinf";
+  case Builtin::BI__builtin_opencl_isnan:
+    return "isnan";
+  case Builtin::BI__builtin_opencl_isnormal:
+    return "isnormal";
+  case Builtin::BI__builtin_opencl_signbit:
+    return "signbit";
+  default:
+    llvm_unreachable("Unknown OpenCL builtin");
+  }
+}
+
+static RValue emitOpenCLMathBuiltin(CodeGenFunction &CGF, unsigned BuiltinID,
+                                    const CallExpr *E) {
+  ASTContext &Context = CGF.getContext();
+  StringRef Name = getOpenCLMathBuiltinName(BuiltinID);
+
+  SmallString<128> MangledName;
+  SmallVector<QualType, 8> ArgTys;
+  for (const auto Arg : E->arguments()) {
+    ArgTys.push_back(Arg->getType());
+  }
+  mangleOpenCLBuiltin(Context, Name, E->getType(), ArgTys, MangledName);
+
+  CallArgList Args;
+  for (const auto Arg : E->arguments()) {
+    Args.add(RValue::get(CGF.EmitScalarExpr(Arg)), Arg->getType());
+  }
+
+  Attribute::AttrKind Attrs[] = {
+    Attribute::Convergent,
+    Attribute::NoUnwind,
+    Attribute::ReadNone
+  };
+
+  return emitOpenCLBuiltin(CGF, MangledName, E->getType(), Args, Attrs);
+}
+
 RValue CodeGenFunction::EmitBuiltinExpr(const FunctionDecl *FD,
                                         unsigned BuiltinID, const CallExpr *E,
                                         ReturnValueSlot ReturnValue) {
   // See if we can constant fold this builtin.  If so, don't emit it at all.
   Expr::EvalResult Result;
   if (E->EvaluateAsRValue(Result, CGM.getContext()) &&
       !Result.hasSideEffects()) {
     if (Result.Val.isInt())
case Builtin::BI__builtin_ms_va_copy: {
     SrcAddr = Address(Builder.CreateBitCast(SrcAddr.getPointer(), BPP, "ap"),
                       SrcAddr.getAlignment());
 
     Value *ArgPtr = Builder.CreateLoad(SrcAddr, "ap.val");
     return RValue::get(Builder.CreateStore(ArgPtr, DestAddr));
   }
   case Builtin::BI__builtin_opencl_convert:
     return emitOpenCLConvert(*this, getContext(), E);
+
+  case Builtin::BI__builtin_opencl_acos:
+  case Builtin::BI__builtin_opencl_acosh:
+  case Builtin::BI__builtin_opencl_acospi:
+  case Builtin::BI__builtin_opencl_asin:
+  case Builtin::BI__builtin_opencl_asinh:
+  case Builtin::BI__builtin_opencl_asinpi:
+  case Builtin::BI__builtin_opencl_atan:
+  case Builtin::BI__builtin_opencl_atanh:
+  case Builtin::BI__builtin_opencl_atanpi:
+  case Builtin::BI__builtin_opencl_cbrt:
+  case Builtin::BI__builtin_opencl_ceil:
+  case Builtin::BI__builtin_opencl_copysign:
+  case Builtin::BI__builtin_opencl_cos:
+  case Builtin::BI__builtin_opencl_cosh:
+  case Builtin::BI__builtin_opencl_cospi:
+  case Builtin::BI__builtin_opencl_erfc:
+  case Builtin::BI__builtin_opencl_erf:
+  case Builtin::BI__builtin_opencl_exp:
+  case Builtin::BI__builtin_opencl_exp2:
+  case Builtin::BI__builtin_opencl_exp10:
+  case Builtin::BI__builtin_opencl_expm1:
+  case Builtin::BI__builtin_opencl_fabs:
+  case Builtin::BI__builtin_opencl_floor:
+  case Builtin::BI__builtin_opencl_ilogb:
+  case Builtin::BI__builtin_opencl_lgamma:
+  case Builtin::BI__builtin_opencl_log:
+  case Builtin::BI__builtin_opencl_log2:
+  case Builtin::BI__builtin_opencl_log10:
+  case Builtin::BI__builtin_opencl_log1p:
+  case Builtin::BI__builtin_opencl_logb:
+  case Builtin::BI__builtin_opencl_rint:
+  case Builtin::BI__builtin_opencl_round:
+  case Builtin::BI__builtin_opencl_rsqrt:
+  case Builtin::BI__builtin_opencl_sin:
+  case Builtin::BI__builtin_opencl_sinh:
+  case Builtin::BI__builtin_opencl_sinpi:
+  case Builtin::BI__builtin_opencl_sqrt:
+  case Builtin::BI__builtin_opencl_tan:
+  case Builtin::BI__builtin_opencl_tanh:
+  case Builtin::BI__builtin_opencl_tanpi:
+  case Builtin::BI__builtin_opencl_tgamma:
+  case Builtin::BI__builtin_opencl_trunk:
+  case Builtin::BI__builtin_opencl_degrees:
+  case Builtin::BI__builtin_opencl_radians:
+  case Builtin::BI__builtin_opencl_sign:
+  case Builtin::BI__builtin_opencl_isfinite:
+  case Builtin::BI__builtin_opencl_isinf:
+  case Builtin::BI__builtin_opencl_isnan:
+  case Builtin::BI__builtin_opencl_isnormal:
+  case Builtin::BI__builtin_opencl_signbit:
+    return emitOpenCLMathBuiltin(*this, BuiltinID, E);
   }
 
   // If this is an alias for a lib function (e.g. __builtin_sin), emit
   // the call using the normal call path, but using the unmangled
   // version of the function name.
   if (getContext().BuiltinInfo.isLibFunction(BuiltinID))
     return emitLibraryCall(*this, FD, E,
                            CGM.getBuiltinLibFunction(FD, BuiltinID));

lib/Sema/SemaChecking.cpp

static bool checkOpenCLConvertBuiltin(Sema &S, CallExpr *Call) {
 
   // Now fixup the return value type, otherwise Clang will attempt to cast from
   // nothing (since the builtin signature is meaningless) to DstTy.
   Call->setType(DstTy);
 
   return false;
 }
 
+static bool checkOpenCLFPScalarOrVectorArg(Sema &S, CallExpr *Call,
+                                           unsigned ArgNum) {
+  Expr *Arg = Call->getArg(0);
+
+  // Perform an implicit conversion from integer
+  if (Arg->getType()->isIntegerType()) {
+    ExprResult ArgExpr(Arg);
+    Sema::AssignConvertType ConvertType = S.CheckSingleAssignmentConstraints(
+        S.getASTContext().FloatTy, ArgExpr);
+    if (Sema::Incompatible != ConvertType) {
+      // Replace an argument with an implicit conversion to float.
+      Call->setArg(0, ArgExpr.get());
+      Arg = ArgExpr.get();
+    }
+  }
+
+  QualType Ty = Arg->getType();
+
+  const BuiltinType *BaseTy = nullptr;
+  if (auto *BuiltinTy = Ty->getAs<BuiltinType>()) {
+    BaseTy = BuiltinTy;
+  } else if (auto *VecTy = Ty->getAs<ExtVectorType>()) {
+    BaseTy = VecTy->getElementType()->getAs<BuiltinType>();
+  }
+
+  if (!BaseTy || !BaseTy->isFloatingPoint()) {
+    S.Diag(Arg->getBeginLoc(), diag::err_opencl_builtin_expected_type)
+        << Call->getDirectCallee() << "scalar or vector floating point";
+    return true;
+  }
+
+  return false;
+}
+
+static bool checkOpenCLMathBuiltin(Sema &S, unsigned BuiltinID,
+                                   CallExpr *Call) {
+  switch (BuiltinID) {
+  case Builtin::BI__builtin_opencl_acos:
+  case Builtin::BI__builtin_opencl_acosh:
+  case Builtin::BI__builtin_opencl_acospi:
+  case Builtin::BI__builtin_opencl_asin:
+  case Builtin::BI__builtin_opencl_asinh:
+  case Builtin::BI__builtin_opencl_asinpi:
+  case Builtin::BI__builtin_opencl_atan:
+  case Builtin::BI__builtin_opencl_atanh:
+  case Builtin::BI__builtin_opencl_atanpi:
+  case Builtin::BI__builtin_opencl_cbrt:
+  case Builtin::BI__builtin_opencl_ceil:
+  case Builtin::BI__builtin_opencl_copysign:
+  case Builtin::BI__builtin_opencl_cos:
+  case Builtin::BI__builtin_opencl_cosh:
+  case Builtin::BI__builtin_opencl_cospi:
+  case Builtin::BI__builtin_opencl_erfc:
+  case Builtin::BI__builtin_opencl_erf:
+  case Builtin::BI__builtin_opencl_exp:
+  case Builtin::BI__builtin_opencl_exp2:
+  case Builtin::BI__builtin_opencl_exp10:
+  case Builtin::BI__builtin_opencl_expm1:
+  case Builtin::BI__builtin_opencl_fabs:
+  case Builtin::BI__builtin_opencl_floor:
+  case Builtin::BI__builtin_opencl_ilogb:
+  case Builtin::BI__builtin_opencl_lgamma:
+  case Builtin::BI__builtin_opencl_log:
+  case Builtin::BI__builtin_opencl_log2:
+  case Builtin::BI__builtin_opencl_log10:
+  case Builtin::BI__builtin_opencl_log1p:
+  case Builtin::BI__builtin_opencl_logb:
+  case Builtin::BI__builtin_opencl_rint:
+  case Builtin::BI__builtin_opencl_round:
+  case Builtin::BI__builtin_opencl_rsqrt:
+  case Builtin::BI__builtin_opencl_sin:
+  case Builtin::BI__builtin_opencl_sinh:
+  case Builtin::BI__builtin_opencl_sinpi:
+  case Builtin::BI__builtin_opencl_sqrt:
+  case Builtin::BI__builtin_opencl_tan:
+  case Builtin::BI__builtin_opencl_tanh:
+  case Builtin::BI__builtin_opencl_tanpi:
+  case Builtin::BI__builtin_opencl_tgamma:
+  case Builtin::BI__builtin_opencl_trunk:
+  case Builtin::BI__builtin_opencl_degrees:
+  case Builtin::BI__builtin_opencl_radians:
+  case Builtin::BI__builtin_opencl_sign:
+  case Builtin::BI__builtin_opencl_isfinite:
+  case Builtin::BI__builtin_opencl_isinf:
+  case Builtin::BI__builtin_opencl_isnan:
+  case Builtin::BI__builtin_opencl_isnormal:
+  case Builtin::BI__builtin_opencl_signbit: {
+    if (checkArgCount(S, Call, 1) ||
+        checkOpenCLFPScalarOrVectorArg(S, Call, 0))
+      return true;
+    Call->setType(Call->getArg(0)->getType());
+    break;
+  }
+  default:
+    llvm_unreachable("Unexpected builtin");
+  }
+
+  return false;
+}
+
 // Emit an error and return true if the current architecture is not in the list
 // of supported architectures.
 static bool
 CheckBuiltinTargetSupport(Sema &S, unsigned BuiltinID, CallExpr *TheCall,
                           ArrayRef<llvm::Triple::ArchType> SupportedArchs) {
   llvm::Triple::ArchType CurArch =
       S.getASTContext().getTargetInfo().getTriple().getArch();
   if (llvm::is_contained(SupportedArchs, CurArch))
case Builtin::BI__builtin_os_log_format_buffer_size:
     if (SemaBuiltinOSLogFormat(TheCall))
       return ExprError();
     break;
   case Builtin::BI__builtin_opencl_convert: {
     if (checkOpenCLConvertBuiltin(*this, TheCall))
       return ExprError();
     break;
   }
+  case Builtin::BI__builtin_opencl_acos:
+  case Builtin::BI__builtin_opencl_acosh:
+  case Builtin::BI__builtin_opencl_acospi:
+  case Builtin::BI__builtin_opencl_asin:
+  case Builtin::BI__builtin_opencl_asinh:
+  case Builtin::BI__builtin_opencl_asinpi:
+  case Builtin::BI__builtin_opencl_atan:
+  case Builtin::BI__builtin_opencl_atanh:
+  case Builtin::BI__builtin_opencl_atanpi:
+  case Builtin::BI__builtin_opencl_cbrt:
+  case Builtin::BI__builtin_opencl_ceil:
+  case Builtin::BI__builtin_opencl_copysign:
+  case Builtin::BI__builtin_opencl_cos:
+  case Builtin::BI__builtin_opencl_cosh:
+  case Builtin::BI__builtin_opencl_cospi:
+  case Builtin::BI__builtin_opencl_erfc:
+  case Builtin::BI__builtin_opencl_erf:
+  case Builtin::BI__builtin_opencl_exp:
+  case Builtin::BI__builtin_opencl_exp2:
+  case Builtin::BI__builtin_opencl_exp10:
+  case Builtin::BI__builtin_opencl_expm1:
+  case Builtin::BI__builtin_opencl_fabs:
+  case Builtin::BI__builtin_opencl_floor:
+  case Builtin::BI__builtin_opencl_ilogb:
+  case Builtin::BI__builtin_opencl_lgamma:
+  case Builtin::BI__builtin_opencl_log:
+  case Builtin::BI__builtin_opencl_log2:
+  case Builtin::BI__builtin_opencl_log10:
+  case Builtin::BI__builtin_opencl_log1p:
+  case Builtin::BI__builtin_opencl_logb:
+  case Builtin::BI__builtin_opencl_rint:
+  case Builtin::BI__builtin_opencl_round:
+  case Builtin::BI__builtin_opencl_rsqrt:
+  case Builtin::BI__builtin_opencl_sin:
+  case Builtin::BI__builtin_opencl_sinh:
+  case Builtin::BI__builtin_opencl_sinpi:
+  case Builtin::BI__builtin_opencl_sqrt:
+  case Builtin::BI__builtin_opencl_tan:
+  case Builtin::BI__builtin_opencl_tanh:
+  case Builtin::BI__builtin_opencl_tanpi:
+  case Builtin::BI__builtin_opencl_tgamma:
+  case Builtin::BI__builtin_opencl_trunk:
+  case Builtin::BI__builtin_opencl_degrees:
+  case Builtin::BI__builtin_opencl_radians:
+  case Builtin::BI__builtin_opencl_sign:
+  case Builtin::BI__builtin_opencl_isfinite:
+  case Builtin::BI__builtin_opencl_isinf:
+  case Builtin::BI__builtin_opencl_isnan:
+  case Builtin::BI__builtin_opencl_isnormal:
+  case Builtin::BI__builtin_opencl_signbit: {
+    if (checkOpenCLMathBuiltin(*this, BuiltinID, TheCall))
+      return ExprError();
+    break;
+  }
   }
 
   // Since the target specific builtins for each arch overlap, only check those
   // of the arch we are compiling for.
   if (Context.BuiltinInfo.isTSBuiltin(BuiltinID)) {
     switch (Context.getTargetInfo().getTriple().getArch()) {
       case llvm::Triple::arm:
       case llvm::Triple::armeb:

test/CodeGenOpenCL/builtin-math.cl

+// RUN: %clang_cc1 -triple spir -cl-std=CL1.2 -O0 -emit-llvm -o - %s | FileCheck %s
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void acos_good() {
+  float i1 = 0.42f;
+  float o1 = __builtin_opencl_acos(i1);
+  // CHECK: call float @_Z4acosf(float %{{.*}}) #[[ATTR:[0-9+]]]
+
+  int   i2 = 42;
+  float o2 = __builtin_opencl_acos(i2);
+  // CHECK: %[[INTCONV:.*]] = sitofp i32 %{{.*}} to float
+  // CHECK: call float @_Z4acosf(float %[[INTCONV]]) #[[ATTR]]
+
+  half  i3 = 42.0;
+  half  o3 = __builtin_opencl_acos(i3);
+  // CHECK: call half @_Z4acosDh(half %{{.*}}) #[[ATTR]]
+
+  double i4 = 42.0;
+  double o4 = __builtin_opencl_acos(i4);
+  // CHECK: call double @_Z4acosd(double %{{.*}}) #[[ATTR]]
+}
+

test/SemaOpenCL/builtin-math.cl

+// RUN: %clang_cc1 -verify -pedantic -fsyntax-only -triple spir -cl-std=CL1.2 %s
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void acos_good() {
+  float i1 = 0.42f;
+  float o1 = __builtin_opencl_acos(i1);
+
+  int   i2 = 42;
+  float o2 = __builtin_opencl_acos(i2);
+
+  half  i3 = 42.0;
+  half  o3 = __builtin_opencl_acos(i3);
+
+  double i4 = 42.0;
+  double o4 = __builtin_opencl_acos(i4);
+}
+
+__kernel void acos_bad() {
+  __builtin_opencl_acos(0.0f, 0.0f); // expected-error{{too many arguments to function call, expected 1, have 2}}
+  __builtin_opencl_acos();           // expected-error{{too few arguments to function call, expected 1, have 0}}
+
+  struct {float f;} i1 = {42.0f};
+  float o1 = __builtin_opencl_acos(i1); // expected-error{{illegal call to '__builtin_opencl_acos', expected scalar or vector floating point argument type}}
+}
+