[PATCH] D62584: [OpenCL][PR42033] Deducing addr space with template parameter types

[Anastasia Stulova via Phabricator via cfe-commits]

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();


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D62584.202753.patch
Type: text/x-patch
Size: 2989 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190603/8f40b702/attachment-0001.bin>