Author: Michael Liao Date: 2019-12-04T09:37:50-05:00 New Revision: fa9dd410a9a9aa65ce6731cbe1ee12c5941eb3e8
URL: https://github.com/llvm/llvm-project/commit/fa9dd410a9a9aa65ce6731cbe1ee12c5941eb3e8 DIFF: https://github.com/llvm/llvm-project/commit/fa9dd410a9a9aa65ce6731cbe1ee12c5941eb3e8.diff LOG: [opencl] Fix address space deduction on array variables. Summary: - The deduced address space needs applying to its element type as well. Reviewers: Anastasia Subscribers: yaxunl, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D70981 Added: Modified: clang/lib/Sema/SemaDecl.cpp clang/test/SemaOpenCL/address-spaces.cl Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index d35037273106..660be458a698 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -6128,7 +6128,26 @@ void Sema::deduceOpenCLAddressSpace(ValueDecl *Decl) { if ((getLangOpts().OpenCLCPlusPlus || getLangOpts().OpenCLVersion >= 200) && Var->hasGlobalStorage()) ImplAS = LangAS::opencl_global; + // If the original type from a decayed type is an array type and that array + // type has no address space yet, deduce it now. + if (auto DT = dyn_cast<DecayedType>(Type)) { + auto OrigTy = DT->getOriginalType(); + if (!OrigTy.getQualifiers().hasAddressSpace() && OrigTy->isArrayType()) { + // Add the address space to the original array type and then propagate + // that to the element type through `getAsArrayType`. + OrigTy = Context.getAddrSpaceQualType(OrigTy, ImplAS); + OrigTy = QualType(Context.getAsArrayType(OrigTy), 0); + // Re-generate the decayed type. + Type = Context.getDecayedType(OrigTy); + } + } Type = Context.getAddrSpaceQualType(Type, ImplAS); + // Apply any qualifiers (including address space) from the array type to + // the element type. This implements C99 6.7.3p8: "If the specification of + // an array type includes any type qualifiers, the element type is so + // qualified, not the array type." + if (Type->isArrayType()) + Type = QualType(Context.getAsArrayType(Type), 0); Decl->setType(Type); } } diff --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl index 55a55dc75050..09a6dd0ba53f 100644 --- a/clang/test/SemaOpenCL/address-spaces.cl +++ b/clang/test/SemaOpenCL/address-spaces.cl @@ -241,3 +241,10 @@ void func_multiple_addr(void) { __private private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}} __private private_int_t *var6;// expected-warning {{multiple identical address spaces specified for type}} } + +void func_with_array_param(const unsigned data[16]); + +__kernel void k() { + unsigned data[16]; + func_with_array_param(data); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
