[RFC][OpenCL] Add builtin types and functions from the standard headers implicitly in the driver


I would like to check if there is any feedback to the following proposal for
improving the interface of standard type and function includes.


Currently the default clang driver provides incomplete functionality for OpenCL
because the headers with builtin function declarations are not included by
default. The header can only be added using frontend flags requiring ‘-cc1’ or
other frontend forwarding options
(Clang Compiler User’s Manual — Clang 16.0.0git documentation). This means it
is inaccessible to users in a conventional way.

We propose to add the implicit header by default when a source is compiled in
OpenCL mode. A review for this has been uploaded by Matt a few months ago:
https://reviews.llvm.org/D78979. Note that the standard library functionality is
added by default in OpenCL C without using include directives in the compiling
sources. This means all builtin function declarations (there are 17000 of
them) are to be parsed every time the source is compiled because which builtins
are used by the kernels is not known beforehand. This impacts the compilation
speed. For example, parsing a simple kernel with the builtin function
declarations is 138 times slower in a Debug build and 13 times slower in a
Release build than parsing the same kernel without those declarations.

To mitigate the overhead of parsing the full header, a fast Tablegen based
solution has been developed
(https://llvm.org/devmtg/2019-10/talk-abstracts.html#lit5). The parsing speed with
this mechanism for builtin function declarations is only 1.3 times slower in a Debug
build and 1.05 times slower in a Release build compared to clang without the
builtins. While this mechanism covers most of OpenCL standard
functions it lacks two main classes of builtins: builtins defined by
vendor extensions and builtins with enum arguments.


We propose the following changes in the clang driver interface for OpenCL:

  • Enable the fast Tablegen based builtin function declaration mechanism by
    default in the clang driver. This makes the majority of OpenCL builtin
    functions available.
  • In addition, include the small header opencl-c-base.h by default in the clang
    driver. This header provides basic types and constants.
    No frontend or driver flags will be needed to allow using the majority of OpenCL
    types and functions from the standard, at a very low parsing speed increase.

Since the Tablegen mechanism has some small overhead and it is not fully
complete, we propose to add the following additional clang driver flags:

  1. Add a new clang driver flag -cl-no-stdinc (*) that disables such extra includes to
    minimize further compilation speed or for the use cases that don’t require
    standard libraries. The majority of OpenCL clang tests will use this option.
  2. [Optionally, if there is enough interest] Add a new clang driver flag
    -cl-all-stdinc (*) that will include the full header instead of using the Tablegen
    mechanism, at the cost of a significant increase in parsing time.

At present we propose no change to the ‘cc1’ interface, but in the future it is
expected that the functionality will be aligned between driver and frontend
command line interfaces for the OpenCL headers.

(*) The exact spelling of command line options is to be discussed.


This proposal enhances the clang driver with full functionality of OpenCL C by
adding builtin function declarations implicitly without the need for any
frontend flags to be given in the command line.

The default clang behavior proposed is not expected to negatively impact users
of clang as the parsing speed difference remains within the same order of
magnitude. While the fast header mechanism matures, a fallback mechanism will
be provided if needed that would allow switching to the slow header with the
full functionality using a new driver flag. For backward compatibility, another
flag is provided to disable all OpenCL declarations that are not builtin into
the compiler.

The solution proposed improves the driver interface and reduces risks of
forcing the OpenCL community to update their use of clang due to significant
regression of the compilation speed.

The whole tablegen thing seems like a sad path to have gone down, although I can certainly see the practical benefits. Substantially the same problem of compilation-speed exists for <x86intrin.h> (and friends), and I really don’t think we want to start defining all of those with a tablegen rule.

It would be really nice to instead somehow take advantage of the modules infrastructure to fix this problem – I’d really love it if we could somehow start shipping a pre-built module artifact for our giant intrinsics headers, included with the compiler distribution. And then use that by default, regardless of whether users are otherwise enabling modules. If we got that to work, we could use the same solution for both X86 and opencl.

However, independent of that discussion – and more to the immediate issue you’re trying to raise – your proposal seems like it’s resulting in a very complex set of options, and I’m not sure what the purpose is.

IIUC, the overall desire is to have, by default, these tens-of-thousands of prototypes available to all OpenCL compilations. But, I don’t see any reason why users should care exactly HOW these are provided. I’d expect that a given prototype should be provided either by the tablegen-based builtin code, OR an auto-included header file – not both, and not neither.

What difference does it make if the builtin-tablegen code doesn’t provide 100% of the declarations, so long as the remainder are provided by an automatically-included header? Why do you want to make users choose between getting a half-baked set of function prototypes (tablegen version) and slow compilation (auto-including giant header), when you can have the fully correct set of functions AND nearly-as-fast compilation, by simply supplementing tablegen with an auto-included header providing the remainder?

And then you need just a single user-visible option: to disable the automatic declarations (via both tablegen + autoinclude).

Thanks for your input James! I am wondering what your concerns are regarding the Tablegen-based declarations. It would be great if you could elaborate on that.

To give some background: we have looked into modules and PCH some time ago and found Tablegen is the only acceptable solution currently. The problem with modules is that the serialized format (same as PCH) occupies a lot of memory, which is a concern for devices where the compiler is part of the device driver. We have investigated various ways to reduce the size but it was not very easy due to the generality of the serialization format. The techniques we looked at had low chances to be accepted by a wider community and they were not extremely effective either. Another drawback with modules is that a separate cached binary would be needed for each combination of command-line options: a language version + target configurations including extensions. This would mean that a lot of memory would be needed for just declaring the OpenCL functions not only per one module but different build configurations too.

After multiple investigations, we have concluded that the more targeted Tablegen approach was the best solution to satisfy the requirements. Of course we could have a module-based approach for offline compilations and a custom one for online compilations. But then we would have to develop and maintain two mechanisms. We do have two mechanisms now but hopefully, we can migrate to Tablegen one very soon. Some background about the issues is summarized in the original proposal: https://lists.llvm.org/pipermail/cfe-dev/2018-November/060041.html.

Regarding your suggestion of the command line interface: I like the idea of mixing the regular header with the Tablegen declarations. The only concern is that the header with regular declarations could still be long to parse but I believe it would not be significant at the moment. Aside from simplifying the command line interface, we can use it for the future development to mitigate the complexity i.e. some simple and fast-to-parse declarations can just be added in the header to bypass the Tablegen syntax. This should also be simple enough to implement because we are already adding an extra header file declaring the types with the Tablegen declarations that can be reused for missing functionality too or additionally we could add a separate header for those.


My short, not carefully considered, input is that the suggested approach of tablegen + a header seems reasonable.

The modules / PCH approach has been considered for the ever-growing immintrin.h on x86, but I don’t believe it got much traction:

To me the tablegen proposal seems a clever and generalizable means of basically “instantiating” non-templatable builtins and related declarations only as they are needed, avoiding not only parsing as a module/PCH would, but also storage of unused AST nodes.

I have little familiarity with x86intrin.h et al, but why couldn’t the same approach be fruitful there? To the extent their content is not manually written/maintained but rather is generated from much smaller data via some program (is it? At a glance seems possible), couldn’t tablegen be that program, generating only the bare-bones data needed to create all the necessary declarations, and leaving it to the compiler to actually create those declarations as needed, just as Anastasia proposes for OpenCL?

Modules need not be the ne plus ultra solution for large, universally-#included headers, given the speed + storage improvements the OpenCL folks have apparently realized with this alternative approach.

OpenCL Tablegen builtin implementation is fairly generic but right now it only
supports function declarations.

The X86 header seems to contain the function definition instead. However
the definitions seem to follow regular patterns, so it feels like using a
similar Tablegen solution should work, but some infrastructural changes will
be needed first. However, if there are no requirements for the memory size
to store the headers perhaps the modules could be good enough.

Thanks to everyone for the feedback. FYI I have created the following review:
https://reviews.llvm.org/D96515 that enabled the Tablegen based declarations
by default for OpenCL and adds an extra flag to disable those explicitly.