diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h index 8ff04cf89a6b91..7640f5a31aaf61 100644 --- a/clang/include/clang/AST/Type.h +++ b/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 diff --git a/clang/test/Misc/diag-overload-cand-ranges.cpp b/clang/test/Misc/diag-overload-cand-ranges.cpp index 080ca484d4b746..e98e915e0f47d2 100644 --- a/clang/test/Misc/diag-overload-cand-ranges.cpp +++ b/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(); } diff --git a/clang/test/Sema/address_space_print_macro.c b/clang/test/Sema/address_space_print_macro.c index e01fcf428270bf..77c1ffcbe78b62 100644 --- a/clang/test/Sema/address_space_print_macro.c +++ b/clang/test/Sema/address_space_print_macro.c @@ -19,14 +19,14 @@ 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) { @@ -34,9 +34,9 @@ void func(void) { 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) { diff --git a/clang/test/Sema/address_spaces.c b/clang/test/Sema/address_spaces.c index 7dbeac71195408..3f17a9e6a728bc 100644 --- a/clang/test/Sema/address_spaces.c +++ b/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); diff --git a/clang/test/Sema/conditional-expr.c b/clang/test/Sema/conditional-expr.c index b54b689ec4f055..99ea3a13b1a037 100644 --- a/clang/test/Sema/conditional-expr.c +++ b/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) { diff --git a/clang/test/Sema/wasm-refs-and-tables.c b/clang/test/Sema/wasm-refs-and-tables.c index dd8536c52cd031..bc01a437ce103f 100644 --- a/clang/test/Sema/wasm-refs-and-tables.c +++ b/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]')}} diff --git a/clang/test/SemaCXX/address-space-conversion.cpp b/clang/test/SemaCXX/address-space-conversion.cpp index b1fb69816334df..720713b6b89928 100644 --- a/clang/test/SemaCXX/address-space-conversion.cpp +++ b/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(vp2); // Ill-formed upcasts - (void)static_cast(bp1); // expected-error{{is not allowed}} - (void)static_cast(bp2); // expected-error{{is not allowed}} + (void)static_cast(bp1); + (void)static_cast(bp2); (void)static_cast(bp); // expected-error{{is not allowed}} (void)static_cast(bp2); // expected-error{{is not allowed}} (void)static_cast(bp); // expected-error{{is not allowed}} (void)static_cast(bp1); // expected-error{{is not allowed}} // Ill-formed downcasts - (void)static_cast(ap1); // expected-error{{casts away qualifiers}} - (void)static_cast(ap2); // expected-error{{casts away qualifiers}} + (void)static_cast(ap1); + (void)static_cast(ap2); (void)static_cast(ap); // expected-error{{casts away qualifiers}} (void)static_cast(ap2); // expected-error{{casts away qualifiers}} (void)static_cast(ap); // expected-error{{casts away qualifiers}} (void)static_cast(ap1); // expected-error{{casts away qualifiers}} // Ill-formed cast to/from void - (void)static_cast(ap1); // expected-error{{is not allowed}} - (void)static_cast(ap2); // expected-error{{is not allowed}} + (void)static_cast(ap1); + (void)static_cast(ap2); (void)static_cast(ap); // expected-error{{is not allowed}} (void)static_cast(ap2); // expected-error{{is not allowed}} (void)static_cast(ap); // expected-error{{is not allowed}} (void)static_cast(ap1); // expected-error{{is not allowed}} - (void)static_cast(vp1); // expected-error{{casts away qualifiers}} - (void)static_cast(vp2); // expected-error{{casts away qualifiers}} + (void)static_cast(vp1); + (void)static_cast(vp2); (void)static_cast(vp); // expected-error{{casts away qualifiers}} (void)static_cast(vp2); // expected-error{{casts away qualifiers}} (void)static_cast(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(ap2); // Ill-formed upcasts - (void)dynamic_cast(bp1); // expected-error{{casts away qualifiers}} - (void)dynamic_cast(bp2); // expected-error{{casts away qualifiers}} + (void)dynamic_cast(bp1); + (void)dynamic_cast(bp2); (void)dynamic_cast(bp); // expected-error{{casts away qualifiers}} (void)dynamic_cast(bp2); // expected-error{{casts away qualifiers}} (void)dynamic_cast(bp); // expected-error{{casts away qualifiers}} (void)dynamic_cast(bp1); // expected-error{{casts away qualifiers}} // Ill-formed downcasts - (void)dynamic_cast(ap1); // expected-error{{casts away qualifiers}} - (void)dynamic_cast(ap2); // expected-error{{casts away qualifiers}} + (void)dynamic_cast(ap1); + (void)dynamic_cast(ap2); (void)dynamic_cast(ap); // expected-error{{casts away qualifiers}} (void)dynamic_cast(ap2); // expected-error{{casts away qualifiers}} (void)dynamic_cast(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(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(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(ap1); + (void)reinterpret_cast(ap2); (void)reinterpret_cast(bp); - (void)reinterpret_cast(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(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(bp1); + (void)reinterpret_cast(bp2); (void)reinterpret_cast(vp); - (void)reinterpret_cast(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(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(vp1); + (void)reinterpret_cast(vp2); (void)reinterpret_cast(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(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(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}} } diff --git a/clang/test/SemaCXX/address-space-ctor.cpp b/clang/test/SemaCXX/address-space-ctor.cpp index b872b5a5a84f2d..71ed220d0037cf 100644 --- a/clang/test/SemaCXX/address-space-ctor.cpp +++ b/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); diff --git a/clang/test/SemaOpenCL/func.cl b/clang/test/SemaOpenCL/func.cl index 233e82f244975f..bbe5ba912fd96d 100644 --- a/clang/test/SemaOpenCL/func.cl +++ b/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); diff --git a/clang/test/SemaOpenCLCXX/address-space-lambda.clcpp b/clang/test/SemaOpenCLCXX/address-space-lambda.clcpp index 54c7c88087be8c..3ec91ef2fbfa7b 100644 --- a/clang/test/SemaOpenCLCXX/address-space-lambda.clcpp +++ b/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(); diff --git a/clang/test/SemaTemplate/address_space-dependent.cpp b/clang/test/SemaTemplate/address_space-dependent.cpp index 2ca9b8007ab418..518444d788c052 100644 --- a/clang/test/SemaTemplate/address_space-dependent.cpp +++ b/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::value == 3, "address space value has been incorrectly deduced");