[OpenCL] Add builtin declarations by default.

This change enables the builtin function declarations
in clang driver 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 with clang to activate the old behavior.

Tags: #clang

Differential Revision: https://reviews.llvm.org/D96515

clang/include/clang/Driver/Options.td

@@ -818,6 +818,8 @@ def cl_fp32_correctly_rounded_divide_sqrt : Flag<["-"], "cl-fp32-correctly-round
 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">;

clang/include/clang/Driver/Types.h

@@ -81,6 +81,9 @@ namespace types {
   /// 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);
 

clang/lib/Driver/ToolChains/Clang.cpp

@@ -3193,7 +3193,8 @@ static void RenderTrivialAutoVarInitOptions(const Driver &D,
   }
 }
 
-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[] = {
@@ -3218,6 +3219,13 @@ static void RenderOpenCLOptions(const ArgList &Args, ArgStringList &CmdArgs) {
   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,
@@ -5726,7 +5734,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   }
 
   // Forward -cl options to -cc1
-  RenderOpenCLOptions(Args, CmdArgs);
+  RenderOpenCLOptions(Args, CmdArgs, InputType);
 
   if (IsHIP) {
     if (Args.hasFlag(options::OPT_fhip_new_launch_api,

clang/lib/Driver/Types.cpp

@@ -160,6 +160,8 @@ bool types::isObjC(ID Id) {
   }
 }
 
+bool types::isOpenCL(ID Id) { return Id == TY_CL; }
+
 bool types::isCXX(ID Id) {
   switch (Id) {
   default:

clang/test/Driver/default-includes.cl

@@ -0,0 +1,13 @@
+// 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

@@ -117,6 +117,7 @@ MatchVerifier<NodeType>::match(const std::string &Code,
     FileName = "input.cc";
     break;
   case Lang_OpenCL:
+    Args.push_back("-cl-no-stdinc");
     FileName = "input.cl";
     break;
   case Lang_OBJCXX: