Generalize calls to ImplicitlyDefineFunction

In C++ and C2x, we would avoid calling ImplicitlyDefineFunction at all,
but in OpenCL mode we would still call the function and have it produce
an error diagnostic. Instead, we now have a helper function to
determine when implicit function definitions are allowed and we use
that to determine whether to call ImplicitlyDefineFunction so that the
behavior is more consistent across language modes.

This changes the diagnostic behavior from telling the users that an
implicit function declaration is not allowed in OpenCL to reporting use
of an unknown identifier and going through typo correction, as done in
C++ and C2x.

clang-tools-extra/clangd/IncludeFixer.cpp

@@ -198,7 +198,6 @@ std::vector<Fix> IncludeFixer::fix(DiagnosticsEngine::Level DiagLevel,
   case diag::err_no_member_template_suggest:
   case diag::warn_implicit_function_decl:
   case diag::ext_implicit_function_decl_c99:
-  case diag::err_opencl_implicit_function_decl:
     dlog("Unresolved name at {0}, last typo was {1}",
          Info.getLocation().printToString(Info.getSourceManager()),
          LastUnresolvedName

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10198,8 +10198,6 @@ def err_opencl_scalar_type_rank_greater_than_vector_type : Error<
     "element. (%0 and %1)">;
 def err_bad_kernel_param_type : Error<
   "%0 cannot be used as the type of a kernel parameter">;
-def err_opencl_implicit_function_decl : Error<
-  "implicit declaration of function %0 is invalid in OpenCL">;
 def err_record_with_pointers_kernel_param : Error<
   "%select{struct|union}0 kernel parameters may not contain pointers">;
 def note_within_field_of_type : Note<

clang/include/clang/Basic/LangOptions.h

@@ -527,6 +527,12 @@ class LangOptions : public LangOptionsBase {
     return CPlusPlus || C2x || DisableKNRFunctions;
   }
 
+  /// Returns true if implicit function declarations are allowed in the current
+  /// language mode.
+  bool implicitFunctionsAllowed() const {
+    return !requiresStrictPrototypes() && !OpenCL;
+  }
+
   /// Check if return address signing is enabled.
   bool hasSignReturnAddress() const {
     return getSignReturnAddressScope() != SignReturnAddressScopeKind::None;

clang/lib/Sema/SemaDecl.cpp

@@ -934,8 +934,9 @@ Sema::NameClassification Sema::ClassifyName(Scope *S, CXXScopeSpec &SS,
       //   appeared.
       //
       // We also allow this in C99 as an extension. However, this is not
-      // allowed in C2x as there are no functions without prototypes there.
-      if (!getLangOpts().C2x) {
+      // allowed in all language modes as functions without prototypes may not
+      // be supported.
+      if (getLangOpts().implicitFunctionsAllowed()) {
         if (NamedDecl *D = ImplicitlyDefineFunction(NameLoc, *Name, S))
           return NameClassification::NonType(D);
       }
@@ -15273,7 +15274,8 @@ void Sema::ActOnFinishDelayedAttribute(Scope *S, Decl *D,
 NamedDecl *Sema::ImplicitlyDefineFunction(SourceLocation Loc,
                                           IdentifierInfo &II, Scope *S) {
   // It is not valid to implicitly define a function in C2x.
-  assert(!LangOpts.C2x && "Cannot implicitly define a function in C2x");
+  assert(LangOpts.implicitFunctionsAllowed() &&
+         "Implicit function declarations aren't allowed in this language mode");
 
   // Find the scope in which the identifier is injected and the corresponding
   // DeclContext.
@@ -15318,8 +15320,6 @@ NamedDecl *Sema::ImplicitlyDefineFunction(SourceLocation Loc,
   if (II.getName().startswith("__builtin_"))
     diag_id = diag::warn_builtin_unknown;
   // OpenCL v2.0 s6.9.u - Implicit function declaration is not supported.
-  else if (getLangOpts().OpenCL)
-    diag_id = diag::err_opencl_implicit_function_decl;
   else if (getLangOpts().C99)
     diag_id = diag::ext_implicit_function_decl_c99;
   else

clang/lib/Sema/SemaExpr.cpp

@@ -2553,10 +2553,10 @@ Sema::ActOnIdExpression(Scope *S, CXXScopeSpec &SS,
   if (R.isAmbiguous())
     return ExprError();
 
-  // This could be an implicitly declared function reference (legal in C90,
-  // extension in C99, forbidden in C++ and C2x).
-  if (R.empty() && HasTrailingLParen && II && !getLangOpts().CPlusPlus &&
-      !getLangOpts().C2x) {
+  // This could be an implicitly declared function reference if the language
+  // mode allows it as a feature.
+  if (R.empty() && HasTrailingLParen && II &&
+      getLangOpts().implicitFunctionsAllowed()) {
     NamedDecl *D = ImplicitlyDefineFunction(NameLoc, *II, S);
     if (D) R.addDecl(D);
   }

clang/test/Driver/default-includes.cl

@@ -8,6 +8,6 @@
 void test() {
 int i = get_global_id(0);
 #ifdef NOINC
-//expected-error@-2{{implicit declaration of function 'get_global_id' is invalid in OpenCL}}
+//expected-error@-2{{use of undeclared identifier 'get_global_id'}}
 #endif
 }

clang/test/Preprocessor/macro_variadic.cl

@@ -17,11 +17,6 @@ int printf(__constant const char *st, ...);
 
 void foo(void) {
   NO_VAR_FUNC(1, 2, 3);
-  VAR_FUNC(1, 2, 3);
-#if !__OPENCL_CPP_VERSION__
-// expected-error@-2{{implicit declaration of function 'func' is invalid in OpenCL}}
-#else
-// expected-error@-4{{use of undeclared identifier 'func'}}
-#endif
+  VAR_FUNC(1, 2, 3); // expected-error {{use of undeclared identifier 'func'}}
   VAR_PRINTF("%i", 1);
 }

clang/test/SemaOpenCL/arm-integer-dot-product.cl

@@ -7,13 +7,13 @@ void test_negative() {
     short2 sa16, sb16;
     uint ur;
     int sr;
-    ur = arm_dot(ua8, ub8); // expected-error{{implicit declaration of function 'arm_dot' is invalid in OpenCL}}
-    sr = arm_dot(sa8, sb8);
-    ur = arm_dot_acc(ua8, ub8, ur); // expected-error{{implicit declaration of function 'arm_dot_acc' is invalid in OpenCL}}
-    sr = arm_dot_acc(sa8, sb8, sr);
-    ur = arm_dot_acc(ua16, ub16, ur);
-    sr = arm_dot_acc(sa16, sb16, sr);
-    ur = arm_dot_acc_sat(ua8, ub8, ur); // expected-error{{implicit declaration of function 'arm_dot_acc_sat' is invalid in OpenCL}}
-    sr = arm_dot_acc_sat(sa8, sb8, sr);
+    ur = arm_dot(ua8, ub8); // expected-error{{use of undeclared identifier 'arm_dot'}}
+    sr = arm_dot(sa8, sb8); // expected-error{{use of undeclared identifier 'arm_dot'}}
+    ur = arm_dot_acc(ua8, ub8, ur); // expected-error{{use of undeclared identifier 'arm_dot_acc'}}
+    sr = arm_dot_acc(sa8, sb8, sr); // expected-error{{use of undeclared identifier 'arm_dot_acc'}}
+    ur = arm_dot_acc(ua16, ub16, ur); // expected-error{{use of undeclared identifier 'arm_dot_acc'}}
+    sr = arm_dot_acc(sa16, sb16, sr); // expected-error{{use of undeclared identifier 'arm_dot_acc'}}
+    ur = arm_dot_acc_sat(ua8, ub8, ur); // expected-error{{use of undeclared identifier 'arm_dot_acc_sat'}}
+    sr = arm_dot_acc_sat(sa8, sb8, sr); // expected-error{{use of undeclared identifier 'arm_dot_acc_sat'}}
 }
 

clang/test/SemaOpenCL/clang-builtin-version.cl

@@ -6,50 +6,60 @@
 
 kernel void dse_builtins(void) {
   int tmp;
-  enqueue_kernel(tmp, tmp, tmp, ^(void) { // expected-error{{implicit declaration of function 'enqueue_kernel' is invalid in OpenCL}}
+  enqueue_kernel(tmp, tmp, tmp, ^(void) { // expected-error{{use of undeclared identifier 'enqueue_kernel'}}
     return;
   });
-  unsigned size = get_kernel_work_group_size(^(void) { // expected-error{{implicit declaration of function 'get_kernel_work_group_size' is invalid in OpenCL}}
+  unsigned size = get_kernel_work_group_size(^(void) { // expected-error{{use of undeclared identifier 'get_kernel_work_group_size'}}
     return;
   });
-  size = get_kernel_preferred_work_group_size_multiple(^(void) { // expected-error{{implicit declaration of function 'get_kernel_preferred_work_group_size_multiple' is invalid in OpenCL}}
+  size = get_kernel_preferred_work_group_size_multiple(^(void) { // expected-error{{use of undeclared identifier 'get_kernel_preferred_work_group_size_multiple'}}
     return;
   });
 #if (__OPENCL_C_VERSION__ >= CL_VERSION_3_0) && !defined(__opencl_c_device_enqueue)
 // expected-error@-10{{support disabled - compile with -fblocks or for OpenCL C 2.0 or OpenCL C 3.0 with __opencl_c_device_enqueue feature}}
-// expected-error@-8{{support disabled - compile with -fblocks or for OpenCL C 2.0 or OpenCL C 3.0 with __opencl_c_device_enqueue feature}}
-// expected-error@-6{{support disabled - compile with -fblocks or for OpenCL C 2.0 or OpenCL C 3.0 with __opencl_c_device_enqueue feature}}
+// FIXME: the typo correction for the undeclared identifiers finds alternative
+// suggestions, but instantiating the typo correction causes us to
+// re-instantiate the argument to the call, which triggers the support
+// diagnostic a second time.
+// expected-error@-12 2{{support disabled - compile with -fblocks or for OpenCL C 2.0 or OpenCL C 3.0 with __opencl_c_device_enqueue feature}}
+// expected-error@-10 2{{support disabled - compile with -fblocks or for OpenCL C 2.0 or OpenCL C 3.0 with __opencl_c_device_enqueue feature}}
 #endif
 }
 
 void pipe_builtins(void) {
   int tmp;
 
-  foo(void); // expected-error{{implicit declaration of function 'foo' is invalid in OpenCL}}
-  // expected-error@-1{{expected expression}}
-  boo(); // expected-error{{implicit declaration of function 'boo' is invalid in OpenCL}}
+  // FIXME: the typo correction for this case goes off the rails and tries to
+  // convert this mistake into a for loop instead of a local function
+  // declaration.
+  foo(void); // expected-error{{use of undeclared identifier 'foo'; did you mean 'for'?}}
+             // expected-error@-1{{expected identifier or '('}}
+             // expected-note@-2{{to match this '('}}
+  boo(); // expected-error{{use of undeclared identifier 'boo'}}
+         // expected-error@-1{{expected ';' in 'for' statement specifier}}
 
-  read_pipe(tmp, tmp);  // expected-error{{implicit declaration of function 'read_pipe' is invalid in OpenCL}}
-  write_pipe(tmp, tmp); // expected-error{{implicit declaration of function 'write_pipe' is invalid in OpenCL}}
+  read_pipe(tmp, tmp);  // expected-error{{use of undeclared identifier 'read_pipe'}}
+                        // expected-error@-1{{expected ')'}}
+  write_pipe(tmp, tmp); // expected-error{{use of undeclared identifier 'write_pipe'}}
 
-  reserve_read_pipe(tmp, tmp);  // expected-error{{implicit declaration of function 'reserve_read_pipe' is invalid in OpenCL}}
-  reserve_write_pipe(tmp, tmp); // expected-error{{implicit declaration of function 'reserve_write_pipe' is invalid in OpenCL}}
+  reserve_read_pipe(tmp, tmp);  // expected-error{{use of undeclared identifier 'reserve_read_pipe'}}
+  reserve_write_pipe(tmp, tmp); // expected-error{{use of undeclared identifier 'reserve_write_pipe'}}
 
-  work_group_reserve_read_pipe(tmp, tmp);  // expected-error{{implicit declaration of function 'work_group_reserve_read_pipe' is invalid in OpenCL}}
-  work_group_reserve_write_pipe(tmp, tmp); // expected-error{{implicit declaration of function 'work_group_reserve_write_pipe' is invalid in OpenCL}}
+  work_group_reserve_read_pipe(tmp, tmp);  // expected-error{{use of undeclared identifier 'work_group_reserve_read_pipe'}}
+  work_group_reserve_write_pipe(tmp, tmp); // expected-error{{use of undeclared identifier 'work_group_reserve_write_pipe'}}
 
-  sub_group_reserve_write_pipe(tmp, tmp); // expected-error{{implicit declaration of function 'sub_group_reserve_write_pipe' is invalid in OpenCL}}
-  sub_group_reserve_read_pipe(tmp, tmp);  // expected-error{{implicit declaration of function 'sub_group_reserve_read_pipe' is invalid in OpenCL}}
+  sub_group_reserve_write_pipe(tmp, tmp); // expected-error{{use of undeclared identifier 'sub_group_reserve_write_pipe'}}
+  sub_group_reserve_read_pipe(tmp, tmp);  // expected-error{{use of undeclared identifier 'sub_group_reserve_read_pipe'}}
 
-  commit_read_pipe(tmp, tmp);  // expected-error{{implicit declaration of function 'commit_read_pipe' is invalid in OpenCL}}
-  commit_write_pipe(tmp, tmp); // expected-error{{implicit declaration of function 'commit_write_pipe' is invalid in OpenCL}}
+  commit_read_pipe(tmp, tmp);  // expected-error{{use of undeclared identifier 'commit_read_pipe'}}
+  commit_write_pipe(tmp, tmp); // expected-error{{use of undeclared identifier 'commit_write_pipe'}}
 
-  work_group_commit_read_pipe(tmp, tmp);  // expected-error{{implicit declaration of function 'work_group_commit_read_pipe' is invalid in OpenCL}}
-  work_group_commit_write_pipe(tmp, tmp); // expected-error{{implicit declaration of function 'work_group_commit_write_pipe' is invalid in OpenCL}}
+  work_group_commit_read_pipe(tmp, tmp);  // expected-error{{use of undeclared identifier 'work_group_commit_read_pipe'}}
+  work_group_commit_write_pipe(tmp, tmp); // expected-error{{use of undeclared identifier 'work_group_commit_write_pipe'}}
 
-  sub_group_commit_write_pipe(tmp, tmp); // expected-error{{implicit declaration of function 'sub_group_commit_write_pipe' is invalid in OpenCL}}
-  sub_group_commit_read_pipe(tmp, tmp);  // expected-error{{implicit declaration of function 'sub_group_commit_read_pipe' is invalid in OpenCL}}
+  sub_group_commit_write_pipe(tmp, tmp); // expected-error{{use of undeclared identifier 'sub_group_commit_write_pipe'}}
+  sub_group_commit_read_pipe(tmp, tmp);  // expected-error{{use of undeclared identifier 'sub_group_commit_read_pipe'}}
 
-  get_pipe_num_packets(tmp); // expected-error{{implicit declaration of function 'get_pipe_num_packets' is invalid in OpenCL}}
-  get_pipe_max_packets(tmp); // expected-error{{implicit declaration of function 'get_pipe_max_packets' is invalid in OpenCL}}
+  get_pipe_num_packets(tmp); // expected-error{{use of undeclared identifier 'get_pipe_num_packets'}}
+  get_pipe_max_packets(tmp); // expected-error{{use of undeclared identifier 'get_pipe_max_packets'}}
 }

clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl

@@ -194,7 +194,7 @@ void test_atomics_without_scope_device(volatile __generic atomic_int *a_int) {
   int d;
 
   atomic_exchange(a_int, d);
-  // expected-error@-1{{implicit declaration of function 'atomic_exchange' is invalid in OpenCL}}
+  // expected-error@-1{{use of undeclared identifier 'atomic_exchange'}}
 
   atomic_exchange_explicit(a_int, d, memory_order_seq_cst);
   // expected-error@-1{{no matching function for call to 'atomic_exchange_explicit'}}
@@ -225,7 +225,7 @@ kernel void basic_conversion(void) {
 
 #ifdef NO_FP64
   (void)convert_double_rtp(f);
-  // expected-error@-1{{implicit declaration of function 'convert_double_rtp' is invalid in OpenCL}}
+  // expected-error@-1{{use of undeclared identifier 'convert_double_rtp'}}
 #else
   double d;
   f = convert_float(d);
@@ -240,11 +240,10 @@ kernel void basic_conversion_neg(void) {
 
   f = convert_float_sat(i);
 #if !defined(__OPENCL_CPP_VERSION__)
-  // expected-error@-2{{implicit declaration of function 'convert_float_sat' is invalid in OpenCL}}
-  // expected-error@-3{{implicit conversion from 'int' to 'float' may lose precision}}
+  // expected-error@-2{{use of undeclared identifier 'convert_float_sat'}}
 #else
-  // expected-error@-5{{use of undeclared identifier 'convert_float_sat'; did you mean 'convert_float'?}}
-  // expected-note@-6{{'convert_float' declared here}}
+  // expected-error@-4{{use of undeclared identifier 'convert_float_sat'; did you mean 'convert_float'?}}
+  // expected-note@-5{{'convert_float' declared here}}
 #endif
 }
 
@@ -312,8 +311,7 @@ kernel void basic_image_writeonly(write_only image1d_buffer_t image_write_only_i
 kernel void basic_subgroup(global uint *out) {
   out[0] = get_sub_group_size();
 #if __OPENCL_C_VERSION__ <= CL_VERSION_1_2 && !defined(__OPENCL_CPP_VERSION__)
-  // expected-error@-2{{implicit declaration of function 'get_sub_group_size' is invalid in OpenCL}}
-  // expected-error@-3{{implicit conversion changes signedness}}
+  // expected-error@-2{{use of undeclared identifier 'get_sub_group_size'}}
 #endif
 
 // Only test when the base header is included, because we need the enum declarations.
@@ -328,12 +326,10 @@ kernel void extended_subgroup(global uint4 *out, global int *scalar, global char
   scalar[1] = sub_group_clustered_reduce_logical_xor(2, 4);
   *c2 = sub_group_broadcast(*c2, 2);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 && !defined(__OPENCL_CPP_VERSION__)
-  // expected-error@-5{{implicit declaration of function 'get_sub_group_eq_mask' is invalid in OpenCL}}
-  // expected-error@-6{{implicit conversion changes signedness}}
-  // expected-error@-6{{implicit declaration of function 'sub_group_non_uniform_scan_inclusive_or' is invalid in OpenCL}}
-  // expected-error@-6{{implicit declaration of function 'sub_group_clustered_reduce_logical_xor' is invalid in OpenCL}}
-  // expected-error@-6{{implicit declaration of function 'sub_group_broadcast' is invalid in OpenCL}}
-  // expected-error@-7{{implicit conversion loses integer precision}}
+  // expected-error@-5{{use of undeclared identifier 'get_sub_group_eq_mask'}}
+  // expected-error@-5{{use of undeclared identifier 'sub_group_non_uniform_scan_inclusive_or'}}
+  // expected-error@-5{{use of undeclared identifier 'sub_group_clustered_reduce_logical_xor'}}
+  // expected-error@-5{{use of undeclared identifier 'sub_group_broadcast'}}
 #endif
 }
 
@@ -376,7 +372,7 @@ kernel void basic_work_item(void) {
 
   get_enqueued_local_size(ui);
 #if !defined(__OPENCL_CPP_VERSION__) && __OPENCL_C_VERSION__ < CL_VERSION_2_0
-// expected-error@-2{{implicit declaration of function 'get_enqueued_local_size' is invalid in OpenCL}}
+// expected-error@-2{{use of undeclared identifier 'get_enqueued_local_size'}}
 #endif
 }
 

clang/test/SemaOpenCL/invalid-block.cl

@@ -35,7 +35,7 @@ bl_t f3b(bl_t bl);
 void f3c(void) {
   // Block with a block argument.
   int (^const bl2)(bl_t block_arg) = ^(void) { // expected-error{{declaring function parameter of type '__private bl_t' (aka 'int (__generic ^const __private)(void)') is not allowed}}
-    return block_arg(); // expected-error{{implicit declaration of function 'block_arg' is invalid in OpenCL}}
+    return block_arg(); // expected-error{{use of undeclared identifier 'block_arg'}}
   };
 }