https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/112248
Summary: Currently, we want to use the OpenCL attributes to indicate the address space. Languages like SYCL, OpenMP, HIP, and CUDA allow their address space qualifiers to be implicitly converted to generic, as does CL2.0 (except for __constant). We want this behavior when targeting C/C++ directly with the OpenCL attributes or when using CUDA/OpenMP and want to qualify pointers with the types. The current CL1.0 rules are unnecessarily strict when the GPU targets are expected to handle flat pointers. This patch changes the logic to allow any cast if the target is Generic. For OpenCL every global will have `opencl_generic` or `opencl_private` attributes unless it's some kind of function object. I'm not sure if this is the best and most correct solution. If we want to leave OpenCL untouched we could just check the language before checking the rules instead of just the address spaces. Alternatively, we could make an entirely new set of address space attributes that drops the `opencl` name and use those (but then we'd need to duplicate the same sema checking everywhere). Fixes: https://github.com/llvm/llvm-project/issues/112233 >From f4641945a30a71e928141d717646bbc3dd4c922d Mon Sep 17 00:00:00 2001 From: Joseph Huber <hube...@outlook.com> Date: Mon, 14 Oct 2024 14:13:03 -0500 Subject: [PATCH] [Clang] Allow all address spaces to be converted to the default Summary: Currently, we want to use the OpenCL attributes to indicate the address space. Languages like SYCL, OpenMP, HIP, and CUDA allow their address space qualifiers to be implicitly converted to generic, as does CL2.0 (except for __constant). We want this behavior when targeting C/C++ directly with the OpenCL attributes or when using CUDA/OpenMP and want to qualify pointers with the types. The current CL1.0 rules are unnecessarily strict when the GPU targets are expected to handle flat pointers. This patch changes the logic to allow any cast if the target is Generic. For OpenCL every global will have `opencl_generic` or `opencl_private` attributes unless it's some kind of function object. I'm not sure if this is the best and most correct solution. If we want to leave OpenCL untouched we could just check the language before checking the rules instead of just the address spaces. Alternatively, we could make an entirely new set of address space attributes that drops the `opencl` name and use those (but then we'd need to duplicate the same sema checking everywhere). Fixes: https://github.com/llvm/llvm-project/issues/112233 --- clang/include/clang/AST/Type.h | 4 +- clang/test/Misc/diag-overload-cand-ranges.cpp | 4 +- clang/test/Sema/address_space_print_macro.c | 14 +++---- clang/test/Sema/address_spaces.c | 6 +-- clang/test/Sema/conditional-expr.c | 2 +- clang/test/Sema/wasm-refs-and-tables.c | 5 +-- .../test/SemaCXX/address-space-conversion.cpp | 38 +++++++++---------- clang/test/SemaCXX/address-space-ctor.cpp | 14 ++----- clang/test/SemaOpenCL/func.cl | 6 --- .../SemaOpenCLCXX/address-space-lambda.clcpp | 27 +++---------- .../SemaTemplate/address_space-dependent.cpp | 10 ++--- 11 files changed, 51 insertions(+), 79 deletions(-) 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<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}} } 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<int __attribute__((address_space(3))) *>::value == 3, "address space value has been incorrectly deduced"); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
