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();
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();
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits