diff options
| -rw-r--r-- | clang/lib/Sema/SemaDecl.cpp | 5 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaType.cpp | 4 | ||||
| -rw-r--r-- | clang/lib/Sema/TreeTransform.h | 7 | ||||
| -rw-r--r-- | clang/test/SemaOpenCLCXX/address-space-deduction.cl | 15 | ||||
| -rw-r--r-- | clang/test/SemaOpenCLCXX/address-space-templates.cl | 2 |
5 files changed, 23 insertions, 10 deletions
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index cc91ec59468..72b4f6bbcda 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7491,7 +7491,10 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) { 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(); diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index bb71db7609f..29acf6177eb 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -7419,7 +7419,9 @@ static void deduceOpenCLImplicitAddrSpace(TypeProcessingState &State, (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() || + // Except 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() || diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 3b841ec649a..8df18b5c278 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -5392,13 +5392,6 @@ QualType TreeTransform<Derived>::TransformFunctionProtoType( if (ResultType.isNull()) return QualType(); - // Return type can not be qualified with an address space. - if (ResultType.getAddressSpace() != LangAS::Default) { - SemaRef.Diag(TL.getReturnLoc().getBeginLoc(), - diag::err_attribute_address_function_type); - return QualType(); - } - if (getDerived().TransformFunctionTypeParams( TL.getBeginLoc(), TL.getParams(), TL.getTypePtr()->param_type_begin(), diff --git a/clang/test/SemaOpenCLCXX/address-space-deduction.cl b/clang/test/SemaOpenCLCXX/address-space-deduction.cl index 08668951dbc..4283ec41e63 100644 --- a/clang/test/SemaOpenCLCXX/address-space-deduction.cl +++ b/clang/test/SemaOpenCLCXX/address-space-deduction.cl @@ -63,3 +63,18 @@ public: //CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &){{( __attribute__.*)?}} __generic' template <typename T> x3<T>::x3(const x3<T> &t) {} + +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]); +} diff --git a/clang/test/SemaOpenCLCXX/address-space-templates.cl b/clang/test/SemaOpenCLCXX/address-space-templates.cl index 48fbdc7642d..3fb935766e9 100644 --- a/clang/test/SemaOpenCLCXX/address-space-templates.cl +++ b/clang/test/SemaOpenCLCXX/address-space-templates.cl @@ -3,7 +3,7 @@ template <typename T> struct S { T a; // expected-error{{field may not be qualified with an address space}} - T f1(); // expected-error{{function type may not be qualified with an address space}} + T f1(); // we ignore address space on a return types. void f2(T); // expected-error{{parameter may not be qualified with an address space}} }; |
