[OpenCL] Implement OpenCL convert builtin

[asavonic]

Summary

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.

include/clang/Basic/Builtins.def

 BUILTIN(__xray_customevent, "vcC*z", "")
 BUILTIN(__xray_typedevent, "vzcC*z", "")
 
 // 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)
+
 #undef BUILTIN
 #undef LIBBUILTIN
 #undef LANGBUILTIN

include/clang/Basic/DiagnosticSemaKinds.td

 def err_opencl_enqueue_kernel_blocks_non_local_void_args : Error<
   "blocks used in enqueue_kernel call are expected to have parameters of type 'local void*'">;
 def err_opencl_enqueue_kernel_blocks_no_args : Error<
   "blocks with parameters are not accepted in this prototype of enqueue_kernel call">;
 
 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">,
   InGroup<OpenCLUnsupportedRGBA>;
 } // end of sema category
 
 let CategoryName = "OpenMP Issue" in {
 // OpenMP support.

lib/CodeGen/CGBuiltin.cpp

RValue CodeGenFunction::emitRotate(const CallExpr *E, bool IsRotateRight) {
   ShiftAmt = Builder.CreateIntCast(ShiftAmt, Ty, false);
 
   // Rotate is a special case of LLVM funnel shift - 1st 2 args are the same.
   unsigned IID = IsRotateRight ? Intrinsic::fshr : Intrinsic::fshl;
   Value *F = CGM.getIntrinsic(IID, Ty);
   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) {
   // 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: {
     DestAddr = Address(Builder.CreateBitCast(DestAddr.getPointer(), BPP, "cp"),
                        DestAddr.getAlignment());
     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);
   }
 
   // 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

default:
     llvm_unreachable("Invalid builtin function");
   }
   Call->setType(S.Context.getPointerType(S.Context.getQualifiedType(
       RT.getUnqualifiedType(), Qual)));
 
   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
 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::BIget_kernel_sub_group_count_for_ndrange:
     if (SemaOpenCLBuiltinNDRangeAndBlock(*this, TheCall))
       return ExprError();
     break;
   case Builtin::BI__builtin_os_log_format:
   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;
+  }
   }
 
   // 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-convert.cl

+// 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]]
+}

test/SemaOpenCL/builtin-convert.cl

+// 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}}
+}
+