asavonic created this revision.
Herald added subscribers: cfe-commits, kristina, Anastasia, yaxunl.
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.
Repository:
rC Clang
https://reviews.llvm.org/D52458
Files:
include/clang/Basic/Builtins.def
lib/CodeGen/CGBuiltin.cpp
lib/Sema/SemaChecking.cpp
test/CodeGenOpenCL/builtin-math.cl
test/SemaOpenCL/builtin-math.cl
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}}
+}
+
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: 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<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
@@ -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: 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<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) {
@@ -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: 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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
