Index: include/clang/Basic/Builtins.def =================================================================== --- include/clang/Basic/Builtins.def +++ include/clang/Basic/Builtins.def @@ -1514,6 +1514,9 @@ 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) + #undef BUILTIN #undef LIBBUILTIN #undef LANGBUILTIN Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -8676,6 +8676,11 @@ def err_opencl_builtin_expected_type : Error< "illegal call to %0, expected %1 argument type">; +def err_opencl_builtin_convert_type : Error< + "conversion is available only from OpenCL built-in scalar or vector type">; +def err_opencl_builtin_convert_num_elements : Error< + "operand and result type must have the same number of elements">; + // OpenCL v2.2 s2.1.2.3 - Vector Component Access def ext_opencl_ext_vector_type_rgba_selector: ExtWarn< "vector component name '%0' is an OpenCL version 2.2 feature">, Index: lib/CodeGen/CGBuiltin.cpp =================================================================== --- lib/CodeGen/CGBuiltin.cpp +++ lib/CodeGen/CGBuiltin.cpp @@ -1267,6 +1267,189 @@ return RValue::get(Builder.CreateCall(F, { Src, Src, ShiftAmt })); } +/// Create a dummy function decl that is suitable for Mangler. +static FunctionDecl * +createOpenCLBuiltinFunctionDecl(ASTContext &Context, StringRef Name, + QualType RetTy, ArrayRef ArgTys) { + QualType FTy = + Context.getFunctionType(RetTy, ArgTys, + FunctionProtoType::ExtProtoInfo()); + + auto *FD = + FunctionDecl::Create(Context, + Context.getTranslationUnitDecl(), + /*StartLoc=*/ SourceLocation(), + /*NLoc=*/ SourceLocation(), + &Context.Idents.get(Name), + FTy, + /*TInfo=*/ nullptr, + SC_Extern, + /*isInlineSpecified=*/ false, + /*hasWrittenPrototype=*/ true); + + SmallVector Params; + for (auto ArgTy : ArgTys) { + Params.push_back( + ParmVarDecl::Create( + Context, FD, + /*StartLoc=*/ SourceLocation(), + /*IdLoc=*/ SourceLocation(), + /*Id=*/ nullptr, + ArgTy, + /*TInfo=*/ nullptr, + SC_None, + /*DefArg=*/ nullptr)); + } + + FD->setParams(Params); + FD->addAttr(OverloadableAttr::CreateImplicit(Context)); + return FD; +} + +static void mangleOpenCLBuiltin(ASTContext &Context, + StringRef Name, + QualType RetTy, ArrayRef ArgTys, + SmallVectorImpl &Result) { + const FunctionDecl *FD = + createOpenCLBuiltinFunctionDecl(Context, Name, RetTy, ArgTys); + + std::unique_ptr MC(Context.createMangleContext()); + llvm::raw_svector_ostream Out(Result); + + MC->mangleName(FD, Out); +} + +static RValue emitOpenCLBuiltin(CodeGenFunction &CGF, StringRef Name, + QualType RetTy, CallArgList Args, + ArrayRef Attrs) { + + const CGFunctionInfo &FuncInfo = + CGF.CGM.getTypes().arrangeBuiltinFunctionCall(RetTy, Args); + + llvm::FunctionType *FTy = CGF.CGM.getTypes().GetFunctionType(FuncInfo); + llvm::Function *Func = + cast(CGF.CGM.CreateRuntimeFunction(FTy, Name)); + + for (auto Attr : Attrs) { + Func->addFnAttr(Attr); + } + + unsigned DefaultCC = CGF.CGM.getTypes().ClangCallConvToLLVMCallConv( + CGF.getContext().getDefaultCallingConvention( + /*IsVariadic=*/false, /*IsCXXMethod=*/false)); + + Func->setCallingConv(DefaultCC); + + return CGF.EmitCall(FuncInfo, CGCallee::forDirect(Func), + ReturnValueSlot(), Args); +} + +enum class OpenCLConvertRounding { + Default = 0, + RTE = 1, + RTZ = 2, + RTP = 3, + RTN = 4, + SAT = 5, + SATRTE = 6, + SATRTZ = 7, + SATRTP = 8, + SATRTN = 9, +}; + +static StringRef getOpenCLConvertRoundingSuffix(OpenCLConvertRounding R) { + switch (R) { + case OpenCLConvertRounding::Default : return ""; + case OpenCLConvertRounding::RTE : return "_rte"; + case OpenCLConvertRounding::RTZ : return "_rtz"; + case OpenCLConvertRounding::RTP : return "_rtp"; + case OpenCLConvertRounding::RTN : return "_rtn"; + case OpenCLConvertRounding::SAT : return "_sat"; + case OpenCLConvertRounding::SATRTE : return "_sat_rte"; + case OpenCLConvertRounding::SATRTZ : return "_sat_rtz"; + case OpenCLConvertRounding::SATRTP : return "_sat_rtp"; + case OpenCLConvertRounding::SATRTN : return "_sat_rtn"; + } + llvm_unreachable("Invalid rounding mode"); +} + +static bool checkOpenCLConvertTypeName(StringRef Name) { + StringRef AllowedTypes[] = { + "char", "uchar", "short", "ushort", "int", "uint", "long", "ulong", + "float", "double", "half", + }; + StringRef AllowedSizes[] = { + "", "2", "3", "4", "8", "16" + }; + + SmallString<16> Allowed; + for (auto Ty : AllowedTypes) { + for (auto Size : AllowedSizes) { + Allowed = Ty; + Allowed += Size; + if (Name == Allowed) { + return true; + } + } + } + + return false; +} + +static RValue emitOpenCLConvert(CodeGenFunction &CGF, ASTContext &Context, + const CallExpr *E) { + const Expr *Src = E->getArg(0); + const Expr *DstTypePlaceholder = E->getArg(1); + const Expr *RoundingMode = E->getArg(2); + + const QualType SrcTy = Src->getType(); + const QualType DstTy = DstTypePlaceholder->getType(); + + // Find a destination type and its name: it is expected to be one of OpenCL + // builtin types. + StringRef TypeName; + if (const auto *Ty = DstTy->getAs()) { + TypeName = Ty->getName(Context.getPrintingPolicy()); + } else { + const auto *Typedef = DstTy->castAs(); + NamedDecl *D = cast(Typedef->getDecl()); + TypeName = D->getName(); + } + assert(checkOpenCLConvertTypeName(TypeName) && "Unexpected type name"); + + llvm::APSInt IntRoundingMode; + if (!RoundingMode->isIntegerConstantExpr(IntRoundingMode, Context)) { + llvm_unreachable("Rounding mode should be an integer constant"); + } + StringRef RoundingSuffix = getOpenCLConvertRoundingSuffix( + static_cast(IntRoundingMode.getZExtValue())); + + // Format a function name to match the one defined in OpenCL + // specification. For example, for DstTy == uchar4 and RoundingSuffix == RTZ + // we want function to be named: convert_uchar4_rtz(...) + SmallString<32> Name("convert_"); + Name += TypeName; + Name += RoundingSuffix; + + // Function `convert_xxx()' can take parameters of different types, so it must + // be mangled. We use a full featured mangler here, so that `convert' is + // mangled the same way as other OpenCL built-ins (which are defined in + // opencl-c.h). + SmallString<128> MangledName; + mangleOpenCLBuiltin(Context, Name, DstTy, {SrcTy}, MangledName); + + CallArgList Args; + Args.add(RValue::get(CGF.EmitScalarExpr(Src)), SrcTy); + + 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) { @@ -3693,6 +3876,8 @@ 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); } // 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 @@ -877,6 +877,54 @@ return false; } +// Check if an argument is a built-in scalar or vector data type with correct +// vector size. Fixup return type to match the placeholder argument. +static bool checkOpenCLConvertBuiltin(Sema &S, CallExpr *Call) { + if (checkArgCount(S, Call, 3)) + return true; + + const Expr *SrcArg = Call->getArg(0); + const Expr *DstTypePlaceholder = Call->getArg(1); + const Expr *RoundingMode = Call->getArg(2); + + const QualType SrcTy = SrcArg->getType(); + const QualType DstTy = DstTypePlaceholder->getType(); + + QualType SrcElem = SrcTy; + unsigned SrcVecSize = 1; + if (const auto *SrcVecTy = SrcTy->getAs()) { + SrcElem = SrcVecTy->getElementType(); + SrcVecSize = SrcVecTy->getNumElements(); + } + + unsigned DstVecSize = 1; + if (const auto *DstVecTy = DstTy->getAs()) { + DstVecSize = DstVecTy->getNumElements(); + } + + if (!SrcElem->isBuiltinType() || SrcElem->isOpenCLSpecificType()) { + S.Diag(SrcArg->getBeginLoc(), diag::err_opencl_builtin_convert_type); + return true; + } + + if (SrcVecSize != DstVecSize) { + S.Diag(SrcArg->getBeginLoc(), diag::err_opencl_builtin_convert_num_elements); + return true; + } + + if (!RoundingMode->getType()->isIntegerType()) { + S.Diag(RoundingMode->getBeginLoc(), diag::err_opencl_builtin_expected_type) + << Call->getDirectCallee() << "integer"; + return true; + } + + // 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; +} + // Emit an error and return true if the current architecture is not in the list // of supported architectures. static bool @@ -1367,6 +1415,11 @@ if (SemaBuiltinOSLogFormat(TheCall)) return ExprError(); break; + case Builtin::BI__builtin_opencl_convert: { + if (checkOpenCLConvertBuiltin(*this, TheCall)) + return ExprError(); + break; + } } // Since the target specific builtins for each arch overlap, only check those Index: test/CodeGenOpenCL/builtin-convert.cl =================================================================== --- /dev/null +++ test/CodeGenOpenCL/builtin-convert.cl @@ -0,0 +1,69 @@ +// RUN: %clang_cc1 -emit-llvm -o - -triple spir -disable-llvm-passes %s | FileCheck %s + +typedef char char2 __attribute__((ext_vector_type(2))); +typedef int int2 __attribute__((ext_vector_type(2))); + +#define __OPENCL_CONVERT_DEFAULT__ 0 +#define __OPENCL_CONVERT_RTE__ 1 +#define __OPENCL_CONVERT_RTZ__ 2 +#define __OPENCL_CONVERT_RTP__ 3 +#define __OPENCL_CONVERT_RTN__ 4 +#define __OPENCL_CONVERT_SAT__ 5 +#define __OPENCL_CONVERT_SATRTE__ 6 +#define __OPENCL_CONVERT_SATRTZ__ 7 +#define __OPENCL_CONVERT_SATRTP__ 8 +#define __OPENCL_CONVERT_SATRTN__ 9 + +#define convert_char(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_DEFAULT__) +#define convert_char2(x) __builtin_opencl_convert((x), (char2)0, __OPENCL_CONVERT_DEFAULT__) + +#define convert_char_rte(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_RTE__) +#define convert_char_rtz(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_RTZ__) +#define convert_char_rtp(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_RTP__) +#define convert_char_rtn(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_RTN__) +#define convert_char_sat(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_SAT__) +#define convert_char_sat_rte(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_SATRTE__) +#define convert_char_sat_rtz(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_SATRTZ__) +#define convert_char_sat_rtp(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_SATRTP__) +#define convert_char_sat_rtn(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_SATRTN__) + +__kernel void good() { + int i = 0; + char o; + + o = convert_char_rte(i); + o = convert_char_rtz(i); + o = convert_char_rtp(i); + o = convert_char_rtn(i); + o = convert_char_sat(i); + o = convert_char_sat_rte(i); + o = convert_char_sat_rtz(i); + o = convert_char_sat_rtp(i); + o = convert_char_sat_rtn(i); + + // CHECK: call signext i8 @_Z16convert_char_rtei(i32 %{{.*}}) + // CHECK: call signext i8 @_Z16convert_char_rtzi(i32 %{{.*}}) + // CHECK: call signext i8 @_Z16convert_char_rtpi(i32 %{{.*}}) + // CHECK: call signext i8 @_Z16convert_char_rtni(i32 %{{.*}}) + // CHECK: call signext i8 @_Z16convert_char_sati(i32 %{{.*}}) + // CHECK: call signext i8 @_Z20convert_char_sat_rtei(i32 %{{.*}}) + // CHECK: call signext i8 @_Z20convert_char_sat_rtzi(i32 %{{.*}}) + // CHECK: call signext i8 @_Z20convert_char_sat_rtpi(i32 %{{.*}}) + // CHECK: call signext i8 @_Z20convert_char_sat_rtni(i32 %{{.*}}) + + int i1 = 42; + long i2 = 43; + int2 i3 = 44; + + char o1 = convert_char(i1); + // CHECK-DAG: call signext i8 @_Z12convert_chari(i32 %{{.*}}) + // CHECK-DAG: declare spir_func i8 @_Z12convert_chari(i32) #[[ATTR:[0-9]+]] + + char o2 = convert_char(i2); + // CHECK-DAG: call signext i8 @_Z12convert_charl(i64 %{{.*}}) + // CHECK-DAG: declare spir_func i8 @_Z12convert_charl(i64) #[[ATTR]] + + char2 o3 = convert_char2(i3); + // CHECK-DAG: call <2 x i8> @_Z13convert_char2Dv2_i(<2 x i32> %{{.*}}) + // CHECK-DAG: declare spir_func <2 x i8> @_Z13convert_char2Dv2_i(<2 x i32>) #[[ATTR]] +} Index: test/SemaOpenCL/builtin-convert.cl =================================================================== --- /dev/null +++ test/SemaOpenCL/builtin-convert.cl @@ -0,0 +1,38 @@ +// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only + +typedef char char2 __attribute__((ext_vector_type(2))); +typedef int int2 __attribute__((ext_vector_type(2))); +typedef int int3 __attribute__((ext_vector_type(3))); + +#define __OPENCL_CONVERT_DEFAULT__ 0 + +#define convert_char(x) __builtin_opencl_convert((x), (char) 0, __OPENCL_CONVERT_DEFAULT__) +#define convert_char2(x) __builtin_opencl_convert((x), (char2)0, __OPENCL_CONVERT_DEFAULT__) + +__kernel void good() { + int i1 = 42; + long i2 = 43; + int2 i3 = 44; + + char o1 = convert_char(i1); + char o2 = convert_char(i2); + char2 o3 = convert_char2(i3); +} + +__kernel void bad(image1d_t image) { + __builtin_opencl_convert(1); // expected-error{{too few arguments to function call, expected 3, have 1}} + __builtin_opencl_convert(1, 2, 3, 4); // expected-error{{too many arguments to function call, expected 3, have 4}} + + char o1 = convert_char(image); // expected-error{{conversion is available only from OpenCL built-in scalar or vector type}} + + struct st { int i; }; + struct st i2; + char o2 = convert_char(i2); // expected-error{{conversion is available only from OpenCL built-in scalar or vector type}} + + int2 i3 = 42; + char2 o3 = convert_char(i3); // expected-error{{operand and result type must have the same number of elements}} + + int3 i4 = 42; + char2 o4 = convert_char2(i4); // expected-error{{operand and result type must have the same number of elements}} +} +