hliao created this revision. hliao added a reviewer: Anastasia. Herald added subscribers: cfe-commits, yaxunl. Herald added a project: clang.
- The deduced address space needs applying to its element type as well. Repository: rG LLVM Github Monorepo https://reviews.llvm.org/D70981 Files: clang/lib/Sema/SemaDecl.cpp clang/test/SemaOpenCL/address-spaces.cl Index: clang/test/SemaOpenCL/address-spaces.cl =================================================================== --- clang/test/SemaOpenCL/address-spaces.cl +++ clang/test/SemaOpenCL/address-spaces.cl @@ -241,3 +241,10 @@ __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); +} Index: clang/lib/Sema/SemaDecl.cpp =================================================================== --- clang/lib/Sema/SemaDecl.cpp +++ clang/lib/Sema/SemaDecl.cpp @@ -6128,7 +6128,23 @@ 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()) { + OrigTy = Context.getAddrSpaceQualType(OrigTy, ImplAS); + OrigTy = QualType(Context.getAsArrayType(OrigTy), 0); + 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); } }
Index: clang/test/SemaOpenCL/address-spaces.cl =================================================================== --- clang/test/SemaOpenCL/address-spaces.cl +++ clang/test/SemaOpenCL/address-spaces.cl @@ -241,3 +241,10 @@ __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); +} Index: clang/lib/Sema/SemaDecl.cpp =================================================================== --- clang/lib/Sema/SemaDecl.cpp +++ clang/lib/Sema/SemaDecl.cpp @@ -6128,7 +6128,23 @@ 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()) { + OrigTy = Context.getAddrSpaceQualType(OrigTy, ImplAS); + OrigTy = QualType(Context.getAsArrayType(OrigTy), 0); + 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); } }
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits