Author: Anastasia Stulova Date: 2020-07-13T11:44:38+01:00 New Revision: 6050c156ab4f13a3c54ca6ec297a72ece95966d7
URL: https://github.com/llvm/llvm-project/commit/6050c156ab4f13a3c54ca6ec297a72ece95966d7 DIFF: https://github.com/llvm/llvm-project/commit/6050c156ab4f13a3c54ca6ec297a72ece95966d7.diff LOG: [OpenCL] Defer addr space deduction for dependent type. This patch removes deduction of address spaces in parsing for types that depend on template parameter even if an address space is already known. Deducing it early interferes with template instantiation/specialization logic that uses source address space where address space is not present. Address space deduction for templates is therefore fully moved to the template instantiation/specialization phase. Patch by Ole Strohm (olestrohm)! Tags: #clang Differential Revision: https://reviews.llvm.org/D82781 Added: Modified: clang/lib/Sema/SemaDecl.cpp clang/lib/Sema/SemaTemplateInstantiateDecl.cpp clang/test/SemaOpenCLCXX/address-space-deduction.cl Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index f5e375134c29..3e2b61ae8cdf 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -6290,6 +6290,8 @@ bool Sema::inferObjCARCLifetime(ValueDecl *decl) { 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()) @@ -7859,6 +7861,7 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) { 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 || diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 85adc4ef2dbd..2efb7acb9724 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -3625,6 +3625,9 @@ Decl *TemplateDeclInstantiator::VisitVarTemplateSpecializationDecl( 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; @@ -4895,6 +4898,9 @@ VarTemplateSpecializationDecl *Sema::CompleteVarTemplateSpecializationDecl( // Instantiate the initializer. InstantiateVariableInitializer(VarSpec, PatternDecl, TemplateArgs); + if (getLangOpts().OpenCL) + deduceOpenCLAddressSpace(VarSpec); + return VarSpec; } diff --git a/clang/test/SemaOpenCLCXX/address-space-deduction.cl b/clang/test/SemaOpenCLCXX/address-space-deduction.cl index 6a81a8b2d7c7..ddfdb6da4347 100644 --- a/clang/test/SemaOpenCLCXX/address-space-deduction.cl +++ b/clang/test/SemaOpenCLCXX/address-space-deduction.cl @@ -5,6 +5,11 @@ //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' @@ -30,7 +35,7 @@ struct c2 { 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; @@ -41,7 +46,7 @@ struct x1 { 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; @@ -57,10 +62,10 @@ void bar(__global x1<int> *xx, __global x2<int> *bar) { 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) {} @@ -68,7 +73,8 @@ 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 = ⅈ @@ -111,4 +117,5 @@ __kernel void k() { t3(&x); t4(&p); t5(&p); + long f1 = foo1<long>; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits