[OpenCL][PR42033] Deducing addr space with template parameter types

[Anastasia]

Summary

If dependent types appear in pointers or references we have to allow the addr space deduction because the addr space in template argument will belong to the pointee and not the pointer or reference itself. Hence we end up with un-deduced i.e. Default address space.

We also need to change somke diagnostic to account for addr space logic with templates.

[rjmccall]

I think the right approach here is probably to make sure you're applying deduction during instantiation as well.

[Anastasia]

• Added more test cases
• Improved diagnostics to account for use of templates with addr spaces

[Anastasia]

In D62584#1522438, @rjmccall wrote:

I think the right approach here is probably to make sure you're applying deduction during instantiation as well.

I agree I think we might need to extend the template instantiation logic to account for some corner cases. However, all local variables in OpenCL are to be deduced to __private, therefore would it be better to deduce them already as we parse the template definition instead of doing it multiple times on each instantiation? It doesn't seem the template argument should affect this rule...

[Anastasia]

lib/Sema/TreeTransform.h

5363 ↗

(On Diff #202753)

I am trying to be a bit more helpful here although I am not sure if we should instead require explicit template parameter and fail the template deduction instead.

Basically, do we want the following code to always require specifying template argument explicitly:

template <class T>
T xxx(T *in) {
  T *i = in;
  return *i;
}

__kernel void test() {
  int foo[10];
  xxx<int>(&foo[0]); // if we deduce type from foo, it ends up being qualified by __private that we currently diagnose. However private is default (implicit) address space for return type so technically there is no danger in just allowing xxx(&foo[0])
}

[rjmccall]

lib/Sema/TreeTransform.h

5363 ↗

(On Diff #202753)

Implicitly ignoring all address-space qualifiers on the return type seems like the right thing to do; I don't think it needs to be limited to __private. That's probably also the right thing to do for locals, but I'm less certain about it.

[Anastasia]

lib/Sema/TreeTransform.h

5363 ↗

(On Diff #202753)

Just to clarify by "Implicitly ignoring" you mean ignoring if the templ parameters were deduced?

Although I am a bit concerned about allowing other than __private address spaces in return types as we reject them in return types of functions generally. Would it not be somehow inconsistent?

[Anastasia]

• Removed diagnostic for address space in return type in tree transforms

[Anastasia]

lib/Sema/TreeTransform.h

5363 ↗

(On Diff #202753)

Ok, I have removed the diagnostic completely. At least we don't seem to generate any incorrect IR.

[rjmccall]

lib/Sema/TreeTransform.h

5363 ↗

(On Diff #202753)

They should be diagnosed somehow when written explicitly on a return type, but if you just do that on the parser path you don't have to worry about it during template instantiation. They should probably otherwise be ignored no matter where they came from — if someone typedefs private_int_t to __private int, you should just treat that as int in a return type. Stripping the qualifier from the type is probably the right thing to do so that it doesn't further impact semantic analysis.

I definitely don't think you want a model where the qualifier actually means that the return is somehow done via an object in that address space.

[Anastasia]

lib/Sema/TreeTransform.h

5363 ↗

(On Diff #202753)

They should be diagnosed somehow when written explicitly on a return type, but if you just do that on the parser path you don't have to worry about it during template instantiation.

Ok, this seems to be working currently. The following won't compile:

template <typename T>
struct S {
  __global T f1();     // error: return value cannot be qualified with address space
};

They should probably otherwise be ignored no matter where they came from — if someone typedefs private_int_t to __private int, you should just treat that as int in a return type.

We produce diagnostic for this case currently. I can update this in a separate patch if you think it's more helpful behavior. Although I feel a bit worried about it. However it would align with what we are doing with templates here... so perhaps it's logical change...

Stripping the qualifier from the type is probably the right thing to do so that it doesn't further impact semantic analysis.

I guess you mean stripping quals for the case of typedef or others where we will accept the code and ignore quals on the return type? I have tried to do this just for the template case but there are some assertions firing because we are expecting the original return type to be the same before and after template instantiation. So it feels like I would have to do it for everything together. Maybe it should rather go into separate review?

[rjmccall]

lib/Sema/TreeTransform.h

5363 ↗

(On Diff #202753)

We produce diagnostic for this case currently. I can update this in a separate patch if you think it's more helpful behavior. Although I feel a bit worried about it. However it would align with what we are doing with templates here... so perhaps it's logical change...

Well, it's debatable. C doesn't say that qualifiers are erased in function return types in general. On the other hand, qualifiers are semantically meaningless there, and C has few rules that rely on exact type identity. On reflection, we should probably keep diagnosing that at parse-time, even for typedefs, and just ignore it in templates.

Can you point out which assertion is firing?

[Anastasia]

lib/Sema/TreeTransform.h

5363 ↗

(On Diff #202753)

On reflection, we should probably keep diagnosing that at parse-time, even for typedefs, and just ignore it in templates.

In that case we don't need to fix anything in parsing.

Can you point out which assertion is firing?

I have tried this:

diff --git a/lib/Sema/TreeTransform.h b/lib/Sema/TreeTransform.h
index 878795bb70..db4ae6e56e 100644
--- a/lib/Sema/TreeTransform.h
+++ b/lib/Sema/TreeTransform.h
@@ -5365,6 +5365,11 @@ QualType TreeTransform<Derived>::TransformFunctionProtoType(
     if (ResultType.isNull())
       return QualType();
 
+    if (ResultType.hasQualifiers()) {
+      ResultType = ResultType.getUnqualifiedType();
+      TLB.TypeWasModifiedSafely(ResultType);
+    }
+
     if (getDerived().TransformFunctionTypeParams(
             TL.getBeginLoc(), TL.getParams(),
             TL.getTypePtr()->param_type_begin(),

and I am getting some errors in tests mainly because of the differences in the return type. I.e.

error: 'error' diagnostics expected but not seen: 
  File llvm/tools/clang/test/CXX/dcl.decl/dcl.init/dcl.init.ref/p5-var.cpp Line 126: binding reference of type 'const Base' to value of type 'const volatile Base' drops 'volatile' qualifier
  File llvm/tools/clang/test/CXX/dcl.decl/dcl.init/dcl.init.ref/p5-var.cpp Line 127: binding reference of type 'const Base' to value of type 'const volatile Derived' drops 'volatile' qualifier

Practically we no longer get diagnostics for cases like:

48 template<typename T> T create();

126  const Base &br5 = create<const volatile Base>();
127  const Base &br6 = create<const volatile Derived>();

Also the qualifiers are not printed in the type of return type.

I guess this is expected. I am however struggling with this assert in a couple tests:

llvm/tools/clang/lib/Sema/TypeLocBuilder.cpp:160: clang::TypeLoc clang::TypeLocBuilder::pushImpl(clang::QualType, size_t, unsigned int): Assertion `Capacity - Index == TypeLoc::getFullDataSizeForType(T) && "incorrect data size provided to CreateTypeSourceInfo!"' failed.

I am not sure if I am modifying TypeLocBuilder in a wrong way or perhaps something isn't working well in the Index calculation. If this is the direction we would like to go I can spend more time debugging this and upload the review. Not sure if that might be too disruptive change for the release but I would at least like to see the current patch in Clang9 because it fixes a couple of useful cases. :)

[rjmccall]

Oh, yes, it definitely can't be done to class types. I suppose we should just forget about it.

[Anastasia]

In D62584#1583340, @rjmccall wrote:

Oh, yes, it definitely can't be done to class types. I suppose we should just forget about it.

Ok, regarding address space qualifiers in template instantiation - is it still ok that we ignore them here in this patch?

[rjmccall]

In D62584#1585091, @Anastasia wrote:

In D62584#1583340, @rjmccall wrote:

Oh, yes, it definitely can't be done to class types. I suppose we should just forget about it.

Ok, regarding address space qualifiers in template instantiation - is it still ok that we ignore them here in this patch?

I think ignoring them in templates rather than diagnosing is the right thing to do, since it means that reasonable things, e.g. template functions returning T, won't be erroneous if you instantiate T = __global int just like they wouldn't be with T = const int.

[Anastasia]

In D62584#1586640, @rjmccall wrote:

In D62584#1585091, @Anastasia wrote:

In D62584#1583340, @rjmccall wrote:

Oh, yes, it definitely can't be done to class types. I suppose we should just forget about it.

Ok, regarding address space qualifiers in template instantiation - is it still ok that we ignore them here in this patch?

I think ignoring them in templates rather than diagnosing is the right thing to do, since it means that reasonable things, e.g. template functions returning T, won't be erroneous if you instantiate T = __global int just like they wouldn't be with T = const int.

Ok cool, I have removed the diagnostic btw. :)

[rjmccall]

Minor comment then LGTM

lib/Sema/SemaType.cpp

7418 ↗

(On Diff #207562)

"Except"

[Anastasia]

Fixed typo in the comment

[rjmccall]

Thanks!

cfe/trunk/lib/Sema/SemaDecl.cpp

if (NewVD->isFileVarDecl() || NewVD->isStaticLocal() ||
                   << "constant";
             else
               Diag(NewVD->getLocation(), diag::err_opencl_addrspace_scope)
                   << "local";
             NewVD->setInvalidDecl();
             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();
         return;
       }
     }
   }
 

cfe/trunk/lib/Sema/SemaType.cpp

if ( // Do not deduce addr space for function return type and function type,
       // to define the typedef.
       (D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_typedef &&
        !IsPointee) ||
       // Do not deduce addr space of the void type, e.g. in f(void), otherwise
       // it will fail some sema check.
       (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() ||
       // OpenCL spec v2.0 s6.9.b:
       // The sampler type cannot be used with the __local and __global address
       // space qualifiers.
       // OpenCL spec v2.0 s6.13.14:
       // Samplers can also be declared as global constants in the program

cfe/trunk/lib/Sema/TreeTransform.h

if (getDerived().TransformFunctionTypeParams(
         return QualType();
     }
   }
   else {
     ResultType = getDerived().TransformType(TLB, TL.getReturnLoc());
     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(),
             T->getExtParameterInfosOrNull(),
             ParamTypes, &ParamDecls, ExtParamInfos))
       return QualType();
   }
 

cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl

 class x3 : public T {
 public:
   //CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &){{( __attribute__.*)?}} __generic'
   x3(const x3 &t);
 };
 //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]);
+}

cfe/trunk/test/SemaOpenCLCXX/address-space-templates.cl

 //RUN: %clang_cc1 %s -cl-std=c++ -pedantic -verify -fsyntax-only
 
 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}}
 };
 
 template <typename T>
 T foo1(__global T *i) { // expected-note{{candidate template ignored: substitution failure [with T = __local int]: conflicting address space qualifiers are provided between types '__global T' and '__local int'}}
   return *i;
 }