[Sema] Change addr space diagnostics in casts to follow C++ style.

This change adds a new diagnostic for mismatching address spaces
to be used for C++ casts (only enabled in C style cast for now,
the rest will follow!).

The change extends C-style cast rules to account for address spaces.
It also adds a separate function for address space cast checking that
can be used to map from a separate address space cast operator
addrspace_cast (to be added as a follow up patch).

Note, that after this change clang will no longer allows arbitrary
address space conversions in reinterpret_casts because they can lead
to accidental errors. The implicit safe conversions would still be
allowed.

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

llvm-svn: 355609

2 files changed, 140 insertions, 28 deletions

diff --git a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
index 619ecc4e477..bbd3919b154 100644
--- a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
+++ b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
@@ -129,27 +129,47 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
 
   AS int *var_cast1 = (AS int *)arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-5{{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   AS int *var_cast2 = (AS int *)arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{C-style cast from '__local int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   AS int *var_cast3 = (AS int *)arg_const;
 #ifndef CONSTANT
-// expected-error-re@-2{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{C-style cast from '__constant int *' to '__{{global|generic}} int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   AS int *var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   AS int *var_cast5 = (AS int *)arg_gen;
 #ifdef CONSTANT
-// expected-error@-2{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-5{{C-style cast from '__generic int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   AS int *var_impl;
@@ -200,27 +220,47 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
 
   var_cast1 = (AS int *)arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-5{{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   var_cast2 = (AS int *)arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{C-style cast from '__local int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   var_cast3 = (AS int *)arg_const;
 #ifndef CONSTANT
-// expected-error-re@-2{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{C-style cast from '__constant int *' to '__{{global|generic}} int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   var_cast5 = (AS int *)arg_gen;
 #ifdef CONSTANT
-// expected-error@-2{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-5{{C-style cast from '__generic int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
 #endif
 
   AS int *var_cmp;
diff --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl
index 30f311d6ef1..f5fa43d9222 100644
--- a/clang/test/SemaOpenCL/address-spaces.cl
+++ b/clang/test/SemaOpenCL/address-spaces.cl
@@ -26,24 +26,96 @@ __kernel void foo(__global int *gip) {
 }
 
 void explicit_cast(__global int *g, __local int *l, __constant int *c, __private int *p, const __constant int *cc) {
-  g = (__global int *)l;  // expected-error {{casting '__local int *' to type '__global int *' changes address space of pointer}}
-  g = (__global int *)c;  // expected-error {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
-  g = (__global int *)cc; // expected-error {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
-  g = (__global int *)p;  // expected-error {{casting 'int *' to type '__global int *' changes address space of pointer}}
-
-  l = (__local int *)g;  // expected-error {{casting '__global int *' to type '__local int *' changes address space of pointer}}
-  l = (__local int *)c;  // expected-error {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
-  l = (__local int *)cc; // expected-error {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
-  l = (__local int *)p;  // expected-error {{casting 'int *' to type '__local int *' changes address space of pointer}}
-
-  c = (__constant int *)g; // expected-error {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
-  c = (__constant int *)l; // expected-error {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
-  c = (__constant int *)p; // expected-error {{casting 'int *' to type '__constant int *' changes address space of pointer}}
-
-  p = (__private int *)g;  // expected-error {{casting '__global int *' to type 'int *' changes address space of pointer}}
-  p = (__private int *)l;  // expected-error {{casting '__local int *' to type 'int *' changes address space of pointer}}
-  p = (__private int *)c;  // expected-error {{casting '__constant int *' to type 'int *' changes address space of pointer}}
-  p = (__private int *)cc; // expected-error {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
+  g = (__global int *)l;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting '__local int *' to type '__global int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from '__local int *' to '__global int *' converts between mismatching address spaces}}
+#endif
+  g = (__global int *)c;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from '__constant int *' to '__global int *' converts between mismatching address spaces}}
+#endif
+  g = (__global int *)cc;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from 'const __constant int *' to '__global int *' converts between mismatching address spaces}}
+#endif
+  g = (__global int *)p;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting 'int *' to type '__global int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from 'int *' to '__global int *' converts between mismatching address spaces}}
+#endif
+  l = (__local int *)g;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting '__global int *' to type '__local int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from '__global int *' to '__local int *' converts between mismatching address spaces}}
+#endif
+  l = (__local int *)c;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from '__constant int *' to '__local int *' converts between mismatching address spaces}}
+#endif
+  l = (__local int *)cc;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from 'const __constant int *' to '__local int *' converts between mismatching address spaces}}
+#endif
+  l = (__local int *)p;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting 'int *' to type '__local int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from 'int *' to '__local int *' converts between mismatching address spaces}}
+#endif
+  c = (__constant int *)g;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
+  c = (__constant int *)l;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from '__local int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
+  c = (__constant int *)p;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting 'int *' to type '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from 'int *' to '__constant int *' converts between mismatching address spaces}}
+#endif
+  p = (__private int *)g;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting '__global int *' to type 'int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from '__global int *' to 'int *' converts between mismatching address spaces}}
+#endif
+  p = (__private int *)l;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting '__local int *' to type 'int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from '__local int *' to 'int *' converts between mismatching address spaces}}
+#endif
+  p = (__private int *)c;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting '__constant int *' to type 'int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from '__constant int *' to 'int *' converts between mismatching address spaces}}
+#endif
+  p = (__private int *)cc;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{C-style cast from 'const __constant int *' to 'int *' converts between mismatching address spaces}}
+#endif
 }
 
 void ok_explicit_casts(__global int *g, __global int *g2, __local int *l, __local int *l2, __private int *p, __private int *p2) {