[Clang] Allow all address spaces to be converted to the default

clang/include/clang/AST/Type.h

@@ -721,7 +721,9 @@ class Qualifiers {
            // to implicitly cast into the default address space.
            (A == LangAS::Default &&
             (B == LangAS::cuda_constant || B == LangAS::cuda_device ||
-             B == LangAS::cuda_shared));
+             B == LangAS::cuda_shared)) ||
+           // Otherwise, assume the default address space is compatible.
+           (A == LangAS::Default);
   }
 
   /// Returns true if the address space in these qualifiers is equal to or

clang/test/Misc/diag-overload-cand-ranges.cpp

@@ -14,8 +14,8 @@ void baz(__attribute__((opencl_private)) int *Data) {}
 void fizz() {
   int *Nop;
   baz(Nop);
-  // CHECK:    error: no matching function
-  // CHECK:    :[[@LINE+1]]:53: note: {{.*}}: 'this' object is in address space '__private'
+
+
   __attribute__((opencl_private)) static auto err = [&]() {};
   err();
 }

clang/test/Sema/address_space_print_macro.c

@@ -19,24 +19,24 @@ char *cmp(AS1 char *x, AS2 char *y) {
 
 __attribute__((address_space(1))) char test_array[10];
 void test3(void) {
-  extern void test3_helper(char *p); // expected-note{{passing argument to parameter 'p' here}}
-  test3_helper(test_array);          // expected-error{{passing '__attribute__((address_space(1))) char *' to parameter of type 'char *' changes address space of pointer}}
+  extern void test3_helper(char *p);
+  test3_helper(test_array);
 }
 
 char AS2 *test4_array;
 void test4(void) {
-  extern void test3_helper(char *p); // expected-note{{passing argument to parameter 'p' here}}
-  test3_helper(test4_array);         // expected-error{{passing 'AS2 char *' to parameter of type 'char *' changes address space of pointer}}
+  extern void test3_helper(char *p);
+  test3_helper(test4_array);
 }
 
 void func(void) {
   char AS1 *x;
   char AS3 *x2;
   AS5 *x3;
   char *y;
-  y = x;  // expected-error{{assigning 'AS1 char *' to 'char *' changes address space of pointer}}
-  y = x2; // expected-error{{assigning 'AS3 char *' to 'char *' changes address space of pointer}}
-  y = x3; // expected-error{{assigning '__attribute__((address_space(5))) char *' to 'char *' changes address space of pointer}}
+  y = x;
+  y = x2;
+  y = x3;
 }
 
 void multiple_attrs(AS_ND int *x) {

clang/test/Sema/address_spaces.c

@@ -35,13 +35,13 @@ struct _st {
 
 __attribute__((address_space(256))) void * * const base = 0;
 void * get_0(void) {
-  return base[0];  // expected-error {{returning '__attribute__((address_space(256))) void *' from a function with result type 'void *' changes address space of pointer}}
+  return base[0];
 }
 
 __attribute__((address_space(1))) char test3_array[10];
 void test3(void) {
-  extern void test3_helper(char *p); // expected-note {{passing argument to parameter 'p' here}}
-  test3_helper(test3_array); // expected-error {{changes address space of pointer}}
+  extern void test3_helper(char *p);
+  test3_helper(test3_array);
 }
 
 typedef void ft(void);

clang/test/Sema/conditional-expr.c

@@ -81,7 +81,7 @@ void foo(void) {
   test0 ? adr2 : adr3; // expected-error{{conditional operator with the second and third operands of type  ('__attribute__((address_space(2))) int *' and '__attribute__((address_space(3))) int *') which are pointers to non-overlapping address spaces}}
 
   // Make sure address-space mask ends up in the result type
-  (test0 ? (test0 ? adr2 : adr2) : nonconst_int); // expected-error{{conditional operator with the second and third operands of type  ('__attribute__((address_space(2))) int *' and 'int *') which are pointers to non-overlapping address spaces}}
+  (void)(test0 ? (test0 ? adr2 : adr2) : nonconst_int);
 }
 
 int Postgresql(void) {

clang/test/Sema/wasm-refs-and-tables.c

@@ -85,9 +85,8 @@ __externref_t func(__externref_t ref) {
   static __externref_t lt2[0];    // expected-error {{WebAssembly table cannot be declared within a function}}
   static __externref_t lt3[0][0]; // expected-error {{multi-dimensional arrays of WebAssembly references are not allowed}}
   static __externref_t(*lt4)[0];  // expected-error {{cannot form a pointer to a WebAssembly table}}
-  // conly-error@+2 {{cannot use WebAssembly table as a function parameter}}
-  // cpp-error@+1 {{no matching function for call to 'illegal_argument_1'}}
-  illegal_argument_1(table);
+
+  illegal_argument_1(table);      // expected-error {{cannot use WebAssembly table as a function parameter}}
   varargs(1, table);              // expected-error {{cannot use WebAssembly table as a function parameter}}
   table == 1;                     // expected-error {{invalid operands to binary expression ('__attribute__((address_space(1))) __externref_t[0]' and 'int')}}
   1 >= table;                     // expected-error {{invalid operands to binary expression ('int' and '__attribute__((address_space(1))) __externref_t[0]')}}

clang/test/SemaCXX/address-space-conversion.cpp

@@ -69,30 +69,30 @@ void test_static_cast(void_ptr vp, void_ptr_1 vp1, void_ptr_2 vp2,
   (void)static_cast<A_ptr_2>(vp2);
 
   // Ill-formed upcasts
-  (void)static_cast<A_ptr>(bp1); // expected-error{{is not allowed}}
-  (void)static_cast<A_ptr>(bp2); // expected-error{{is not allowed}}
+  (void)static_cast<A_ptr>(bp1);
+  (void)static_cast<A_ptr>(bp2);
   (void)static_cast<A_ptr_1>(bp); // expected-error{{is not allowed}}
   (void)static_cast<A_ptr_1>(bp2); // expected-error{{is not allowed}}
   (void)static_cast<A_ptr_2>(bp); // expected-error{{is not allowed}}
   (void)static_cast<A_ptr_2>(bp1); // expected-error{{is not allowed}}
 
   // Ill-formed downcasts
-  (void)static_cast<B_ptr>(ap1); // expected-error{{casts away qualifiers}}
-  (void)static_cast<B_ptr>(ap2); // expected-error{{casts away qualifiers}}
+  (void)static_cast<B_ptr>(ap1);
+  (void)static_cast<B_ptr>(ap2);
   (void)static_cast<B_ptr_1>(ap); // expected-error{{casts away qualifiers}}
   (void)static_cast<B_ptr_1>(ap2); // expected-error{{casts away qualifiers}}
   (void)static_cast<B_ptr_2>(ap); // expected-error{{casts away qualifiers}}
   (void)static_cast<B_ptr_2>(ap1); // expected-error{{casts away qualifiers}}
 
   // Ill-formed cast to/from void
-  (void)static_cast<void_ptr>(ap1); // expected-error{{is not allowed}}
-  (void)static_cast<void_ptr>(ap2); // expected-error{{is not allowed}}
+  (void)static_cast<void_ptr>(ap1);
+  (void)static_cast<void_ptr>(ap2);
   (void)static_cast<void_ptr_1>(ap); // expected-error{{is not allowed}}
   (void)static_cast<void_ptr_1>(ap2); // expected-error{{is not allowed}}
   (void)static_cast<void_ptr_2>(ap); // expected-error{{is not allowed}}
   (void)static_cast<void_ptr_2>(ap1); // expected-error{{is not allowed}}
-  (void)static_cast<A_ptr>(vp1); // expected-error{{casts away qualifiers}}
-  (void)static_cast<A_ptr>(vp2); // expected-error{{casts away qualifiers}}
+  (void)static_cast<A_ptr>(vp1);
+  (void)static_cast<A_ptr>(vp2);
   (void)static_cast<A_ptr_1>(vp); // expected-error{{casts away qualifiers}}
   (void)static_cast<A_ptr_1>(vp2); // expected-error{{casts away qualifiers}}
   (void)static_cast<A_ptr_2>(vp); // expected-error{{casts away qualifiers}}
@@ -112,16 +112,16 @@ void test_dynamic_cast(A_ptr ap, A_ptr_1 ap1, A_ptr_2 ap2,
   (void)dynamic_cast<B_ptr_2>(ap2);
 
   // Ill-formed upcasts
-  (void)dynamic_cast<A_ptr>(bp1); // expected-error{{casts away qualifiers}}
-  (void)dynamic_cast<A_ptr>(bp2); // expected-error{{casts away qualifiers}}
+  (void)dynamic_cast<A_ptr>(bp1);
+  (void)dynamic_cast<A_ptr>(bp2);
   (void)dynamic_cast<A_ptr_1>(bp); // expected-error{{casts away qualifiers}}
   (void)dynamic_cast<A_ptr_1>(bp2); // expected-error{{casts away qualifiers}}
   (void)dynamic_cast<A_ptr_2>(bp); // expected-error{{casts away qualifiers}}
   (void)dynamic_cast<A_ptr_2>(bp1); // expected-error{{casts away qualifiers}}
 
   // Ill-formed downcasts
-  (void)dynamic_cast<B_ptr>(ap1); // expected-error{{casts away qualifiers}}
-  (void)dynamic_cast<B_ptr>(ap2); // expected-error{{casts away qualifiers}}
+  (void)dynamic_cast<B_ptr>(ap1);
+  (void)dynamic_cast<B_ptr>(ap2);
   (void)dynamic_cast<B_ptr_1>(ap); // expected-error{{casts away qualifiers}}
   (void)dynamic_cast<B_ptr_1>(ap2); // expected-error{{casts away qualifiers}}
   (void)dynamic_cast<B_ptr_2>(ap); // expected-error{{casts away qualifiers}}
@@ -133,14 +133,14 @@ void test_reinterpret_cast(void_ptr vp, void_ptr_1 vp1, void_ptr_2 vp2,
                            B_ptr bp, B_ptr_1 bp1, B_ptr_2 bp2,
                            const void __attribute__((address_space(1))) * cvp1) {
   // reinterpret_cast can't be used to cast to a different address space unless they are matching (i.e. overlapping).
-  (void)reinterpret_cast<A_ptr>(ap1); // expected-error{{reinterpret_cast from 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') to 'A_ptr' (aka 'A *') is not allowed}}
-  (void)reinterpret_cast<A_ptr>(ap2); // expected-error{{reinterpret_cast from 'A_ptr_2' (aka '__attribute__((address_space(2))) A *') to 'A_ptr' (aka 'A *') is not allowed}}
+  (void)reinterpret_cast<A_ptr>(ap1);
+  (void)reinterpret_cast<A_ptr>(ap2);
   (void)reinterpret_cast<A_ptr>(bp);
-  (void)reinterpret_cast<A_ptr>(bp1); // expected-error{{reinterpret_cast from 'B_ptr_1' (aka '__attribute__((address_space(1))) B *') to 'A_ptr' (aka 'A *') is not allowed}}
-  (void)reinterpret_cast<A_ptr>(bp2); // expected-error{{reinterpret_cast from 'B_ptr_2' (aka '__attribute__((address_space(2))) B *') to 'A_ptr' (aka 'A *') is not allowed}}
+  (void)reinterpret_cast<A_ptr>(bp1);
+  (void)reinterpret_cast<A_ptr>(bp2);
   (void)reinterpret_cast<A_ptr>(vp);
-  (void)reinterpret_cast<A_ptr>(vp1);   // expected-error{{reinterpret_cast from 'void_ptr_1' (aka '__attribute__((address_space(1))) void *') to 'A_ptr' (aka 'A *') is not allowed}}
-  (void)reinterpret_cast<A_ptr>(vp2);   // expected-error{{reinterpret_cast from 'void_ptr_2' (aka '__attribute__((address_space(2))) void *') to 'A_ptr' (aka 'A *') is not allowed}}
+  (void)reinterpret_cast<A_ptr>(vp1);
+  (void)reinterpret_cast<A_ptr>(vp2);
   (void)reinterpret_cast<A_ptr_1>(ap);  // expected-error{{reinterpret_cast from 'A_ptr' (aka 'A *') to 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') is not allowed}}
   (void)reinterpret_cast<A_ptr_1>(ap2); // expected-error{{reinterpret_cast from 'A_ptr_2' (aka '__attribute__((address_space(2))) A *') to 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') is not allowed}}
   (void)reinterpret_cast<A_ptr_1>(bp);  // expected-error{{reinterpret_cast from 'B_ptr' (aka 'B *') to 'A_ptr_1' (aka '__attribute__((address_space(1))) A *') is not allowed}}
@@ -190,8 +190,6 @@ void test_implicit_conversion(void_ptr vp, void_ptr_1 vp1, void_ptr_2 vp2,
   A_ptr_2 ap_A2 = bp2;
 
   // Ill-formed conversions
-  void_ptr vpB = ap1; // expected-error{{cannot initialize a variable of type}}
   void_ptr_1 vp_1B = ap2; // expected-error{{cannot initialize a variable of type}}
-  A_ptr ap_B = bp1; // expected-error{{cannot initialize a variable of type}}
   A_ptr_1 ap_B1 = bp2; // expected-error{{cannot initialize a variable of type}}
 }

clang/test/SemaCXX/address-space-ctor.cpp

@@ -1,18 +1,12 @@
 // RUN: %clang_cc1 %s -std=c++14 -triple=spir -verify -fsyntax-only
 // RUN: %clang_cc1 %s -std=c++17 -triple=spir -verify -fsyntax-only
 
+// expected-no-diagnostics
+
 struct MyType {
   MyType(int i) : i(i) {}
   int i;
 };
 
-//expected-note@-5{{candidate constructor (the implicit copy constructor) not viable: no known conversion from 'int' to 'const MyType &' for 1st argument}}
-//expected-note@-6{{candidate constructor (the implicit move constructor) not viable: no known conversion from 'int' to 'MyType &&' for 1st argument}}
-//expected-note@-6{{candidate constructor ignored: cannot be used to construct an object in address space '__attribute__((address_space(10)))'}}
-//expected-note@-8{{candidate constructor ignored: cannot be used to construct an object in address space '__attribute__((address_space(10)))'}}
-//expected-note@-9{{candidate constructor ignored: cannot be used to construct an object in address space '__attribute__((address_space(10)))'}}
-//expected-note@-9{{candidate constructor ignored: cannot be used to construct an object in address space '__attribute__((address_space(10)))'}}
-
-// FIXME: We can't implicitly convert between address spaces yet.
-MyType __attribute__((address_space(10))) m1 = 123; //expected-error{{no viable conversion from 'int' to '__attribute__((address_space(10))) MyType'}}
-MyType __attribute__((address_space(10))) m2(123);  //expected-error{{no matching constructor for initialization of '__attribute__((address_space(10))) MyType'}}
+MyType __attribute__((address_space(10))) m1 = 123;
+MyType __attribute__((address_space(10))) m2(123);

clang/test/SemaOpenCL/func.cl

@@ -57,12 +57,6 @@ void bar()
   foo((void*)foo);
 #ifndef FUNCPTREXT
   // expected-error@-2{{taking address of function is not allowed}}
-#else
-  // FIXME: Functions should probably be in the address space defined by the
-  // implementation. It might make sense to put them into the Default address
-  // space that is bind to a physical segment by the target rather than fixing
-  // it to any of the concrete OpenCL address spaces during parsing.
-  // expected-error@-8{{casting 'void (*)(__private void *__private)' to type '__private void *' changes address space}}
 #endif
 
   foo(&foo);

clang/test/SemaOpenCLCXX/address-space-lambda.clcpp

@@ -32,28 +32,13 @@ __kernel void test_qual() {
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () {{.*}}const __generic'
   auto priv2 = []() __generic {};
   priv2();
-  auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}}
-#if defined(_WIN32) && !defined(_WIN64)
-  //expected-note@35{{conversion candidate of type 'void (*)() __attribute__((thiscall))'}}
-#else
-  //expected-note@35{{conversion candidate of type 'void (*)()'}}
-#endif
-  priv3(); //expected-error{{no matching function for call to object of type}}
+  auto priv3 = []() __global {};
+  priv3();
 
-  __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__private'}}
-#if defined(_WIN32) && !defined(_WIN64)
-  //expected-note@43{{conversion candidate of type 'void (*)() __attribute__((thiscall))'}}
-#else
-  //expected-note@43{{conversion candidate of type 'void (*)()'}}
-#endif
-  const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
-  __constant auto const2 = []() __generic{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__generic'}}
-#if defined(_WIN32) && !defined(_WIN64)
-  //expected-note@50{{conversion candidate of type 'void (*)() __attribute__((thiscall))'}}
-#else
-  //expected-note@50{{conversion candidate of type 'void (*)()'}}
-#endif
-  const2(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
+  __constant auto const1 = []() __private{};
+  const1();
+  __constant auto const2 = []() __generic{};
+  const2();
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () {{.*}}const __constant'
   __constant auto const3 = []() __constant{};
   const3();

clang/test/SemaTemplate/address_space-dependent.cpp

@@ -63,15 +63,15 @@ struct fooFunction {
   __attribute__((address_space(I))) void **const base = 0;
 
   void *get_0(void) {
-    return base[0]; // expected-error {{cannot initialize return object of type 'void *' with an lvalue of type '__attribute__((address_space(1))) void *}}
+    return base[0];
   }
 
   __attribute__((address_space(I))) ft qf; // expected-error {{function type may not be qualified with an address space}}
   __attribute__((address_space(I))) char *test3_val;
 
   void test3(void) {
-    extern void test3_helper(char *p); // expected-note {{passing argument to parameter 'p' here}}
-    test3_helper(test3_val);           // expected-error {{cannot initialize a parameter of type 'char *' with an lvalue of type '__attribute__((address_space(1))) char *'}}
+    extern void test3_helper(char *p);
+    test3_helper(test3_val);
   }
 };
 
@@ -109,9 +109,9 @@ int main() {
   cmp<1, 2>(x, y); // expected-note {{in instantiation of function template specialization 'cmp<1, 2>' requested here}}
 
   fooFunction<1> ff;
-  ff.get_0(); // expected-note {{in instantiation of member function 'fooFunction<1>::get_0' requested here}}
+  ff.get_0();
   ff.qf();
-  ff.test3(); // expected-note {{in instantiation of member function 'fooFunction<1>::test3' requested here}}
+  ff.test3();
 
   static_assert(partial_spec_deduce_as<int __attribute__((address_space(3))) *>::value == 3, "address space value has been incorrectly deduced");