[OpenCL] Fix missing address space deduction in template variables

[olestrohm]

Summary

This patch fixes two instances of not deducing address spaces for global template variables.
I've added OpenCL address space deduction to the functions responsible for specializing template variable declarations,
and while this does give the correct address spaces, it is not optimal, as this same address space is deduced at least three times for a single variable.
It would be better if the correct type was correctly passed forward through the phases, which I've explained in the comments.
Since the variable is templated I'm not entirely certain if this is feasible or worthwhile compared to the simple solution of rededucing the address space several times.

I've also included this an example of this in the test for address space deductions

[Anastasia]

So if I understand this correctly when the template is instantiated we don't have QualType with the address space deduced initially during parsing of a template global variable?

[olestrohm]

Essentially, in each pass the type for the variable starts off being extract from the SourceTypeInfo, but since the address space isn't specified, this type is then missing an address space.
Adding the address space deduction then reintroduces the address space to the type.

This is then done at least three times (pre-specialization, in the middle of specialization and when completing specialization), but since this is a template it may not be possible to correctly
deduce the type earlier in case the template type introduces an address space in template specialization?

[rjmccall]

Seems like you shouldn't do it earlier if the type is dependent.

[olestrohm]

Added a guard in deduceOpenCLAddressSpace to stop it from deducing address space for dependent types, as requested.

[olestrohm]

Disregard the last comment, rolling back that diff for now.

[olestrohm]

I have added a check in deduceOpenCLAddressSpace() to check if the type is dependent, and not deduce the address space if it is. This is a big change in the behaviour of address space deduction, meaning that template variables only receive address spaces when they are being specialized. This is good, but caused a lot of changes in the current test file, which has been included in this patch.

There also had to be added a check for dependent types in CheckVariableDeclarationType when checking whether global variables have the correct address space. With this new change dependent types don't necessarily have address spaces yet, and thus then lets it pass until the variable declaration of the specialized template variable is checked later.

[Anastasia]

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

3628

Is this still an issue with your new patch? I presume deducing the address space early would still be better in general but it seems not to play well with the way the template instantiation is implemented currently. Although it might generally be safer not to deduce anything for templates until the final instantiation is done.

[olestrohm]

I've removed the comments calling for a fix because I no longer feel that this approach needs that. Given the code that already exists, and without changing too much of it, adding address space deduction in both cases seems like the right choice.

[Anastasia]

LGTM! Thanks!

clang/lib/Sema/SemaDecl.cpp

bool Sema::inferObjCARCLifetime(ValueDecl *decl) {
   }
 
   return false;
 }
 
 void Sema::deduceOpenCLAddressSpace(ValueDecl *Decl) {
   if (Decl->getType().hasAddressSpace())
     return;
+  if (Decl->getType()->isDependentType())
+    return;
   if (VarDecl *Var = dyn_cast<VarDecl>(Decl)) {
     QualType Type = Var->getType();
     if (Type->isSamplerT() || Type->isVoidType())
       return;
     LangAS ImplAS = LangAS::opencl_private;
     if ((getLangOpts().OpenCLCPlusPlus || getLangOpts().OpenCLVersion >= 200) &&
         Var->hasGlobalStorage())
       ImplAS = LangAS::opencl_global;
if (getLangOpts().OpenCL) {
     // OpenCL C v2.0 s6.5.1 - Variables defined at program scope and static
     // variables inside a function can also be declared in the global
     // address space.
     // C++ for OpenCL inherits rule from OpenCL C v2.0.
     // FIXME: Adding local AS in C++ for OpenCL might make sense.
     if (NewVD->isFileVarDecl() || NewVD->isStaticLocal() ||
         NewVD->hasExternalStorage()) {
       if (!T->isSamplerT() &&
+          !T->isDependentType() &&
           !(T.getAddressSpace() == LangAS::opencl_constant ||
             (T.getAddressSpace() == LangAS::opencl_global &&
              (getLangOpts().OpenCLVersion == 200 ||
               getLangOpts().OpenCLCPlusPlus)))) {
         int Scope = NewVD->isStaticLocal() | NewVD->hasExternalStorage() << 1;
         if (getLangOpts().OpenCLVersion == 200 || getLangOpts().OpenCLCPlusPlus)
           Diag(NewVD->getLocation(), diag::err_opencl_global_invalid_addr_space)
               << Scope << "global or constant";

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Decl *TemplateDeclInstantiator::VisitVarTemplateSpecializationDecl(
   // Build the instantiated declaration
   VarTemplateSpecializationDecl *Var = VarTemplateSpecializationDecl::Create(
       SemaRef.Context, Owner, D->getInnerLocStart(), D->getLocation(),
       VarTemplate, DI->getType(), DI, D->getStorageClass(), Converted);
   Var->setTemplateArgsInfo(TemplateArgsInfo);
   if (InsertPos)
     VarTemplate->AddSpecialization(Var, InsertPos);
 
+  if (SemaRef.getLangOpts().OpenCL)
+    SemaRef.deduceOpenCLAddressSpace(Var);
+
   // Substitute the nested name specifier, if any.
   if (SubstQualifier(D, Var))
     return nullptr;
 
   SemaRef.BuildVariableInstantiation(Var, D, TemplateArgs, LateAttrs, Owner,
                                      StartingScope, false, PrevDecl);
 
   return Var;
VarTemplateSpecializationDecl *Sema::CompleteVarTemplateSpecializationDecl(
   VarSpec->setType(DI->getType());
 
   // Convert the declaration into a definition now.
   VarSpec->setCompleteDefinition();
 
   // Instantiate the initializer.
   InstantiateVariableInitializer(VarSpec, PatternDecl, TemplateArgs);
 
+  if (getLangOpts().OpenCL)
+    deduceOpenCLAddressSpace(VarSpec);
+
   return VarSpec;
 }
 
 /// BuildVariableInstantiation - Used after a new variable has been created.
 /// Sets basic variable data and decides whether to postpone the
 /// variable instantiation.
 void Sema::BuildVariableInstantiation(
     VarDecl *NewVar, VarDecl *OldVar,

clang/test/SemaOpenCLCXX/address-space-deduction.cl

 //RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
 
 //expected-no-diagnostics
 
 //CHECK: |-VarDecl {{.*}} foo 'const __global int'
 constexpr int foo = 0;
 
+//CHECK: |-VarDecl {{.*}} foo1 'T' cinit
+//CHECK: `-VarTemplateSpecializationDecl {{.*}} used foo1 '__global long':'__global long' cinit
+template <typename T>
+T foo1 = 0;
+
 class c {
 public:
   //CHECK: `-VarDecl {{.*}} foo2 'const __global int'
   static constexpr int foo2 = 0;
 };
 
 struct c1 {};
 

 };
 
 
 // Addr spaces for pointee of dependent types are not deduced
 // during parsing but during template instantiation instead.
 
 template <class T>
 struct x1 {
-//CHECK: -CXXMethodDecl {{.*}} operator= 'x1<T> &(const x1<T> &__private){{( __attribute__.*)?}} __generic'
+//CHECK: -CXXMethodDecl {{.*}} operator= 'x1<T> &(const x1<T> &){{( __attribute__.*)?}} __generic'
 //CHECK: -CXXMethodDecl {{.*}} operator= '__generic x1<int> &(const __generic x1<int> &__private){{( __attribute__.*)?}} __generic'
   x1<T>& operator=(const x1<T>& xx) {
     y = xx.y;
     return *this;
   }
   int y;
 };
 
 template <class T>
 struct x2 {
-//CHECK: -CXXMethodDecl {{.*}} foo 'void (x1<T> *__private){{( __attribute__.*)?}} __generic'
+//CHECK: -CXXMethodDecl {{.*}} foo 'void (x1<T> *){{( __attribute__.*)?}} __generic'
 //CHECK: -CXXMethodDecl {{.*}} foo 'void (__generic x1<int> *__private){{( __attribute__.*)?}} __generic'
   void foo(x1<T>* xx) {
     m[0] = *xx;
   }
 //CHECK: -FieldDecl {{.*}}  m 'x1<int> [2]'
   x1<T> m[2];
 };
 
 void bar(__global x1<int> *xx, __global x2<int> *bar) {
   bar->foo(xx);
 }
 
 template <typename T>
 class x3 : public T {
 public:
-  //CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &__private){{( __attribute__.*)?}} __generic'
+  //CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &){{( __attribute__.*)?}} __generic'
   x3(const x3 &t);
 };
-//CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &__private){{( __attribute__.*)?}} __generic'
+//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 *in1, T in2) {
   // This pointer can't be deduced to generic because addr space
   // will be taken from the template argument.
-  //CHECK: `-VarDecl {{.*}} '__private T *__private' cinit
+  //CHECK: `-VarDecl {{.*}} 'T *' cinit
+  //CHECK: `-VarDecl {{.*}} i '__private int *__private' cinit
   T *i = in1;
   T ii;
   __private T *ptr = &ii;
   ptr = &in2;
   return *i;
 }
 
 __kernel void test() {

 __kernel void k() {
   __local float x[2];
   float(*p)[2];
   t1(x);
   t2(&x);
   t3(&x);
   t4(&p);
   t5(&p);
+  long f1 = foo1<long>;
 }