Author: stulova Date: Thu Jul 18 02:12:49 2019 New Revision: 366417 URL: http://llvm.org/viewvc/llvm-project?rev=366417&view=rev Log: [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 Modified: cfe/trunk/lib/Sema/SemaDecl.cpp cfe/trunk/lib/Sema/SemaType.cpp cfe/trunk/lib/Sema/TreeTransform.h cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl cfe/trunk/test/SemaOpenCLCXX/address-space-templates.cl Modified: cfe/trunk/lib/Sema/SemaDecl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=366417&r1=366416&r2=366417&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDecl.cpp (original) +++ cfe/trunk/lib/Sema/SemaDecl.cpp Thu Jul 18 02:12:49 2019 @@ -7491,7 +7491,10 @@ void Sema::CheckVariableDeclarationType( 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(); Modified: cfe/trunk/lib/Sema/SemaType.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=366417&r1=366416&r2=366417&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaType.cpp (original) +++ cfe/trunk/lib/Sema/SemaType.cpp Thu Jul 18 02:12:49 2019 @@ -7419,7 +7419,9 @@ static void deduceOpenCLImplicitAddrSpac (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() || Modified: cfe/trunk/lib/Sema/TreeTransform.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/TreeTransform.h?rev=366417&r1=366416&r2=366417&view=diff ============================================================================== --- cfe/trunk/lib/Sema/TreeTransform.h (original) +++ cfe/trunk/lib/Sema/TreeTransform.h Thu Jul 18 02:12:49 2019 @@ -5392,13 +5392,6 @@ QualType TreeTransform<Derived>::Transfo 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(), Modified: cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl?rev=366417&r1=366416&r2=366417&view=diff ============================================================================== --- cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl (original) +++ cfe/trunk/test/SemaOpenCLCXX/address-space-deduction.cl Thu Jul 18 02:12:49 2019 @@ -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]); +} Modified: cfe/trunk/test/SemaOpenCLCXX/address-space-templates.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCLCXX/address-space-templates.cl?rev=366417&r1=366416&r2=366417&view=diff ============================================================================== --- cfe/trunk/test/SemaOpenCLCXX/address-space-templates.cl (original) +++ cfe/trunk/test/SemaOpenCLCXX/address-space-templates.cl Thu Jul 18 02:12:49 2019 @@ -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}} }; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits