Anastasia updated this revision to Diff 202753.
Anastasia retitled this revision from "[OpenCL][PR42033] Deducing addr space of
pointer/reference with template parameter types" to "[OpenCL][PR42033] Deducing
addr space with template parameter types".
Anastasia edited the summary of this revision.
Anastasia added a comment.
- Added more test cases
- Improved diagnostics to account for use of templates with addr spaces
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D62584/new/
https://reviews.llvm.org/D62584
Files:
lib/Sema/SemaDecl.cpp
lib/Sema/SemaType.cpp
lib/Sema/TreeTransform.h
test/SemaOpenCLCXX/address-space-deduction.cl
Index: test/SemaOpenCLCXX/address-space-deduction.cl
===================================================================
--- test/SemaOpenCLCXX/address-space-deduction.cl
+++ test/SemaOpenCLCXX/address-space-deduction.cl
@@ -10,3 +10,18 @@
//CHECK: `-VarDecl {{.*}} foo2 'const __global int' static constexpr cinit
static constexpr int foo2 = 0;
};
+
+template <class T>
+T xxx(T *in) {
+ // This pointer can't be deduced to generic because addr space
+ // will be taken from the template argument.
+ //CHECK: `-VarDecl {{.*}} i 'T *' cinit
+ T *i = in;
+ T ii;
+ return *i;
+}
+
+__kernel void test() {
+ int foo[10];
+ xxx(&foo[0]);
+}
Index: lib/Sema/TreeTransform.h
===================================================================
--- lib/Sema/TreeTransform.h
+++ lib/Sema/TreeTransform.h
@@ -5355,8 +5355,12 @@
if (ResultType.isNull())
return QualType();
- // Return type can not be qualified with an address space.
- if (ResultType.getAddressSpace() != LangAS::Default) {
+ // Return type can not be qualified with an address space apart from when
it
+ // comes from a template argument in which case we can accept OpenCL
private
+ // address space because similar to Default it is used for an automatic
+ // storage.
+ if (ResultType.getAddressSpace() != LangAS::Default &&
+ (ResultType.getAddressSpace() != LangAS::opencl_private)) {
SemaRef.Diag(TL.getReturnLoc().getBeginLoc(),
diag::err_attribute_address_function_type);
return QualType();
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -7360,7 +7360,9 @@
(T->isVoidType() && !IsPointee) ||
// Do not deduce addr spaces for dependent types because they might end
// up instantiating to a type with an explicit address space qualifier.
- T->isDependentType() ||
+ // Expect for pointer or reference types because the addr space in
+ // template argument can only belong to a pointee.
+ (T->isDependentType() && !T->isPointerType() && !T->isReferenceType()) ||
// Do not deduce addr space of decltype because it will be taken from
// its argument.
T->isDecltypeType() ||
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -7476,7 +7476,10 @@
return;
}
}
- } else if (T.getAddressSpace() != LangAS::opencl_private) {
+ } else if (T.getAddressSpace() != LangAS::opencl_private &&
+ // If we are parsing a template we didn't deduce an addr
+ // space yet.
+ T.getAddressSpace() != LangAS::Default) {
// Do not allow other address spaces on automatic variable.
Diag(NewVD->getLocation(), diag::err_as_qualified_auto_decl) << 1;
NewVD->setInvalidDecl();
Index: test/SemaOpenCLCXX/address-space-deduction.cl
===================================================================
--- test/SemaOpenCLCXX/address-space-deduction.cl
+++ test/SemaOpenCLCXX/address-space-deduction.cl
@@ -10,3 +10,18 @@
//CHECK: `-VarDecl {{.*}} foo2 'const __global int' static constexpr cinit
static constexpr int foo2 = 0;
};
+
+template <class T>
+T xxx(T *in) {
+ // This pointer can't be deduced to generic because addr space
+ // will be taken from the template argument.
+ //CHECK: `-VarDecl {{.*}} i 'T *' cinit
+ T *i = in;
+ T ii;
+ return *i;
+}
+
+__kernel void test() {
+ int foo[10];
+ xxx(&foo[0]);
+}
Index: lib/Sema/TreeTransform.h
===================================================================
--- lib/Sema/TreeTransform.h
+++ lib/Sema/TreeTransform.h
@@ -5355,8 +5355,12 @@
if (ResultType.isNull())
return QualType();
- // Return type can not be qualified with an address space.
- if (ResultType.getAddressSpace() != LangAS::Default) {
+ // Return type can not be qualified with an address space apart from when it
+ // comes from a template argument in which case we can accept OpenCL private
+ // address space because similar to Default it is used for an automatic
+ // storage.
+ if (ResultType.getAddressSpace() != LangAS::Default &&
+ (ResultType.getAddressSpace() != LangAS::opencl_private)) {
SemaRef.Diag(TL.getReturnLoc().getBeginLoc(),
diag::err_attribute_address_function_type);
return QualType();
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -7360,7 +7360,9 @@
(T->isVoidType() && !IsPointee) ||
// Do not deduce addr spaces for dependent types because they might end
// up instantiating to a type with an explicit address space qualifier.
- T->isDependentType() ||
+ // Expect for pointer or reference types because the addr space in
+ // template argument can only belong to a pointee.
+ (T->isDependentType() && !T->isPointerType() && !T->isReferenceType()) ||
// Do not deduce addr space of decltype because it will be taken from
// its argument.
T->isDecltypeType() ||
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -7476,7 +7476,10 @@
return;
}
}
- } else if (T.getAddressSpace() != LangAS::opencl_private) {
+ } else if (T.getAddressSpace() != LangAS::opencl_private &&
+ // If we are parsing a template we didn't deduce an addr
+ // space yet.
+ T.getAddressSpace() != LangAS::Default) {
// Do not allow other address spaces on automatic variable.
Diag(NewVD->getLocation(), diag::err_as_qualified_auto_decl) << 1;
NewVD->setInvalidDecl();
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
