[clang-tools-extra] a9d68a5 - Generalize calls to ImplicitlyDefineFunction

Aaron Ballman via cfe-commits cfe-commits at lists.llvm.org
Sat Apr 30 07:04:03 PDT 2022


Author: Aaron Ballman
Date: 2022-04-30T10:03:51-04:00
New Revision: a9d68a5524dea113cace5983697786599cbdce9a

URL: https://github.com/llvm/llvm-project/commit/a9d68a5524dea113cace5983697786599cbdce9a
DIFF: https://github.com/llvm/llvm-project/commit/a9d68a5524dea113cace5983697786599cbdce9a.diff

LOG: 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.

Added: 
    

Modified: 
    clang-tools-extra/clangd/IncludeFixer.cpp
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/include/clang/Basic/LangOptions.h
    clang/lib/Sema/SemaDecl.cpp
    clang/lib/Sema/SemaExpr.cpp
    clang/test/Driver/default-includes.cl
    clang/test/Preprocessor/macro_variadic.cl
    clang/test/SemaOpenCL/arm-integer-dot-product.cl
    clang/test/SemaOpenCL/clang-builtin-version.cl
    clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
    clang/test/SemaOpenCL/invalid-block.cl
    clang/test/SemaOpenCL/to_addr_builtin.cl

Removed: 
    


################################################################################
diff  --git a/clang-tools-extra/clangd/IncludeFixer.cpp b/clang-tools-extra/clangd/IncludeFixer.cpp
index 69ec27a024e1f..7fcb01e6e957a 100644
--- a/clang-tools-extra/clangd/IncludeFixer.cpp
+++ b/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

diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 69093a6e51dc3..084d67af570e2 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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<

diff  --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index 9d852d72e593f..eb4c7c4c7d93e 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/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;

diff  --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index da6d35ca98a44..5890bbc7d574b 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/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

diff  --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 90f96a3660dea..67012527888d3 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/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);
   }

diff  --git a/clang/test/Driver/default-includes.cl b/clang/test/Driver/default-includes.cl
index 2596ccca96382..5b96052a54cf1 100644
--- a/clang/test/Driver/default-includes.cl
+++ b/clang/test/Driver/default-includes.cl
@@ -8,6 +8,6 @@
 void test() {
 int i = get_global_id(0);
 #ifdef NOINC
-//expected-error at -2{{implicit declaration of function 'get_global_id' is invalid in OpenCL}}
+//expected-error at -2{{use of undeclared identifier 'get_global_id'}}
 #endif
 }

diff  --git a/clang/test/Preprocessor/macro_variadic.cl b/clang/test/Preprocessor/macro_variadic.cl
index 04e96e3c21465..ee710e7b06dec 100644
--- a/clang/test/Preprocessor/macro_variadic.cl
+++ b/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 at -2{{implicit declaration of function 'func' is invalid in OpenCL}}
-#else
-// expected-error at -4{{use of undeclared identifier 'func'}}
-#endif
+  VAR_FUNC(1, 2, 3); // expected-error {{use of undeclared identifier 'func'}}
   VAR_PRINTF("%i", 1);
 }

diff  --git a/clang/test/SemaOpenCL/arm-integer-dot-product.cl b/clang/test/SemaOpenCL/arm-integer-dot-product.cl
index d61c6a1c506a4..bff3ab4aa42f8 100644
--- a/clang/test/SemaOpenCL/arm-integer-dot-product.cl
+++ b/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'}}
 }
 

diff  --git a/clang/test/SemaOpenCL/clang-builtin-version.cl b/clang/test/SemaOpenCL/clang-builtin-version.cl
index 71c6da56f862d..ec6eecee3106c 100644
--- a/clang/test/SemaOpenCL/clang-builtin-version.cl
+++ b/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 at -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 at -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 at -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 at -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 at -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 at -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 at -1{{expected identifier or '('}}
+             // expected-note at -2{{to match this '('}}
+  boo(); // expected-error{{use of undeclared identifier 'boo'}}
+         // expected-error at -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 at -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'}}
 }

diff  --git a/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl b/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
index 416f5342ef22e..737632fdf07b1 100644
--- a/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
+++ b/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 at -1{{implicit declaration of function 'atomic_exchange' is invalid in OpenCL}}
+  // expected-error at -1{{use of undeclared identifier 'atomic_exchange'}}
 
   atomic_exchange_explicit(a_int, d, memory_order_seq_cst);
   // expected-error at -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 at -1{{implicit declaration of function 'convert_double_rtp' is invalid in OpenCL}}
+  // expected-error at -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 at -2{{implicit declaration of function 'convert_float_sat' is invalid in OpenCL}}
-  // expected-error at -3{{implicit conversion from 'int' to 'float' may lose precision}}
+  // expected-error at -2{{use of undeclared identifier 'convert_float_sat'}}
 #else
-  // expected-error at -5{{use of undeclared identifier 'convert_float_sat'; did you mean 'convert_float'?}}
-  // expected-note at -6{{'convert_float' declared here}}
+  // expected-error at -4{{use of undeclared identifier 'convert_float_sat'; did you mean 'convert_float'?}}
+  // expected-note at -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 at -2{{implicit declaration of function 'get_sub_group_size' is invalid in OpenCL}}
-  // expected-error at -3{{implicit conversion changes signedness}}
+  // expected-error at -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 at -5{{implicit declaration of function 'get_sub_group_eq_mask' is invalid in OpenCL}}
-  // expected-error at -6{{implicit conversion changes signedness}}
-  // expected-error at -6{{implicit declaration of function 'sub_group_non_uniform_scan_inclusive_or' is invalid in OpenCL}}
-  // expected-error at -6{{implicit declaration of function 'sub_group_clustered_reduce_logical_xor' is invalid in OpenCL}}
-  // expected-error at -6{{implicit declaration of function 'sub_group_broadcast' is invalid in OpenCL}}
-  // expected-error at -7{{implicit conversion loses integer precision}}
+  // expected-error at -5{{use of undeclared identifier 'get_sub_group_eq_mask'}}
+  // expected-error at -5{{use of undeclared identifier 'sub_group_non_uniform_scan_inclusive_or'}}
+  // expected-error at -5{{use of undeclared identifier 'sub_group_clustered_reduce_logical_xor'}}
+  // expected-error at -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 at -2{{implicit declaration of function 'get_enqueued_local_size' is invalid in OpenCL}}
+// expected-error at -2{{use of undeclared identifier 'get_enqueued_local_size'}}
 #endif
 }
 

diff  --git a/clang/test/SemaOpenCL/invalid-block.cl b/clang/test/SemaOpenCL/invalid-block.cl
index ddc0f80d1da44..6c918d302f801 100644
--- a/clang/test/SemaOpenCL/invalid-block.cl
+++ b/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'}}
   };
 }
 

diff  --git a/clang/test/SemaOpenCL/to_addr_builtin.cl b/clang/test/SemaOpenCL/to_addr_builtin.cl
index 5b1f7839cdbce..5cd74885aa9df 100644
--- a/clang/test/SemaOpenCL/to_addr_builtin.cl
+++ b/clang/test/SemaOpenCL/to_addr_builtin.cl
@@ -14,37 +14,36 @@ void test(void) {
 
   glob = to_global(glob, loc);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-error at -2{{implicit declaration of function 'to_global' is invalid in OpenCL}}
-  // expected-warning at -3{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_global'}}
 #else
-  // expected-error at -5{{too many arguments to function call, expected 1, have 2}}
+  // expected-error at -4{{too many arguments to function call, expected 1, have 2}}
 #endif
 
   int x;
   glob = to_global(x);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_global'}}
 #else
   // expected-error at -4{{invalid argument x to function: 'to_global', expecting a generic pointer argument}}
 #endif
 
   glob = to_global(con);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_global'}}
 #else
   // expected-error at -4{{invalid argument con to function: 'to_global', expecting a generic pointer argument}}
 #endif
 
   glob = to_global(con_typedef);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_global'}}
 #else
   // expected-error at -4{{invalid argument con_typedef to function: 'to_global', expecting a generic pointer argument}}
 #endif
 
   loc = to_global(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_global'}}
 #else
   // expected-error at -4{{assigning '__global int *' to '__local int *__private' changes address space of pointer}}
   // expected-warning at -5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
@@ -52,24 +51,22 @@ void test(void) {
 
   loc = to_private(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-error at -2{{implicit declaration of function 'to_private' is invalid in OpenCL}}
-  // expected-warning at -3{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_private'}}
 #else
-  // expected-error at -5{{assigning '__private int *' to '__local int *__private' changes address space of pointer}}
-  // expected-warning at -6{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
+  // expected-error at -4{{assigning '__private int *' to '__local int *__private' changes address space of pointer}}
+  // expected-warning at -5{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
 #endif
 
   loc = to_local(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-error at -2{{implicit declaration of function 'to_local' is invalid in OpenCL}}
-  // expected-warning at -3{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_local'}}
 #else
-  // expected-warning at -5{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
+  // expected-warning at -4{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
 #endif
 
   priv = to_global(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_global'}}
 #else
   // expected-error at -4{{assigning '__global int *' to '__private int *__private' changes address space of pointer}}
   // expected-warning at -5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
@@ -77,7 +74,7 @@ void test(void) {
 
   priv = to_private(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_private'}}
 #else
   // expected-warning at -4{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
 #endif
@@ -85,7 +82,7 @@ void test(void) {
 
   priv = to_local(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_local'}}
 #else
   // expected-error at -4{{assigning '__local int *' to '__private int *__private' changes address space of pointer}}
   // expected-warning at -5{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
@@ -93,14 +90,14 @@ void test(void) {
 
   glob = to_global(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_global'}}
 #else
   // expected-warning at -4{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
 #endif
 
   glob = to_private(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_private'}}
 #else
   // expected-error at -4{{assigning '__private int *' to '__global int *__private' changes address space of pointer}}
   // expected-warning at -5{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
@@ -108,7 +105,7 @@ void test(void) {
 
   glob = to_local(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_local'}}
 #else
   // expected-error at -4{{assigning '__local int *' to '__global int *__private' changes address space of pointer}}
   // expected-warning at -5{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
@@ -116,7 +113,7 @@ void test(void) {
 
   global char *glob_c = to_global(loc);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion initializing '__global char *__private' with an expression of type 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_global'}}
 #else
   // expected-warning at -4{{incompatible pointer types initializing '__global char *__private' with an expression of type '__global int *'}}
   // expected-warning at -5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
@@ -124,7 +121,7 @@ void test(void) {
 
   glob_wrong_ty = to_global(glob);
 #if (__OPENCL_C_VERSION__ < CL_VERSION_2_0) || (__OPENCL_C_VERSION__ >= CL_VERSION_3_0 && !defined(__opencl_c_generic_address_space))
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global float *__private' from 'int'}}
+  // expected-error at -2{{use of undeclared identifier 'to_global'}}
 #else
   // expected-warning at -4{{incompatible pointer types assigning to '__global float *__private' from '__global int *'}}
   // expected-warning at -5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}


        


More information about the cfe-commits mailing list