[OpenCL][PR42033] Fix addr space deduction with template parameters

If dependent types appear in pointers or references we allow addr
space deduction because the addr space in template argument will
belong to the pointee and not the pointer or reference itself.

We also don't diagnose addr space on a function return type after
template instantiation. If any addr space for the return type was
provided on a template parameter this will be diagnosed during the
parsing of template definition.

Differential Revision: https://reviews.llvm.org/D62584

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@366417 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp
index cc91ec5946844f17348cd2e0acaa7c2ac7b11671..72b4f6bbcdaf2fac0d83ec241f19f09808673489 100644 (file)
--- a/lib/Sema/SemaDecl.cpp
+++ b/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/lib/Sema/SemaType.cpp b/lib/Sema/SemaType.cpp
index bb71db7609f55766e6b3c553986435d2658def77..29acf6177eb9c7a455a23d60db75873624bb22f8 100644 (file)
--- a/lib/Sema/SemaType.cpp
+++ b/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/lib/Sema/TreeTransform.h b/lib/Sema/TreeTransform.h
index 3b841ec649a8f54768904016f8c7d6a57ec65f3a..8df18b5c27844bda6a38b770abd1f48939f5c4de 100644 (file)
--- a/lib/Sema/TreeTransform.h
+++ b/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/test/SemaOpenCLCXX/address-space-deduction.cl b/test/SemaOpenCLCXX/address-space-deduction.cl
index 08668951dbca6827933fbc21a19c83f68bede68a..4283ec41e6356581cf14408954778caa861e0688 100644 (file)
--- a/test/SemaOpenCLCXX/address-space-deduction.cl
+++ b/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/test/SemaOpenCLCXX/address-space-templates.cl b/test/SemaOpenCLCXX/address-space-templates.cl
index 48fbdc7642d6fc7ab7ccec06697ed1f673992a57..3fb935766e941436f63b627d2bdbde1fa06f2688 100644 (file)
--- a/test/SemaOpenCLCXX/address-space-templates.cl
+++ b/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}}
 };