[OpenCL] Add builtin declarations by default.

[Anastasia]

Summary

See original RFC: https://lists.llvm.org/pipermail/cfe-dev/2021-February/067610.html

This patch enabled the Tablegen based declarations by default along with opencl-c-base.h containing complete list of OpenCL types. It also adds a new flag -cl-no-stdinc that disables extra header includes and declarations.

The next step would be to either add missing declarations into Tablegen header or opencl-c-base.h.

The commit message suggestion:

This patch enables the builtin function declarations by default using the Tablegen solution along with the implicit include of opencl-c-base.h header. A new flag '-cl-no-stdinc' disabling all default declarations and header includes is added. If any other mechanisms were used to include the declarations (e.g. with -Xclang -finclude-default-header) and the new default approach is not sufficient the `-cl-no-stdinc` flag has to be used to switch to the old behavior.

[svenvh]

It probably makes sense to update clang/docs/UsersManual.rst as part of this change. In particular the following sentence is no longer true after this patch: "By default the OpenCL headers are not loaded and therefore certain builtin types and most of builtin functions are not declared."

clang/include/clang/Driver/Options.td

822

Typo.

Also, this suggests the option is limited to OpenCL C, is that your intent?

[Anastasia]

Changed command line option text.

[Anastasia]

In D96515#2559424, @svenvh wrote:

It probably makes sense to update clang/docs/UsersManual.rst as part of this change. In particular the following sentence is no longer true after this patch: "By default the OpenCL headers are not loaded and therefore certain builtin types and most of builtin functions are not declared."

Yes, that's right but I think there is a bigger change that needs to be made i.e. I would completely remove -finclude-default-header and let it live on OpenCLSupport page. I would prefer a separate review for docs though.

[Anastasia]

In D96515#2560280, @Anastasia wrote:

In D96515#2559424, @svenvh wrote:

It probably makes sense to update clang/docs/UsersManual.rst as part of this change. In particular the following sentence is no longer true after this patch: "By default the OpenCL headers are not loaded and therefore certain builtin types and most of builtin functions are not declared."

Yes, that's right but I think there is a bigger change that needs to be made i.e. I would completely remove -finclude-default-header and let it live on OpenCLSupport page. I would prefer a separate review for docs though.

FYI here is the review for the docs update: https://reviews.llvm.org/D96616

[svenvh]

LGTM

[thakis]

Looks like this breaks tests: http://45.33.8.238/linux/39988/step_8.txt

Please take a look, and revert for now if it takes a while to fix.

[Anastasia]

In D96515#2578629, @thakis wrote:

Looks like this breaks tests: http://45.33.8.238/linux/39988/step_8.txt

Please take a look, and revert for now if it takes a while to fix.

Thanks for pointing this out. I think the following change should fix the issue, but I am still trying to reproduce it locally as it is taking time to build.

diff --git a/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp b/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
index 8ba26c3b7396..31f61027994f 100644
--- a/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
+++ b/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
@@ -1,4 +1,4 @@
-// RUN: pp-trace -callbacks '*,-FileChanged,-MacroDefined' %s -- -x cl | FileCheck --strict-whitespace %s
+// RUN: pp-trace -callbacks '*,-FileChanged,-MacroDefined' %s -- -x cl -cl-no-stdinc | FileCheck --strict-whitespace %s
 
 #pragma OPENCL EXTENSION all : disable
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : disable

FYI, @thakis would you be able to approve the fix? Otherwise, I am also ok to revert it until I figure out how to commit the fix.

[Anastasia]

In D96515#2578666, @Anastasia wrote:

In D96515#2578629, @thakis wrote:

Looks like this breaks tests: http://45.33.8.238/linux/39988/step_8.txt

Please take a look, and revert for now if it takes a while to fix.

Thanks for pointing this out. I think the following change should fix the issue, but I am still trying to reproduce it locally as it is taking time to build.

diff --git a/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp b/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
index 8ba26c3b7396..31f61027994f 100644
--- a/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
+++ b/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
@@ -1,4 +1,4 @@
-// RUN: pp-trace -callbacks '*,-FileChanged,-MacroDefined' %s -- -x cl | FileCheck --strict-whitespace %s
+// RUN: pp-trace -callbacks '*,-FileChanged,-MacroDefined' %s -- -x cl -cl-no-stdinc | FileCheck --strict-whitespace %s
 
 #pragma OPENCL EXTENSION all : disable
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : disable

FYI, @thakis would you be able to approve the fix? Otherwise, I am also ok to revert it until I figure out how to commit the fix.

@klimek, I am not very familiar with clang-tools-extra, would it be possible to look at the fix for the broken test? I can create a review if needed.

[thakis]

In D96515#2578666, @Anastasia wrote:

In D96515#2578629, @thakis wrote:

Looks like this breaks tests: http://45.33.8.238/linux/39988/step_8.txt

Please take a look, and revert for now if it takes a while to fix.

Thanks for pointing this out. I think the following change should fix the issue, but I am still trying to reproduce it locally as it is taking time to build.

diff --git a/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp b/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
index 8ba26c3b7396..31f61027994f 100644
--- a/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
+++ b/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
@@ -1,4 +1,4 @@
-// RUN: pp-trace -callbacks '*,-FileChanged,-MacroDefined' %s -- -x cl | FileCheck --strict-whitespace %s
+// RUN: pp-trace -callbacks '*,-FileChanged,-MacroDefined' %s -- -x cl -cl-no-stdinc | FileCheck --strict-whitespace %s
 
 #pragma OPENCL EXTENSION all : disable
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : disable

FYI, @thakis would you be able to approve the fix? Otherwise, I am also ok to revert it until I figure out how to commit the fix.

The test is already broken, so speculatively landing that and seeing if it helps sounds fine. But let's try to not keep the build broken for hours. (Was afk for a bit, and it looks like it's still broken now that I look again.)

[Anastasia]

In D96515#2578765, @thakis wrote:

In D96515#2578666, @Anastasia wrote:

In D96515#2578629, @thakis wrote:

Looks like this breaks tests: http://45.33.8.238/linux/39988/step_8.txt

Please take a look, and revert for now if it takes a while to fix.

Thanks for pointing this out. I think the following change should fix the issue, but I am still trying to reproduce it locally as it is taking time to build.

diff --git a/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp b/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
index 8ba26c3b7396..31f61027994f 100644
--- a/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
+++ b/clang-tools-extra/test/pp-trace/pp-trace-pragma-opencl.cpp
@@ -1,4 +1,4 @@
-// RUN: pp-trace -callbacks '*,-FileChanged,-MacroDefined' %s -- -x cl | FileCheck --strict-whitespace %s
+// RUN: pp-trace -callbacks '*,-FileChanged,-MacroDefined' %s -- -x cl -cl-no-stdinc | FileCheck --strict-whitespace %s
 
 #pragma OPENCL EXTENSION all : disable
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : disable

FYI, @thakis would you be able to approve the fix? Otherwise, I am also ok to revert it until I figure out how to commit the fix.

The test is already broken, so speculatively landing that and seeing if it helps sounds fine. But let's try to not keep the build broken for hours. (Was afk for a bit, and it looks like it's still broken now that I look again.)

I have just committed the fix after verifying that it indeed makes the test pass: https://github.com/llvm/llvm-project/commit/b71add9777bed67e05206fa1fdb665f3e21a13ab

clang/include/clang/Driver/Options.td

 def cl_denorms_are_zero : Flag<["-"], "cl-denorms-are-zero">, Group<opencl_Group>,
   HelpText<"OpenCL only. Allow denormals to be flushed to zero.">;
 def cl_fp32_correctly_rounded_divide_sqrt : Flag<["-"], "cl-fp32-correctly-rounded-divide-sqrt">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Specify that single precision floating-point divide and sqrt used in the program source are correctly rounded.">,
   MarshallingInfoFlag<CodeGenOpts<"CorrectlyRoundedDivSqrt">>;
 def cl_uniform_work_group_size : Flag<["-"], "cl-uniform-work-group-size">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Defines that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel">,
   MarshallingInfoFlag<CodeGenOpts<"UniformWGSize">>;
+def cl_no_stdinc : Flag<["-"], "cl-no-stdinc">, Group<opencl_Group>,
+  HelpText<"OpenCL only. Disables all standard includes containing non-native compiler types and functions.">;
 def client__name : JoinedOrSeparate<["-"], "client_name">;
 def combine : Flag<["-", "--"], "combine">, Flags<[NoXarchOption, Unsupported]>;
 def compatibility__version : JoinedOrSeparate<["-"], "compatibility_version">;
 def config : Separate<["--"], "config">, Flags<[NoXarchOption]>,
   HelpText<"Specifies configuration file">;
 def config_system_dir_EQ : Joined<["--"], "config-system-dir=">, Flags<[NoXarchOption, HelpHidden]>,
   HelpText<"System directory for configuration files">;
 def config_user_dir_EQ : Joined<["--"], "config-user-dir=">, Flags<[NoXarchOption, HelpHidden]>,

clang/include/clang/Driver/Types.h

#undef TYPE
   bool isCuda(ID Id);
 
   /// isHIP - Is this a HIP input.
   bool isHIP(ID Id);
 
   /// isObjC - Is this an "ObjC" input (Obj-C and Obj-C++ sources and headers).
   bool isObjC(ID Id);
 
+  /// isOpenCL - Is this an "OpenCL" input.
+  bool isOpenCL(ID Id);
+
   /// isFortran - Is this a Fortran input.
   bool isFortran(ID Id);
 
   /// isSrcFile - Is this a source file, i.e. something that still has to be
   /// preprocessed. The logic behind this is the same that decides if the first
   /// compilation phase is a preprocessing one.
   bool isSrcFile(ID Id);
 

clang/lib/Driver/ToolChains/Clang.cpp

if (Arg *A =
     StringRef Val = A->getValue();
     if (std::stoi(Val.str()) <= 0)
       D.Diag(diag::err_drv_trivial_auto_var_init_stop_after_invalid_value);
     CmdArgs.push_back(
         Args.MakeArgString("-ftrivial-auto-var-init-stop-after=" + Val));
   }
 }
 
-static void RenderOpenCLOptions(const ArgList &Args, ArgStringList &CmdArgs) {
+static void RenderOpenCLOptions(const ArgList &Args, ArgStringList &CmdArgs,
+                                types::ID InputType) {
   // cl-denorms-are-zero is not forwarded. It is translated into a generic flag
   // for denormal flushing handling based on the target.
   const unsigned ForwardedArguments[] = {
       options::OPT_cl_opt_disable,
       options::OPT_cl_strict_aliasing,
       options::OPT_cl_single_precision_constant,
       options::OPT_cl_finite_math_only,
       options::OPT_cl_kernel_arg_info,
       options::OPT_cl_unsafe_math_optimizations,
       options::OPT_cl_fast_relaxed_math,
       options::OPT_cl_mad_enable,
       options::OPT_cl_no_signed_zeros,
       options::OPT_cl_fp32_correctly_rounded_divide_sqrt,
       options::OPT_cl_uniform_work_group_size
   };
 
   if (Arg *A = Args.getLastArg(options::OPT_cl_std_EQ)) {
     std::string CLStdStr = std::string("-cl-std=") + A->getValue();
     CmdArgs.push_back(Args.MakeArgString(CLStdStr));
   }
 
   for (const auto &Arg : ForwardedArguments)
     if (const auto *A = Args.getLastArg(Arg))
       CmdArgs.push_back(Args.MakeArgString(A->getOption().getPrefixedName()));
+
+  // Only add the default headers if we are compiling OpenCL sources.
+  if ((types::isOpenCL(InputType) || Args.hasArg(options::OPT_cl_std_EQ)) &&
+      !Args.hasArg(options::OPT_cl_no_stdinc)) {
+    CmdArgs.push_back("-finclude-default-header");
+    CmdArgs.push_back("-fdeclare-opencl-builtins");
+  }
 }
 
 static void RenderARCMigrateToolOptions(const Driver &D, const ArgList &Args,
                                         ArgStringList &CmdArgs) {
   bool ARCMTEnabled = false;
   if (!Args.hasArg(options::OPT_fno_objc_arc, options::OPT_fobjc_arc)) {
     if (const Arg *A = Args.getLastArg(options::OPT_ccc_arcmt_check,
                                        options::OPT_ccc_arcmt_modify,
if (A->getOption().matches(options::OPT_mrestrict_it)) {
              (Triple.getArch() == llvm::Triple::arm ||
               Triple.getArch() == llvm::Triple::thumb)) {
     // Windows on ARM expects restricted IT blocks
     CmdArgs.push_back("-mllvm");
     CmdArgs.push_back("-arm-restrict-it");
   }
 
   // Forward -cl options to -cc1
-  RenderOpenCLOptions(Args, CmdArgs);
+  RenderOpenCLOptions(Args, CmdArgs, InputType);
 
   if (IsHIP) {
     if (Args.hasFlag(options::OPT_fhip_new_launch_api,
                      options::OPT_fno_hip_new_launch_api, true))
       CmdArgs.push_back("-fhip-new-launch-api");
     if (Args.hasFlag(options::OPT_fgpu_allow_device_init,
                      options::OPT_fno_gpu_allow_device_init, false))
       CmdArgs.push_back("-fgpu-allow-device-init");

clang/lib/Driver/Types.cpp

bool types::isObjC(ID Id) {
   case TY_ObjC: case TY_PP_ObjC: case TY_PP_ObjC_Alias:
   case TY_ObjCXX: case TY_PP_ObjCXX:
   case TY_ObjCHeader: case TY_PP_ObjCHeader:
   case TY_ObjCXXHeader: case TY_PP_ObjCXXHeader: case TY_PP_ObjCXX_Alias:
     return true;
   }
 }
 
+bool types::isOpenCL(ID Id) { return Id == TY_CL; }
+
 bool types::isCXX(ID Id) {
   switch (Id) {
   default:
     return false;
 
   case TY_CXX: case TY_PP_CXX:
   case TY_ObjCXX: case TY_PP_ObjCXX: case TY_PP_ObjCXX_Alias:
   case TY_CXXHeader: case TY_PP_CXXHeader:

clang/test/Driver/default-includes.cl

+// RUN: %clang %s -Xclang -verify -fsyntax-only
+// RUN: %clang %s -cl-no-stdinc -Xclang -verify -DNOINC -fsyntax-only
+
+#ifndef NOINC
+//expected-no-diagnostics
+#endif
+
+void test() {
+int i = get_global_id(0);
+#ifdef NOINC
+//expected-error@-2{{implicit declaration of function 'get_global_id' is invalid in OpenCL}}
+#endif
+}

clang/unittests/AST/MatchVerifier.h

case Lang_CXX17:
     Args.push_back("-std=c++17");
     FileName = "input.cc";
     break;
   case Lang_CXX20:
     Args.push_back("-std=c++20");
     FileName = "input.cc";
     break;
   case Lang_OpenCL:
+    Args.push_back("-cl-no-stdinc");
     FileName = "input.cl";
     break;
   case Lang_OBJCXX:
     FileName = "input.mm";
     break;
   }
 
   // Default to failure in case callback is never called