Prototype OpenCL BIFs using Tablegen

[joey]

Summary

This is the prototype for the approach that was mentioned by Anastasia in http://lists.llvm.org/pipermail/cfe-dev/2018-September/059529.html

The tablegen file describes the BIFs and all their overloads, in hopefully a concise manner.

There are 3 things generated from the OpenCLBuiltins.td file.

1. OpenCLArgTypes[], this is a table containing all the different types of overloads. This is a separate table so it can be shared by the BIFs.
2. OpenCLBuiltins[], this is a table that contains all the overloads for the BIFs.
3. isOpenCLBuiltin, this is a function that uses a trie-like switch/case to determine if a StringRef is the name of a BIF.

Just a quick snippet of the above:

OpenCLType OpenCLArgTypes[] = {
// 0
{ OCLT_float, 0, 0, clang::LangAS::Default, },
// 1
{ OCLT_float, 2, 0, clang::LangAS::Default, },

OpenCLBuiltinDecl OpenCLBuiltins[] = {
// acos
  { { OCLT_float, 0, 0, clang::LangAS::Default, }, 1, 0, "", 100,  },
  { { OCLT_float, 2, 0, clang::LangAS::Default, }, 1, 1, "", 100,  },

std::pair<unsigned, unsigned> isOpenCLBuiltin(llvm::StringRef name) {
  switch (name.size()) {
  default: break;
  case 3:  // 1 string to match.
    if (memcmp(name.data()+0, "foo", 3) != 0)
      break;
    return std::make_pair(707, 2);   // "foo"

While it's a prototype, I have tried to keep it as clean as possible.

TODO:

1. Bit-pack the tables to reduce the size.
2. Include the return type in the ArgTypes table to reduce the size.
3. Measure the performance / size impact
4. Auto-generate parts of OCL2Qual, to reduce repeated typing
5. OCL2Qual does not support pointers-to-pointers currently, but I believe no BIFs use that.
6. InsertBuiltinDeclarations builds up an AST function declaration manually, perhaps there is a helper function for this.
7. There is a FIXME in SemaDecl.cpp that needs to be implemented.

[joey]

Re-uploading the patch. The previous version of the patch was missing the include and lib prefixes.

[asavonic]

TODO
Measure the performance / size impact

I looks like the total size of OpenCLBuiltinDecl table should be
~450KB (for ~15000 functions). I guess it can be reduced furthermore by:

1. Replacing return type with an index (you've mentioned this in TODO).
2. Replace extension and a version with a single index.
3. Same for NumArgs and ArgTableIndex - they seem to be really tied together, so maybe we can use a single identifier for them?

With these improvements the total size should become ~175KB, which
seems to be ok, considering that libclangSema.so takes 9.4MB.

In my opinion, even 450KB is a vast improvement over what we have now,
because opencl-c.h alone takes 786KB.

lib/Sema/SemaExpr.cpp

2131 ↗

(On Diff #169182)

This dummy identifier name actually makes a diagnostic a bit weird:

t/gid.cl:3:31: error: too many arguments to function call, expected single argument 'a0', have 2 arguments
uint gid = get_global_id(0, 2);

It seems to be better if you pass nullptr instead:

t/gid.cl:4:31: error: too many arguments to function call, expected 1, have 2
uint gid = get_global_id(0, 1);

[asavonic]

lib/Sema/SemaExpr.cpp

2247 ↗

(On Diff #169182)

It will be good to have a driver option to enable this feature.

[joey]

I re-worked where the builtins are actually inserted, now it's in a similar place to the "normal" clang builtins.

I addressed the issue of putting the return type also into the args type (could rename this to signature/proto table to make it more obvious).

I'm planning to write an RFC to get buy-in from the rest of the community before doing too much else to the patch.

[Nicola]

This looks quite useful, thanks for taking time to do it! I can help verifying some of the tablegen'd builtins with internal tests too, if you decide to pursue this further.

[Nicola]

Since this review is just for a prototype, I'll leave some feedback after having played with it for a bit.
It seems to be quite trivial to add new builtins and fix existing ones (conversion builtins sat/rounding modes). One builtin that might have a small issue is the as_type builtin in which source and destination type need to have the same bitwidth. To get that to work I ended up patching TableGen to add !mul support for it (and set VectorSize to 1 for scalars).

[AlexeySotkin]

LGTM! @joey, any idea when it will be landed?

[svenvh]

In D53023#1499942, @AlexeySotkin wrote:

LGTM! @joey, any idea when it will be landed?

This work is being continued in https://reviews.llvm.org/D60763

include/clang/Basic/CMakeLists.txt

clang_tablegen(AttrSubMatchRulesList.inc -gen-clang-attr-subject-match-rule-list
   TARGET ClangAttrSubjectMatchRuleList)
 
 clang_tablegen(AttrHasAttributeImpl.inc -gen-clang-attr-has-attribute-impl
   -I ${CMAKE_CURRENT_SOURCE_DIR}/../../
   SOURCE Attr.td
   TARGET ClangAttrHasAttributeImpl
   )
 
+clang_tablegen(OpenCLBuiltins.inc
+  -I ${CMAKE_CURRENT_SOURCE_DIR}/../../ -gen-clang-opencl-builtins
+  SOURCE OpenCLBuiltins.td
+  TARGET ClangOpenCLBuiltinsImpl
+  )
+
 # ARM NEON
 clang_tablegen(arm_neon.inc -gen-arm-neon-sema
   SOURCE arm_neon.td
   TARGET ClangARMNeon)
 clang_tablegen(arm_fp16.inc -gen-arm-neon-sema
   SOURCE arm_fp16.td
   TARGET ClangARMFP16)

include/clang/Basic/OpenCLBuiltins.td

+//==--- OpenCLBuiltins.td - OpenCL builtin definitions --------------------===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+
+class Type<string Name> {
+  string name = Name;
+  int vecWidth = 0;
+  int isPointer = 0;
+  string as = "clang::LangAS::Default";
+  int has_rounding = 0;
+  int is_integer = 0;
+}
+
+class VectorType<Type Ty, int VecWidth> : Type<Ty.name> {
+  int vecWidth = VecWidth;
+}
+
+class AddressSpace<string _as> {
+  string as = _as;
+}
+
+def Global: AddressSpace<"clang::LangAS::opencl_global">;
+def Constant: AddressSpace<"clang::LangAS::opencl_constant">;
+def Generic: AddressSpace<"clang::LangAS::opencl_generic">;
+
+class PointerType<Type Ty, AddressSpace AS = Global> : Type<Ty.name> {
+  int isPointer = 1;
+  string as = AS.as;
+}
+
+class Version<int v> {
+  int version = v;
+}
+
+def CL10: Version<100>;
+def CL11: Version<110>;
+def CL12: Version<120>;
+def CL20: Version<200>;
+
+def void_t: Type<"void_t">;
+
+def size_t_t: Type<"size_t">;
+
+let is_integer = 1 in {
+  def char_t   : Type<"char">;
+  def uchar_t  : Type<"uchar">;
+  def short_t  : Type<"short">;
+  def ushort_t : Type<"ushort">;
+  def int_t    : Type<"int">;
+  def uint_t   : Type<"uint">;
+  def long_t   : Type<"long">;
+  def ulong_t  : Type<"ulong">;
+}
+def int2 : VectorType<int_t, 2>;
+
+def half_t   : Type<"half">;
+let has_rounding = 1 in {
+  def float_t  : Type<"float">;
+  def double_t : Type<"double">;
+}
+
+def float4_t  : VectorType<float_t, 4>;
+
+def image2d_ro_t : Type<"image2d_ro_t">;
+def image2d_wo_t : Type<"image2d_wo_t">;
+
+class Builtin<string Name, Type ReturnType, list<Type> Args> {
+  string name = Name;
+  string extension;
+  list<Type> args = !listconcat([ReturnType], Args);
+  string extension = "";
+  Version version = CL10;
+}
+
+// Creates builtins for one argument BIFs, taking and returning the same type.
+multiclass bi_vec<string name, Type ReturnType> {
+  def: Builtin<name, ReturnType, [ReturnType]>;
+  foreach v = [2, 3, 4, 8, 16] in {
+    def : Builtin<name, VectorType<ReturnType, v>, [VectorType<ReturnType, v>]>;
+  }
+}
+
+// Creates builtins for two argument BIFs, taking and returning the same type.
+multiclass bi_vec2<string name, Type ReturnType> {
+  def: Builtin<name, ReturnType, [ReturnType]>;
+  foreach v = [2, 3, 4, 8, 16] in {
+    def : Builtin<name, VectorType<ReturnType, v>, [VectorType<ReturnType, v>, VectorType<ReturnType, v>]>;
+  }
+}
+
+// Creates builtins for two argument BIFs, taking a vector and a scale and returning the vector type.
+multiclass bi_vec3<string name, Type ReturnType> {
+  foreach v = [2, 3, 4, 8, 16] in {
+    def : Builtin<name, VectorType<ReturnType, v>, [VectorType<ReturnType, v>, ReturnType]>;
+  }
+}
+
+// 6.12.2
+foreach name = ["acos", "acosh", "acospi", "asin", "asinh", "asinpi", "atan"] in {
+  foreach type = [float_t, double_t] in {
+    defm name#type : bi_vec<name, type>;
+  }
+}
+
+foreach name = ["atan2"] in {
+  foreach type = [float_t, double_t] in {
+    defm name#type : bi_vec2<name, type>;
+  }
+}
+
+foreach name = ["fmax", "fmin"] in {
+  foreach type = [float_t, double_t] in {
+    defm: bi_vec2<name, type>;
+    defm: bi_vec3<name, type>;
+  }
+}
+
+// example 'foo', to show using 'version'
+def: Builtin<"foo_version", int_t, [PointerType<int_t, Global>]>;
+let version = CL20 in {
+  def: Builtin<"foo_version", int_t, [PointerType<int_t, Constant>]>;
+}
+
+// Helper classes for the convert_ BIFs.
+class SatModes<Type ty> {
+  list<string> modes = !if(ty.is_integer, ["", "_sat"], [""]);
+}
+
+class RoundingModes<Type ty, Type ty2> {
+  list<string> modes = !if(!or(ty.has_rounding, ty2.has_rounding), ["", "_rte", "_rtz", "_rtp", "_rtn"], [""]);
+}
+
+// Generate the convert_ builtins.
+foreach type = [float_t, double_t, char_t, uchar_t, short_t, ushort_t,
+                int_t, uint_t, long_t, ulong_t] in {
+  foreach type2 = [float_t, double_t, char_t, uchar_t, short_t, ushort_t,
+                   int_t, uint_t, long_t, ulong_t] in {
+    foreach sat = SatModes<type>.modes in {
+      foreach rte = RoundingModes<type, type2>.modes in {
+        def: Builtin<"convert_"  # type.name # sat # rte, type, [type2]>;
+        foreach v = [2, 3, 4, 8, 16] in {
+          def: Builtin<"convert_" # type.name # v # sat # rte, VectorType<type, v>, [VectorType<type2, v>]>;
+        }
+      }
+    }
+  }
+}
+
+// Example showing 'extension'
+let extension = "cl_khr_subgroups" in {
+  def : Builtin<"get_sub_group_size", uint_t, []>;
+}
+
+// samplerless read image
+def : Builtin<"read_imagef", float4_t, [image2d_ro_t, VectorType<int_t, 2>]>;
+def : Builtin<"write_imagef", void_t, [image2d_wo_t, VectorType<int_t, 2>, VectorType<float_t, 4>]>;
+
+// 6.11.1
+def get_work_dim : Builtin<"get_work_dim", uint_t, []>;
+foreach name = ["get_global_size", "get_global_id", "get_local_size",
+                "get_local_id", "get_num_groups", "get_group_id",
+                "get_global_offset"] in {
+  def name : Builtin<name, size_t_t, [uint_t]>;
+}

lib/Sema/SemaLookup.cpp

 #include "llvm/Support/ErrorHandling.h"
 #include <algorithm>
 #include <iterator>
 #include <list>
 #include <set>
 #include <utility>
 #include <vector>
 
+#include "clang/Basic/OpenCLBuiltins.inc"
+
 using namespace clang;
 using namespace sema;
 
 namespace {
   class UnqualUsingEntry {
     const DeclContext *Nominated;
     const DeclContext *CommonAncestor;
 

 
 LLVM_DUMP_METHOD void LookupResult::dump() {
   llvm::errs() << "lookup results for " << getLookupName().getAsString()
                << ":\n";
   for (NamedDecl *D : *this)
     D->dump();
 }
 
+// TODO: Auto-generate this from tablegen
+static QualType OCL2Qual(ASTContext &Context, OpenCLType Ty) {
+  QualType RT = Context.VoidTy;
+
+  switch (Ty.ID) {
+  case OCLT_size_t:
+    RT = Context.getSizeType();
+    break;
+  case OCLT_int:
+    RT = Context.IntTy;
+    break;
+  case OCLT_uint:
+    RT = Context.UnsignedIntTy;
+    break;
+  case OCLT_char:
+    RT = Context.CharTy;
+    break;
+  case OCLT_uchar:
+    RT = Context.UnsignedCharTy;
+    break;
+  case OCLT_short:
+    RT = Context.ShortTy;
+    break;
+  case OCLT_ushort:
+    RT = Context.UnsignedShortTy;
+    break;
+  case OCLT_long:
+    RT = Context.LongTy;
+    break;
+  case OCLT_ulong:
+    RT = Context.UnsignedLongTy;
+    break;
+  case OCLT_float:
+    RT = Context.FloatTy;
+    break;
+  case OCLT_double:
+    RT = Context.DoubleTy;
+    break;
+  case OCLT_image2d_ro_t:
+    RT = Context.OCLImage2dROTy;
+    break;
+  case OCLT_image2d_wo_t:
+    RT = Context.OCLImage2dWOTy;
+    break;
+  case OCLT_void_t:
+    RT = Context.VoidTy;
+    break;
+  default:
+    assert(0 && "unexpected type!");
+    break;
+  }
+
+  if (Ty.VectorWidth > 0)
+    RT = Context.getExtVectorType(RT, Ty.VectorWidth);
+
+  if (Ty.isPointer != 0) {
+    RT = Context.getAddrSpaceQualType(RT, Ty.AS);
+    RT = Context.getPointerType(RT);
+  }
+  return RT;
+}
+
+static void InsertBuiltinDeclarations(Sema &S, LookupResult &LR,
+                                      IdentifierInfo *II,
+                                      unsigned Index, unsigned Len) {
+  for (unsigned i = 0; i < Len; ++i) {
+    OpenCLBuiltinDecl &Decl = OpenCLBuiltins[Index - 1 + i];
+    ASTContext &Context = S.Context;
+
+    // Ignore this BIF if the the version is incorrect.
+    if (Context.getLangOpts().OpenCLVersion < Decl.Version)
+      continue;
+
+    FunctionProtoType::ExtProtoInfo PI;
+    PI.Variadic = false;
+
+    QualType RT = OCL2Qual(Context, OpenCLArgTypes[Decl.ArgTableIndex]);
+
+    SmallVector<QualType, 5> ArgTypes;
+    for (unsigned i = 1; i < Decl.NumArgs; i++) {
+      QualType Ty = OCL2Qual(Context, OpenCLArgTypes[Decl.ArgTableIndex + i]);
+      ArgTypes.push_back(Ty);
+    }
+
+    QualType R = Context.getFunctionType(RT, ArgTypes, PI);
+    SourceLocation Loc = LR.getNameLoc();
+
+    // TODO: This part is taken from Sema::LazilyCreateBuiltin, maybe refactor it.
+    DeclContext *Parent = Context.getTranslationUnitDecl();
+    FunctionDecl *New = FunctionDecl::Create(Context,
+                                             Parent,
+                                             Loc, Loc, II, R, /*TInfo=*/nullptr,
+                                             SC_Extern,
+                                             false,
+                                             R->isFunctionProtoType());
+    New->setImplicit();
+
+    // Create Decl objects for each parameter, adding them to the
+    // FunctionDecl.
+    if (const FunctionProtoType *FT = dyn_cast<FunctionProtoType>(R)) {
+      SmallVector<ParmVarDecl*, 16> Params;
+      for (unsigned i = 0, e = FT->getNumParams(); i != e; ++i) {
+        ParmVarDecl *parm =
+            ParmVarDecl::Create(Context, New, SourceLocation(), SourceLocation(),
+                                nullptr, FT->getParamType(i), /*TInfo=*/nullptr,
+                                SC_None, nullptr);
+        parm->setScopeInfo(0, i);
+        Params.push_back(parm);
+      }
+      New->setParams(Params);
+    }
+
+    New->addAttr(OverloadableAttr::CreateImplicit(Context));
+
+    if (strlen(Decl.Extension))
+      S.setOpenCLExtensionForDecl(New, Decl.Extension);
+
+    LR.addDecl(New);
+  }
+
+  // If we added overloads, need to resolve the lookup result.
+  if (Len > 1)
+    LR.resolveKind();
+}
+
 /// Lookup a builtin function, when name lookup would otherwise
 /// fail.
 static bool LookupBuiltin(Sema &S, LookupResult &R) {
   Sema::LookupNameKind NameKind = R.getLookupKind();
 
   // If we didn't find a use of this identifier, and if the identifier
   // corresponds to a compiler builtin, create the decl object for the builtin
   // now, injecting it into translation unit scope, and return it.
   if (NameKind == Sema::LookupOrdinaryName ||
       NameKind == Sema::LookupRedeclarationWithLinkage) {
     IdentifierInfo *II = R.getLookupName().getAsIdentifierInfo();
     if (II) {
       if (S.getLangOpts().CPlusPlus && NameKind == Sema::LookupOrdinaryName) {
         if (II == S.getASTContext().getMakeIntegerSeqName()) {
           R.addDecl(S.getASTContext().getMakeIntegerSeqDecl());
           return true;
         } else if (II == S.getASTContext().getTypePackElementName()) {
           R.addDecl(S.getASTContext().getTypePackElementDecl());
           return true;
         }
       }
 
+      // Check if this is an OpenCL Builtin, and if so, insert the declarations.
+      if (S.getLangOpts().OpenCL) {
+        auto Index = isOpenCLBuiltin(II->getName());
+        if (Index.first) {
+          InsertBuiltinDeclarations(S, R, II, Index.first, Index.second);
+          return true;
+        }
+      }
+
       // If this is a builtin on this (or all) targets, create the decl.
       if (unsigned BuiltinID = II->getBuiltinID()) {
         // In C++ and OpenCL (spec v1.2 s6.9.f), we don't have any predefined
         // library functions like 'malloc'. Instead, we'll just error.
         if ((S.getLangOpts().CPlusPlus || S.getLangOpts().OpenCL) &&
             S.Context.BuiltinInfo.isPredefinedLibFunction(BuiltinID))
           return false;
 

test/SemaOpenCL/builtin-new.cl

+// RUN: %clang_cc1 %s -triple spir -verify -pedantic -fsyntax-only -cl-std=CL1.2 -DCL12
+// RUN: %clang_cc1 %s -triple spir -verify -pedantic -fsyntax-only -cl-std=CL2.0 -DCL20
+
+typedef float float4 __attribute__((ext_vector_type(4)));
+typedef int int4 __attribute__((ext_vector_type(4)));
+typedef int int2 __attribute__((ext_vector_type(2)));
+typedef unsigned int uint;
+
+kernel void test(global float4* buf, global int4* res)
+{
+  res[0] = convert_int4(buf[0]);
+}
+
+kernel void test2(global int* bar) {
+  bar[0] = foo_version(bar);
+}
+
+kernel void test3(constant int* bar, global int* baz) {
+  baz[0] = foo_version(bar);
+#ifdef CL12
+// expected-error@-2{{passing '__constant int *' to parameter of type '__global int *' changes address space of pointer}}
+#endif
+}
+
+kernel void test4(image2d_t img, int2 coord, global float4* out) {
+  out[0] = read_imagef(img, coord);
+}
+
+kernel void test5(write_only image2d_t img, int2 coord, float4 colour) {
+  write_imagef(img, coord, colour);
+}
+
+#ifdef CL20
+kernel void test6(global uint* out) {
+  out[0] = get_sub_group_size();
+// expected-error@-1{{use of declaration 'get_sub_group_size' requires cl_khr_subgroups extension to be enabled}}
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+  out[1] = get_sub_group_size();
+}
+#endif

utils/TableGen/CMakeLists.txt

 set(LLVM_LINK_COMPONENTS Support)
 
 add_tablegen(clang-tblgen CLANG
   ClangASTNodesEmitter.cpp
   ClangAttrEmitter.cpp
   ClangCommentCommandInfoEmitter.cpp
   ClangCommentHTMLNamedCharacterReferenceEmitter.cpp
   ClangCommentHTMLTagsEmitter.cpp
   ClangDataCollectorsEmitter.cpp
   ClangDiagnosticsEmitter.cpp
+  ClangOpenCLBuiltinEmitter.cpp
   ClangOptionDocEmitter.cpp
   ClangSACheckersEmitter.cpp
   NeonEmitter.cpp
   TableGen.cpp
   )
 set_target_properties(clang-tblgen PROPERTIES FOLDER "Clang tablegenning")

utils/TableGen/ClangOpenCLBuiltinEmitter.cpp

+//===- ClangOpenCLBuiltinEmitter.cpp - Generate Clang OpenCL Builtin handling
+//=-*- C++ -*--=//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This tablegen backend emits Clang OpenCL Builtin checking code.
+//
+//===----------------------------------------------------------------------===//
+
+#include "llvm/ADT/MapVector.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/SmallString.h"
+#include "llvm/ADT/StringExtras.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/ADT/StringSet.h"
+#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/raw_ostream.h"
+#include "llvm/TableGen/Error.h"
+#include "llvm/TableGen/Record.h"
+#include "llvm/TableGen/StringMatcher.h"
+#include "llvm/TableGen/TableGenBackend.h"
+
+#include <set>
+
+using namespace llvm;
+
+namespace {
+class BuiltinNameEmitter {
+public:
+  BuiltinNameEmitter(RecordKeeper &Records, raw_ostream &OS)
+      : Records(Records), OS(OS) {}
+
+  void Emit();
+
+private:
+  RecordKeeper &Records;
+  raw_ostream &OS;
+
+  void EmitDeclarations();
+  void EmitTable();
+  void GetOverloads();
+
+  MapVector<StringRef, std::vector<std::pair<const Record *, unsigned>>>
+      OverloadInfo;
+  std::vector<std::pair<std::vector<Record *>, unsigned>> ArgTypesSet;
+};
+} // namespace
+
+void BuiltinNameEmitter::GetOverloads() {
+  unsigned CumulativeArgIndex = 0;
+  std::vector<Record *> Builtins = Records.getAllDerivedDefinitions("Builtin");
+  for (const auto *B : Builtins) {
+    StringRef BName = B->getValueAsString("name");
+
+    if (OverloadInfo.find(BName) == OverloadInfo.end()) {
+      OverloadInfo.insert(std::make_pair(
+          BName, std::vector<std::pair<const Record *, unsigned>>{}));
+    }
+
+    auto Args = B->getValueAsListOfDefs("args");
+    auto it =
+        std::find_if(ArgTypesSet.begin(), ArgTypesSet.end(),
+                     [&](const std::pair<std::vector<Record *>, unsigned> &a) {
+                       return a.first == Args;
+                     });
+    unsigned ArgIndex;
+    if (it == ArgTypesSet.end()) {
+      ArgTypesSet.push_back(std::make_pair(Args, CumulativeArgIndex));
+      ArgIndex = CumulativeArgIndex;
+      CumulativeArgIndex += Args.size();
+    } else {
+      ArgIndex = it->second;
+    }
+    OverloadInfo[BName].push_back(std::make_pair(B, ArgIndex));
+  }
+}
+
+void BuiltinNameEmitter::EmitDeclarations() {
+  OS << "enum OpenCLTypeID {\n";
+  std::vector<Record *> Types = Records.getAllDerivedDefinitions("Type");
+  StringMap<bool> TypesSeen;
+  for (const auto *T : Types) {
+    if (TypesSeen.find(T->getValueAsString("name")) == TypesSeen.end())
+      OS << "  OCLT_" + T->getValueAsString("name") << ",\n";
+    TypesSeen.insert(std::make_pair(T->getValueAsString("name"), true));
+  }
+  OS << "};\n";
+
+  OS << R"(
+struct OpenCLType {
+  OpenCLTypeID ID;
+  unsigned VectorWidth;
+  unsigned isPointer;
+  clang::LangAS AS;
+};
+
+struct OpenCLBuiltinDecl {
+  unsigned NumArgs;
+  unsigned ArgTableIndex;
+  const char* Extension;
+  unsigned Version;
+};
+
+)";
+}
+
+void BuiltinNameEmitter::EmitTable() {
+  OS << "OpenCLBuiltinDecl OpenCLBuiltins[] = {\n";
+  for (auto &i : OverloadInfo) {
+    StringRef Name = i.first;
+    OS << "// " << Name << "\n";
+    for (auto &Overload : i.second) {
+      OS << "  { "
+         << Overload.first->getValueAsListOfDefs("args").size() << ", "
+         << Overload.second << ", " << '"'
+         << Overload.first->getValueAsString("extension") << "\", "
+         << Overload.first->getValueAsDef("version")->getValueAsInt("version")
+         << " },\n";
+    }
+  }
+  OS << "};\n\n";
+}
+
+void BuiltinNameEmitter::Emit() {
+  emitSourceFileHeader("OpenCL Builtin handling", OS);
+
+  OS << "#include \"llvm/ADT/StringRef.h\"\n\n";
+
+  EmitDeclarations();
+
+  GetOverloads();
+
+  std::vector<std::vector<Record *>> ArgTypes;
+
+  OS << "OpenCLType OpenCLArgTypes[] = {\n";
+  for (auto &P : ArgTypesSet) {
+    OS << "// " << P.second << "\n";
+    for (Record *R : P.first) {
+      OS << "{ OCLT_" << R->getValueAsString("name") << ", "
+         << R->getValueAsInt("vecWidth") << ", "
+         << R->getValueAsInt("isPointer") << ", " << R->getValueAsString("as")
+         << ", "
+         << "}, ";
+      OS << "\n";
+    }
+  }
+  OS << "};\n\n";
+
+  EmitTable();
+
+  // Construct a StringMatcher.
+  std::vector<StringMatcher::StringPair> validBuiltins;
+  unsigned CumulativeIndex = 1;
+  for (auto &i : OverloadInfo) {
+    auto &Ov = i.second;
+    std::string RetStmt;
+    raw_string_ostream SS(RetStmt);
+    SS << "return std::make_pair(" << CumulativeIndex << ", " << Ov.size()
+       << ");";
+    SS.flush();
+    CumulativeIndex += Ov.size();
+
+    validBuiltins.push_back(StringMatcher::StringPair(i.first, RetStmt));
+  }
+
+  OS << "// Return 0 if name is not a recognized OpenCL builtin, or an index\n"
+        "// into a table of declarations if it is an OpenCL builtin.\n"
+        "std::pair<unsigned, unsigned> isOpenCLBuiltin(llvm::StringRef name) "
+        "{\n";
+
+  StringMatcher("name", validBuiltins, OS).Emit(0, true);
+
+  OS << "  return std::make_pair(0, 0);\n";
+  OS << "}\n";
+}
+
+namespace clang {
+
+void EmitClangOpenCLBuiltins(RecordKeeper &Records, raw_ostream &OS) {
+  BuiltinNameEmitter NameChecker(Records, OS);
+  NameChecker.Emit();
+}
+
+} // end namespace clang

utils/TableGen/TableGen.cpp

enum ActionType {
   GenArmNeon,
   GenArmFP16,
   GenArmNeonSema,
   GenArmNeonTest,
   GenAttrDocs,
   GenDiagDocs,
   GenOptDocs,
   GenDataCollectors,
-  GenTestPragmaAttributeSupportedAttributes
+  GenTestPragmaAttributeSupportedAttributes,
+  GenClangOpenCLBuiltins,
 };
 
 namespace {
 cl::opt<ActionType> Action(
     cl::desc("Action to perform:"),
     cl::values(
         clEnumValN(PrintRecords, "print-records",
                    "Print all records to stdout (default)"),
cl::values(
         clEnumValN(GenDiagDocs, "gen-diag-docs",
                    "Generate diagnostic documentation"),
         clEnumValN(GenOptDocs, "gen-opt-docs", "Generate option documentation"),
         clEnumValN(GenDataCollectors, "gen-clang-data-collectors",
                    "Generate data collectors for AST nodes"),
         clEnumValN(GenTestPragmaAttributeSupportedAttributes,
                    "gen-clang-test-pragma-attribute-supported-attributes",
                    "Generate a list of attributes supported by #pragma clang "
-                   "attribute for testing purposes")));
+                   "attribute for testing purposes"),
+        clEnumValN(GenClangOpenCLBuiltins, "gen-clang-opencl-builtins",
+                   "Generate OpenCL builtin handlers")));
 
 cl::opt<std::string>
 ClangComponent("clang-component",
                cl::desc("Only use warnings from specified component"),
                cl::value_desc("component"), cl::Hidden);
 
 bool ClangTableGenMain(raw_ostream &OS, RecordKeeper &Records) {
   switch (Action) {
case GenOptDocs:
     EmitClangOptDocs(Records, OS);
     break;
   case GenDataCollectors:
     EmitClangDataCollectors(Records, OS);
     break;
   case GenTestPragmaAttributeSupportedAttributes:
     EmitTestPragmaAttributeSupportedAttributes(Records, OS);
     break;
+  case GenClangOpenCLBuiltins:
+    EmitClangOpenCLBuiltins(Records, OS);
+    break;
   }
 
   return false;
 }
 }
 
 int main(int argc, char **argv) {
   sys::PrintStackTraceOnErrorSignal(argv[0]);

utils/TableGen/TableGenBackends.h

 void EmitClangDiagDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
 void EmitClangOptDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
 
 void EmitClangDataCollectors(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
 
 void EmitTestPragmaAttributeSupportedAttributes(llvm::RecordKeeper &Records,
                                                 llvm::raw_ostream &OS);
 
+void EmitClangOpenCLBuiltins(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
 } // end namespace clang
 
 #endif