[Sema][OpenCL] Make address space conversions a bit stricter.

The semantics for converting nested pointers between address
spaces are not very well defined. Some conversions which do not
really carry any meaning only produce warnings, and in some cases
warnings hide invalid conversions, such as 'global int*' to
'local float*'!

This patch changes the logic in checkPointerTypesForAssignment
and checkAddressSpaceCast to fail properly on implicit conversions
that should definitely not be permitted. We also dig deeper into the
pointer types and warn on explicit conversions where the address
space in a nested pointer changes, regardless of whether the address
space is compatible with the corresponding pointer nesting level
on the destination type.

Fixes PR39674!

Patch by ebevhan (Bevin Hansson)!

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



git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@360258 91177308-0d34-0410-b5e6-96231b3b80d8

4 files changed, 103 insertions, 3 deletions

diff --git a/test/SemaOpenCL/address-spaces.cl b/test/SemaOpenCL/address-spaces.cl
index f5fa43d922..13e9442b95 100644
--- a/test/SemaOpenCL/address-spaces.cl
+++ b/test/SemaOpenCL/address-spaces.cl
@@ -124,6 +124,106 @@ void ok_explicit_casts(__global int *g, __global int *g2, __local int *l, __loca
   p = (__private int *)p2;
 }
 
+#if !__OPENCL_CPP_VERSION__
+void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
+  g = gg;    // expected-error {{assigning '__global int **' to '__global int *' changes address space of pointer}}
+  g = l;     // expected-error {{assigning '__local int *' to '__global int *' changes address space of pointer}}
+  g = ll;    // expected-error {{assigning '__local int **' to '__global int *' changes address space of pointer}}
+  g = gg_f;  // expected-error {{assigning '__global float **' to '__global int *' changes address space of pointer}}
+  g = (__global int *)gg_f; // expected-error {{casting '__global float **' to type '__global int *' changes address space of pointer}}
+
+  gg = g;    // expected-error {{assigning '__global int *' to '__global int **' changes address space of pointer}}
+  gg = l;    // expected-error {{assigning '__local int *' to '__global int **' changes address space of pointer}}
+  gg = ll;   // expected-error {{assigning '__local int **' to '__global int **' changes address space of nested pointer}}
+  gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int **' from '__global float **'}}
+  gg = (__global int * __private *)gg_f;
+
+  l = g;     // expected-error {{assigning '__global int *' to '__local int *' changes address space of pointer}}
+  l = gg;    // expected-error {{assigning '__global int **' to '__local int *' changes address space of pointer}}
+  l = ll;    // expected-error {{assigning '__local int **' to '__local int *' changes address space of pointer}}
+  l = gg_f;  // expected-error {{assigning '__global float **' to '__local int *' changes address space of pointer}}
+  l = (__local int *)gg_f; // expected-error {{casting '__global float **' to type '__local int *' changes address space of pointer}}
+
+  ll = g;    // expected-error {{assigning '__global int *' to '__local int **' changes address space of pointer}}
+  ll = gg;   // expected-error {{assigning '__global int **' to '__local int **' changes address space of nested pointer}}
+  ll = l;    // expected-error {{assigning '__local int *' to '__local int **' changes address space of pointer}}
+  ll = gg_f; // expected-error {{assigning '__global float **' to '__local int **' changes address space of nested pointer}}
+  ll = (__local int * __private *)gg_f; // expected-warning {{casting '__global float **' to type '__local int **' discards qualifiers in nested pointer types}}
+
+  gg_f = g;  // expected-error {{assigning '__global int *' to '__global float **' changes address space of pointer}}
+  gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float **' from '__global int **'}}
+  gg_f = l;  // expected-error {{assigning '__local int *' to '__global float **' changes address space of pointer}}
+  gg_f = ll; // expected-error {{assigning '__local int **' to '__global float **' changes address space of nested pointer}}
+  gg_f = (__global float * __private *)gg;
+
+  // FIXME: This doesn't seem right. This should be an error, not a warning.
+  __local int * __global * __private * lll;
+  lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global **' from '__global int **'}}
+
+  typedef __local int * l_t;
+  typedef __global int * g_t;
+  __private l_t * pl;
+  __private g_t * pg;
+  gg = pl;  // expected-error {{assigning 'l_t *' (aka '__local int **') to '__global int **' changes address space of nested pointer}}
+  pl = gg;  // expected-error {{assigning '__global int **' to 'l_t *' (aka '__local int **') changes address space of nested pointer}}
+  gg = pg;
+  pg = gg;
+  pg = pl;  // expected-error {{assigning 'l_t *' (aka '__local int **') to 'g_t *' (aka '__global int **') changes address space of nested pointer}}
+  pl = pg;  // expected-error {{assigning 'g_t *' (aka '__global int **') to 'l_t *' (aka '__local int **') changes address space of nested pointer}}
+
+  ll = (__local int * __private *)(void *)gg;
+  void *vp = ll;
+}
+#else
+void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
+  g = gg;    // expected-error {{assigning to '__global int *' from incompatible type '__global int **'}}
+  g = l;     // expected-error {{assigning to '__global int *' from incompatible type '__local int *'}}
+  g = ll;    // expected-error {{assigning to '__global int *' from incompatible type '__local int **'}}
+  g = gg_f;  // expected-error {{assigning to '__global int *' from incompatible type '__global float **'}}
+  g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__global int *' converts between mismatching address spaces}}
+
+  gg = g;    // expected-error {{assigning to '__global int **' from incompatible type '__global int *'; take the address with &}}
+  gg = l;    // expected-error {{assigning to '__global int **' from incompatible type '__local int *'}}
+  gg = ll;   // expected-error {{assigning to '__global int **' from incompatible type '__local int **'}}
+  gg = gg_f; // expected-error {{assigning to '__global int **' from incompatible type '__global float **'}}
+  gg = (__global int * __private *)gg_f;
+
+  l = g;     // expected-error {{assigning to '__local int *' from incompatible type '__global int *'}}
+  l = gg;    // expected-error {{assigning to '__local int *' from incompatible type '__global int **'}}
+  l = ll;    // expected-error {{assigning to '__local int *' from incompatible type '__local int **'}}
+  l = gg_f;  // expected-error {{assigning to '__local int *' from incompatible type '__global float **'}}
+  l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__local int *' converts between mismatching address spaces}}
+
+  ll = g;    // expected-error {{assigning to '__local int **' from incompatible type '__global int *'}}
+  ll = gg;   // expected-error {{assigning to '__local int **' from incompatible type '__global int **'}}
+  ll = l;    // expected-error {{assigning to '__local int **' from incompatible type '__local int *'; take the address with &}}
+  ll = gg_f; // expected-error {{assigning to '__local int **' from incompatible type '__global float **'}}
+  // FIXME: The below becomes a reinterpret_cast, and therefore does not emit an error
+  // even though the address space mismatches in the nested pointers.
+  ll = (__local int * __private *)gg;
+
+  gg_f = g;  // expected-error {{assigning to '__global float **' from incompatible type '__global int *'}}
+  gg_f = gg; // expected-error {{assigning to '__global float **' from incompatible type '__global int **'}}
+  gg_f = l;  // expected-error {{assigning to '__global float **' from incompatible type '__local int *'}}
+  gg_f = ll; // expected-error {{assigning to '__global float **' from incompatible type '__local int **'}}
+  gg_f = (__global float * __private *)gg;
+
+  typedef __local int * l_t;
+  typedef __global int * g_t;
+  __private l_t * pl;
+  __private g_t * pg;
+  gg = pl;  // expected-error {{assigning to '__global int **' from incompatible type 'l_t *' (aka '__local int **')}}
+  pl = gg;  // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type '__global int **'}}
+  gg = pg;
+  pg = gg;
+  pg = pl;  // expected-error {{assigning to 'g_t *' (aka '__global int **') from incompatible type 'l_t *' (aka '__local int **')}}
+  pl = pg;  // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type 'g_t *' (aka '__global int **')}}
+
+  ll = (__local int * __private *)(void *)gg;
+  void *vp = ll;
+}
+#endif
+
 __private int func_return_priv(void);       //expected-error {{return value cannot be qualified with address space}}
 __global int func_return_global(void);      //expected-error {{return value cannot be qualified with address space}}
 __local int func_return_local(void);        //expected-error {{return value cannot be qualified with address space}}
diff --git a/test/SemaOpenCL/event_t_overload.cl b/test/SemaOpenCL/event_t_overload.cl
index 0bc67c13c7..eaf05a0489 100644
--- a/test/SemaOpenCL/event_t_overload.cl
+++ b/test/SemaOpenCL/event_t_overload.cl
@@ -7,5 +7,5 @@ void kernel ker(__local char *src1, __local float *src2, __global int *src3) {
   event_t evt;
   foo(evt, src1);
   foo(0, src2);
-  foo(evt, src3); // expected-error {{call to 'foo' is ambiguous}}
+  foo(evt, src3); // expected-error {{no matching function for call to 'foo'}}
 }
diff --git a/test/SemaOpenCL/numbered-address-space.cl b/test/SemaOpenCL/numbered-address-space.cl
index 423d03274c..8a24b3ce97 100644
--- a/test/SemaOpenCL/numbered-address-space.cl
+++ b/test/SemaOpenCL/numbered-address-space.cl
@@ -26,6 +26,6 @@ void test_generic_as_to_builtin_parameter_explicit_cast_numeric(__attribute__((a
 
 void test_generic_as_to_builtin_parameterimplicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) {
   generic int* generic_ptr = as3_ptr;
-  volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__local float *'}}
+  volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-error {{passing '__generic int *' to parameter of type '__local float *' changes address space of pointer}}
 }
 
diff --git a/test/SemaOpenCL/queue_t_overload.cl b/test/SemaOpenCL/queue_t_overload.cl
index 0048895ac8..bc3b241bbe 100644
--- a/test/SemaOpenCL/queue_t_overload.cl
+++ b/test/SemaOpenCL/queue_t_overload.cl
@@ -7,6 +7,6 @@ void kernel ker(__local char *src1, __local float *src2, __global int *src3) {
   queue_t q;
   foo(q, src1);
   foo(0, src2);
-  foo(q, src3); // expected-error {{call to 'foo' is ambiguous}}
+  foo(q, src3); // expected-error {{no matching function for call to 'foo'}}
   foo(1, src3); // expected-error {{no matching function for call to 'foo'}}
 }