Page MenuHomePhabricator

Prototype OpenCL BIFs using Tablegen
Needs ReviewPublic

Authored by joey on Oct 9 2018, 7:26 AM.



This is the prototype for the approach that was mentioned by Anastasia in

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

There are 3 things generated from the 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(, "foo", 3) != 0)
    return std::make_pair(707, 2);   // "foo"

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


  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.

Diff Detail

Event Timeline

joey created this revision.Oct 9 2018, 7:26 AM
svenvh added a subscriber: svenvh.Oct 9 2018, 8:21 AM
Nicola added a subscriber: Nicola.Oct 9 2018, 9:20 AM
joey updated this revision to Diff 169182.Oct 11 2018, 3:20 AM

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

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 takes 9.4MB.

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

2131 ↗(On Diff #169182)

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

t/ 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/ error: too many arguments to function call, expected 1, have 2
uint gid = get_global_id(0, 1);
asavonic added inline comments.Oct 26 2018, 8:42 AM
2247 ↗(On Diff #169182)

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

joey updated this revision to Diff 172122.Nov 1 2018, 7:38 AM

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.

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 added a comment.Jan 2 2019, 7:24 AM

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).

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

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

This work is being continued in