Index: include/clang/Basic/Builtins.def =================================================================== --- include/clang/Basic/Builtins.def +++ include/clang/Basic/Builtins.def @@ -1517,6 +1517,58 @@ // 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 Index: lib/CodeGen/CGBuiltin.cpp =================================================================== --- lib/CodeGen/CGBuiltin.cpp +++ lib/CodeGen/CGBuiltin.cpp @@ -1450,6 +1450,139 @@ 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 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) { @@ -3878,6 +4011,58 @@ } 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 Index: lib/Sema/SemaChecking.cpp =================================================================== --- lib/Sema/SemaChecking.cpp +++ lib/Sema/SemaChecking.cpp @@ -925,6 +925,106 @@ 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()) { + BaseTy = BuiltinTy; + } else if (auto *VecTy = Ty->getAs()) { + BaseTy = VecTy->getElementType()->getAs(); + } + + 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 @@ -1420,6 +1520,60 @@ 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 Index: test/CodeGenOpenCL/builtin-math.cl =================================================================== --- /dev/null +++ test/CodeGenOpenCL/builtin-math.cl @@ -0,0 +1,24 @@ +// 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]] +} + Index: test/SemaOpenCL/builtin-math.cl =================================================================== --- /dev/null +++ test/SemaOpenCL/builtin-math.cl @@ -0,0 +1,27 @@ +// 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}} +} +