asavonic created this revision.
Herald added subscribers: cfe-commits, kristina, Anastasia, yaxunl.
OpenCL `convert' is a builtin that takes a builtin scalar or vector
type, and converts it to another builtin scalar or vector type using
different rounding modes (round-to-even, round-to-zero, etc.)
Previously `convert' builtin was declared in opencl-c.h header as 6284
function declarations: all combinations of types are multiplied by 10
modes.
This patch adds a single clang builtin that can codegen into any of
these 6284 builtins.
`__builtin_opencl_convert' takes 3 parameters:
1. Value to convert
2. Return type placeholder that is used to determine a result type of
conversion.
3. Integer constant that denotes a rounding mode.
Repository:
rC Clang
https://reviews.llvm.org/D52457
Files:
include/clang/Basic/Builtins.def
include/clang/Basic/DiagnosticSemaKinds.td
lib/CodeGen/CGBuiltin.cpp
lib/Sema/SemaChecking.cpp
test/CodeGenOpenCL/builtin-convert.cl
test/SemaOpenCL/builtin-convert.cl
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}}
+}
+
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: 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<VectorType>()) {
+ SrcElem = SrcVecTy->getElementType();
+ SrcVecSize = SrcVecTy->getNumElements();
+ }
+
+ unsigned DstVecSize = 1;
+ if (const auto *DstVecTy = DstTy->getAs<VectorType>()) {
+ 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: 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<QualType> 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<ParmVarDecl*, 8> 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<QualType> ArgTys,
+ SmallVectorImpl<char> &Result) {
+ const FunctionDecl *FD =
+ createOpenCLBuiltinFunctionDecl(Context, Name, RetTy, ArgTys);
+
+ std::unique_ptr<MangleContext> 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<Attribute::AttrKind> Attrs) {
+
+ const CGFunctionInfo &FuncInfo =
+ CGF.CGM.getTypes().arrangeBuiltinFunctionCall(RetTy, Args);
+
+ llvm::FunctionType *FTy = CGF.CGM.getTypes().GetFunctionType(FuncInfo);
+ llvm::Function *Func =
+ cast<Function>(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<BuiltinType>()) {
+ TypeName = Ty->getName(Context.getPrintingPolicy());
+ } else {
+ const auto *Typedef = DstTy->castAs<TypedefType>();
+ NamedDecl *D = cast<NamedDecl>(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<OpenCLConvertRounding>(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: 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: 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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
