[Clang] Allow all address spaces to be converted to the default #112248
Conversation
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: llvm#112233
jhuber6
requested review from
alexey-bataev,
AnastasiaStulova,
arsenm,
Artem-B,
bader,
jayfoad,
jdoerfert,
JonChesterfield,
shiltian and
yxsamliu
llvmbot
added
clang
|
@llvm/pr-subscribers-backend-webassembly @llvm/pr-subscribers-clang Author: Joseph Huber (jhuber6) ChangesSummary: This patch changes the logic to allow any cast if the target is Generic. I'm not sure if this is the best and most correct solution. If we want Fixes: #112233 Full diff: https://github.com/llvm/llvm-project/pull/112248.diff 11 Files Affected:
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");
|
| B == LangAS::cuda_shared)); | ||
| B == LangAS::cuda_shared)) || | ||
| // Otherwise, assume the default address space is compatible. | ||
| (A == LangAS::Default); |
There was a problem hiding this comment.
Default is kind of broken, at least for OpenCL. Should avoid attaching any behavior to it
There was a problem hiding this comment.
Yeah I noted there's actually a case where OpenCL touches this for function pointers. I might just have the OpenCL language call a different version of this. Overall the other targets just allow any AS to decay to default and that's likely the behavior we want when not bound by OpenCL semantics.
There was a problem hiding this comment.
@jhuber6, I agree with you, but you might want to check with @AnastasiaStulova.
We had a couple of lengthy discussions about re-using OpenCL attributes in SYCL mode (here and here), but if I recall it correctly the conclusion was that OpenCL attributes inherit OpenCL semantics in non-OpenCL modes as well. The solution for SYCL mode was adding new attributes (review).
@Naghasan, FYI.
There was a problem hiding this comment.
that's likely the behavior we want when not bound by OpenCL semantics.
I kind of agree, but that's up to the target to tell if this makes sense IMO. My mental model of what Default means in the case of CUDA/HIP/SYCL/OpenMP is "the address of a variable maps to the flat address space of the target".
My understanding here is you want a bypass for target address spaces to decay into default. While I think this is desirable in general, this disregards potential target limitations. The one I have in mind is for SPIR/SPIR-V, the target address space for constant shouldn't be allowed to decay into Default. There is some related discussions here as well.
|
I think the default AS is expected to be losslessly converted to other AS but not the other way around, though I understand it is not clearly stated in LangRef or other places and a lot of code just assumes it can. |
All the other language address spaces (cuda, OpenMP, sycl) support converting to the default address space. OpenCL supports it except for constant. |
I don't think this expectation is thoroughly encoded anywhere, and it's definitely not binding today (there are targets who have a mapping of default that breaks this assumption). Whilst the premise is interesting I'm not sure this patch is doing the right thing. If you are in vanilla C/C++, you wouldn't have anything but default, and hopefully the default is mapped to something sane on the target. Running into an observable situation where this is a concern means that either you've messed around with (non C/C++) attributes, or are linking in something exciting. Neither of which constitutes valid C/C++, and I'm leaning towards saying shouldn't be allowed to silently work - we probably should refrain from infecting C/C++ with explicit address spaces a la OpenCL. What are you actually running into @jhuber6 where this is a concern? More specifically, you are saying "we want to do X, in C++" -> why? |
I'm trying to port the OpenMP device runtime to just use C/C++, we use a lot of This is a case where it's actually harder to do something in C/C++ than OpenCL and I don't want it to be that way. See https://godbolt.org/z/1Gn71qqPT for an example of what I'm talking about. TL;DR, every other language allows this and I want to be able to use classes or LDS variables without spamming |
Well it's harder to do OpenCL in C++ than doing OpenCL in OpenCL is not entirely surprising, is it? I understand you do not want to do it that way, however this allows for what is essentially non-standard C++ to sneakily start "working" as if it were standard C++. Every other language in this case == every other GPU offload language modelled after CUDA/OpenCL, but I don't know if that should mean we can simply drop something in an otherwise standard C++ compilation. This is important because these are typefully significant target specific attributes, with consequential semantics in the AST (for example they contribute to overload resolution). It's also plausibly important for quirky cases such as bit-patterns |
I'm thinking we could have a language option like I do understand that there's some weird semantics here, and it does change some surprisingly unrelated tests so I don't think the patch as-is will be sound. Hiding it behind an option makes sense so I can do that. |
I think we should avoid modality / language dependence. I thought the problem you were dealing with was the cuda_ attributes not behaving consistently in C++ mode? |
But this is language dependent, it's not some inconsequential BE/LLVM aspect being bubbled up. Explicit, semantically meaningful ASes impact the language. They should not "just work" if you're compiling for vanilla C/C++ (so not CUDA C++, or SYCL, or OpenCL C/C++ etc.) because neither of those defines such a notion. Furthermore, C++ expects that it's possible for implementation specific attributes to be removable/ignorable without there being an impact on program semantics / behaviour, which doesn't necessarily have to be the case here. Finally, it's not clear that it's even right to have this implicit conversion to default from a linguistic standpoint in C/C++, even if it is convenient. It would require more thought and consideration for how it composes with the rest of the language. So given these, it actually makes sense to have it be opt-in, IMHO, to clearly signal that one is opting into non-standard behaviour. |
a6a237f added an explicit addrspace_cast operator. It would be better to use something like this rather than just freely enabling any implicit cast |
That's only enabled in OpenCLC++ I think? We could probably make a patch to enable it in C++. (Guessing we'd need to call it I tried adding an option to relax the rules but it's much more difficult than I anticipated because it requires forwarding the language options to every 100 or so places that Sema wants to check if the types are legal. OpenCL3.0 lets you enable it for convenience so I'd like that kind of behavior (We probably also do want addrspace cast in general). But it seems like it might be really, really annoying to enable that without becoming beholden to OpenCL rules. The only way I could thing to change these rules optionally would be to introduce a ton of new addrspace rules which are basically just OpenCL but not. |
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_genericoropencl_privateattributes 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
openclname and use those (but then we'd need to duplicate thesame sema checking everywhere).
Fixes: #112233