From 43c9dd09a6fc8eef1914a8fda5488bf9482f118d Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Sun, 9 Oct 2022 16:56:20 -0700 Subject: [PATCH 1/4] [SYCL] Do not decompose arrays with pointers Arrays are no longer decomposed into its elements. Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 274 ++++++++---- .../no_opaque_pointers-in-structs.cpp | 13 +- .../test/CodeGenSYCL/pointers-in-structs.cpp | 14 +- .../test/CodeGenSYCL/pointers-int-header.cpp | 30 +- clang/test/SemaSYCL/array-kernel-param.cpp | 415 ++++++++---------- .../SemaSYCL/built-in-type-kernel-arg.cpp | 93 ++-- clang/test/SemaSYCL/decomposition.cpp | 8 +- 7 files changed, 424 insertions(+), 423 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ec2803b37159b..5c4c0425af06e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1237,10 +1237,6 @@ class KernelObjVisitor { (Handlers.leaveArray(Field, ArrayTy, ET), 0)...}; } - template - void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType ArrayTy, HandlerTys &... Handlers); - template void visitField(const CXXRecordDecl *Owner, FieldDecl *Field, QualType FieldTy, HandlerTys &... Handlers) { @@ -1286,6 +1282,11 @@ class KernelObjVisitor { for (const auto Field : Owner->fields()) visitField(Owner, Field, Field->getType(), Handlers...); } + + template + void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, + QualType ArrayTy, HandlerTys &...Handlers); + #undef KF_FOR_EACH }; @@ -1497,14 +1498,15 @@ void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, visitComplexArray(Owner, Field, ArrayTy, Handlers...); } else if (AnyTrue:: Value) { - assert(!Field->hasAttr() && - "Arrays should trigger decomposition"); - // We are currently in PointerHandler visitor, which implies this is a - // 'simple' array i.e. one that does not include special types or pointers. - // Array of pointers/ array of type containing pointers will be handled in - // a follow-up PR. Currently, they continue to trigger decomposition, and - // will be handled in 'if' statement above. - visitSimpleArray(Owner, Field, ArrayTy, Handlers...); + // We are currently in PointerHandler visitor + if (Field->hasAttr()) { + // This is an array of pointers, or an array of a type containing + // pointers. + visitComplexArray(Owner, Field, ArrayTy, Handlers...); + } else { + // This is an array which does not contain pointers. + visitSimpleArray(Owner, Field, ArrayTy, Handlers...); + } } else { if (!AllTrue::Value) visitSimpleArray( @@ -1841,6 +1843,11 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } bool leaveArray(FieldDecl *FD, QualType ArrayTy, QualType ElementTy) final { + // If an array needs to be decomposed, it is marked with + // SYCLRequiresDecompositionAttr. Else if the array is an array of pointers + // or an array of structs containing pointers, it is marked with + // SYCLGenerateNewTypeAttr. An array will never be marked with both + // attributes. if (CollectionStack.pop_back_val()) { // Cannot assert, since in MD arrays we'll end up marking them multiple // times. @@ -1850,23 +1857,41 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { - // FIXME: Array of pointers/ array of type containing pointers - // will be handled in a follow up PR. Currently, they continue - // to trigger decomposition. - if (!FD->hasAttr()) - FD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( - SemaRef.getASTContext())); - CollectionStack.back() = true; + if (!FD->hasAttr()) + FD->addAttr( + SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + PointerStack.back() = true; } return true; } }; +static QualType ModifyAddressSpace(Sema &SemaRef, QualType Ty) { + // USM allows to use raw pointers instead of buffers/accessors, but these + // pointers point to the specially allocated memory. For pointer fields, + // except for function pointer fields, we add a kernel argument with the + // same type as field but global address space, because OpenCL requires it. + // Function pointers should have program address space. This is set in + // CodeGen. + QualType PointeeTy = Ty->getPointeeType(); + Qualifiers Quals = PointeeTy.getQualifiers(); + auto AS = Quals.getAddressSpace(); + // Leave global_device and global_host address spaces as is to help FPGA + // device in memory allocations + if (!PointeeTy->isFunctionType() && AS != LangAS::sycl_global_device && + AS != LangAS::sycl_global_host) + Quals.setAddressSpace(LangAS::sycl_global); + PointeeTy = SemaRef.getASTContext().getQualifiedType( + PointeeTy.getUnqualifiedType(), Quals); + return SemaRef.getASTContext().getPointerType(PointeeTy); +} + // This visitor is used to traverse a non-decomposed record/array to // generate a new type corresponding to this record/array. class SyclKernelPointerHandler : public SyclKernelFieldHandler { llvm::SmallVector ModifiedRecords; SmallVector ModifiedBases; + SmallVector ModifiedArrayElementsOrArray; IdentifierInfo *getModifiedName(IdentifierInfo *Id) { std::string Name = @@ -1930,24 +1955,35 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return ModifiedRD; } + bool isArrayElement(const FieldDecl *FD, QualType Ty) const { + return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); + } + public: static constexpr const bool VisitInsideSimpleContainersWithPointer = true; + static constexpr const bool VisitNthArrayElement = false; SyclKernelPointerHandler(Sema &S, const CXXRecordDecl *RD) : SyclKernelFieldHandler(S) { createNewType(RD); } + SyclKernelPointerHandler(Sema &S) : SyclKernelFieldHandler(S) {} + bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { createNewType(Ty->getAsCXXRecordDecl()); return true; } bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { + CXXRecordDecl *ModifiedRD = getGeneratedNewRecord(Ty->getAsCXXRecordDecl()); - // Add this record as a field of it's parent record. - if (!ModifiedRecords.empty()) + if (!isArrayElement(FD, Ty)) addField(FD, QualType(ModifiedRD->getTypeForDecl(), 0)); + else + ModifiedArrayElementsOrArray.push_back( + QualType(ModifiedRD->getTypeForDecl(), 0)); + return true; } @@ -1966,22 +2002,39 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return true; } - bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { - QualType PointeeTy = FieldTy->getPointeeType(); - Qualifiers Quals = PointeeTy.getQualifiers(); - LangAS AS = Quals.getAddressSpace(); - // Leave global_device and global_host address spaces as is to help FPGA - // device in memory allocations. - if (!PointeeTy->isFunctionType() && AS != LangAS::sycl_global_device && - AS != LangAS::sycl_global_host) - Quals.setAddressSpace(LangAS::sycl_global); - PointeeTy = SemaRef.getASTContext().getQualifiedType( - PointeeTy.getUnqualifiedType(), Quals); - QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy); - addField(FD, ModTy); + bool leaveArray(FieldDecl *FD, QualType ArrayTy, QualType ET) final { + QualType ModifiedArrayElement = ModifiedArrayElementsOrArray.pop_back_val(); + + const ConstantArrayType *CAT = + SemaRef.getASTContext().getAsConstantArrayType(ArrayTy); + QualType ModifiedArray = SemaRef.getASTContext().getConstantArrayType( + ModifiedArrayElement, CAT->getSize(), + const_cast(CAT->getSizeExpr()), CAT->getSizeModifier(), + CAT->getIndexTypeCVRQualifiers()); + + if (ModifiedRecords.empty()) { + // This is a top-level kernel argument. + ModifiedArrayElementsOrArray.push_back(ModifiedArray); + } else if (!isArrayElement(FD, ArrayTy)) { + // Add this array field as a field of it's parent record. + addField(FD, ModifiedArray); + } else { + // Multi-dimensional array element + ModifiedArrayElementsOrArray.push_back(ModifiedArray); + } + return true; + } + + bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { + QualType ModifiedPointerType = ModifyAddressSpace(SemaRef, FieldTy); + if (!isArrayElement(FD, FieldTy)) + addField(FD, ModifiedPointerType); + else + ModifiedArrayElementsOrArray.push_back(ModifiedPointerType); // We do not need to wrap pointers since this is a pointer inside // non-decomposed struct. + return true; } bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { @@ -2010,10 +2063,6 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return true; } - // FIXME: Array of pointers/ array of types containing pointers - // will be handled in a follow-up PR. Currently they continue to - // trigger decomposition. - public: QualType getNewType() { CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); @@ -2024,6 +2073,9 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return QualType(ModifiedRD->getTypeForDecl(), 0); } + QualType getNewArrayType() { + return ModifiedArrayElementsOrArray.pop_back_val(); + } }; // A type to Create and own the FunctionDecl for the kernel. @@ -2218,8 +2270,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // recursively and modifies the address spaces of any pointer // found as required, thereby generating a new record with all // pointers in 'right' address space. PointerHandler.getNewType() - // returns this generated type, which is then added an openCL - // kernel argument. + // returns this generated type. QualType GenerateNewType(const CXXRecordDecl *RD) { SyclKernelPointerHandler PointerHandler(SemaRef, RD); KernelObjVisitor Visitor{SemaRef}; @@ -2228,6 +2279,18 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return PointerHandler.getNewType(); } + // If the array has been marked with SYCLGenerateNewTypeAttr, + // it implies that this is an array of pointers, or an array + // of a type which contains pointers. This function generates + // a new array with all pointers in the required address space. + QualType GenerateNewArrayType(FieldDecl *FD, QualType FieldTy) { + const CXXRecordDecl *Owner = dyn_cast(FD->getParent()); + SyclKernelPointerHandler PointerHandler(SemaRef); + KernelObjVisitor Visitor{SemaRef}; + Visitor.visitArray(Owner, FD, FieldTy, PointerHandler); + return PointerHandler.getNewArrayType(); + } + public: static constexpr const bool VisitInsideSimpleContainers = false; SyclKernelDeclCreator(Sema &S, SourceLocation Loc, bool IsInline, @@ -2335,23 +2398,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { }; bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { - // USM allows to use raw pointers instead of buffers/accessors, but these - // pointers point to the specially allocated memory. For pointer fields, - // except for function pointer fields, we add a kernel argument with the - // same type as field but global address space, because OpenCL requires it. - // Function pointers should have program address space. This is set in - // CodeGen. - QualType PointeeTy = FieldTy->getPointeeType(); - Qualifiers Quals = PointeeTy.getQualifiers(); - LangAS AS = Quals.getAddressSpace(); - // Leave global_device and global_host address spaces as is to help FPGA - // device in memory allocations. - if (!PointeeTy->isFunctionType() && AS != LangAS::sycl_global_device && - AS != LangAS::sycl_global_host) - Quals.setAddressSpace(LangAS::sycl_global); - PointeeTy = SemaRef.getASTContext().getQualifiedType( - PointeeTy.getUnqualifiedType(), Quals); - QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy); + QualType ModTy = ModifyAddressSpace(SemaRef, FieldTy); // When the kernel is generated, struct type kernel arguments are // decomposed; i.e. the parameters of the kernel are the fields of the // struct, and not the struct itself. This causes an error in the backend @@ -2368,11 +2415,15 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { } bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { - // Arrays are always wrapped in a struct since they cannot be passed - // directly. - RecordDecl *WrappedArray = wrapField(FD, FieldTy); - QualType ModTy = SemaRef.getASTContext().getRecordType(WrappedArray); - addParam(FD, ModTy); + QualType ArrayTy = FieldTy; + + // This is an array of pointers or an array of a type with pointer. + if (FD->hasAttr()) + ArrayTy = GenerateNewArrayType(FD, FieldTy); + + // Arrays are wrapped in a struct since they cannot be passed directly. + RecordDecl *WrappedArray = wrapField(FD, ArrayTy); + addParam(FD, SemaRef.getASTContext().getRecordType(WrappedArray)); return true; } @@ -2690,6 +2741,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { InitializedEntity VarEntity; const CXXRecordDecl *KernelObj; llvm::SmallVector MemberExprBases; + llvm::SmallVector ArrayParamBases; FunctionDecl *KernelCallerFunc; SourceLocation KernelCallerSrcLoc; // KernelCallerFunc source location. // Contains a count of how many containers we're in. This is used by the @@ -3148,6 +3200,70 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { KernelHandlerClone->setInitStyle(VarDecl::CallInit); } + Expr *createArraySubscriptExpr(uint64_t Index, Expr *ArrayRef) { + QualType SizeT = SemaRef.getASTContext().getSizeType(); + + llvm::APInt IndexVal{ + static_cast(SemaRef.getASTContext().getTypeSize(SizeT)), + Index, SizeT->isSignedIntegerType()}; + + auto IndexLiteral = IntegerLiteral::Create( + SemaRef.getASTContext(), IndexVal, SizeT, KernelCallerSrcLoc); + + ExprResult IndexExpr = SemaRef.CreateBuiltinArraySubscriptExpr( + ArrayRef, KernelCallerSrcLoc, IndexLiteral, KernelCallerSrcLoc); + + assert(!IndexExpr.isInvalid()); + return IndexExpr.get(); + } + + void addSimpleArrayInit(FieldDecl *FD, QualType FieldTy) { + Expr *ArrayRef = createSimpleArrayParamReferenceExpr(FieldTy); + InitializationKind InitKind = InitializationKind::CreateDirect({}, {}, {}); + + InitializedEntity Entity = + InitializedEntity::InitializeMember(FD, &VarEntity, /*Implicit*/ true); + + addFieldInit(FD, FieldTy, ArrayRef, InitKind, Entity); + } + + void addArrayElementInit(FieldDecl *FD, QualType T) { + Expr *RCE = createReinterpretCastExpr( + createGetAddressOf(ArrayParamBases.pop_back_val()), + SemaRef.Context.getPointerType(T)); + Expr *Initializer = createDerefOp(RCE); + addFieldInit(FD, T, Initializer); + } + + void createArrayInit(FieldDecl *FD, QualType T) { + const ConstantArrayType *CAT = + SemaRef.getASTContext().getAsConstantArrayType(T); + + if (!CAT) { + addArrayElementInit(FD, T); + return; + } + + QualType ET = CAT->getElementType(); + uint64_t ElemCount = CAT->getSize().getZExtValue(); + enterArray(FD, T, ET); + + for (uint64_t Index = 0; Index < ElemCount; ++Index) { + ArrayInfos.back().second = Index; + Expr *ArraySubscriptExpr = + createArraySubscriptExpr(Index, ArrayParamBases.back()); + ArrayParamBases.push_back(ArraySubscriptExpr); + createArrayInit(FD, ET); + } + + leaveArray(FD, T, ET); + } + + void handleGeneratedArrayType(FieldDecl *FD, QualType FieldTy) { + ArrayParamBases.push_back(createSimpleArrayParamReferenceExpr(FieldTy)); + createArrayInit(FD, FieldTy); + } + public: static constexpr const bool VisitInsideSimpleContainers = false; SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, @@ -3198,13 +3314,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { - Expr *ArrayRef = createSimpleArrayParamReferenceExpr(FieldTy); - InitializationKind InitKind = InitializationKind::CreateDirect({}, {}, {}); - - InitializedEntity Entity = - InitializedEntity::InitializeMember(FD, &VarEntity, /*Implicit*/ true); - - addFieldInit(FD, FieldTy, ArrayRef, InitKind, Entity); + if (FD->hasAttr()) + handleGeneratedArrayType(FD, FieldTy); + else + addSimpleArrayInit(FD, FieldTy); return true; } @@ -3326,21 +3439,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { if (Index != 0) MemberExprBases.pop_back(); - QualType SizeT = SemaRef.getASTContext().getSizeType(); - - llvm::APInt IndexVal{ - static_cast(SemaRef.getASTContext().getTypeSize(SizeT)), - Index, SizeT->isSignedIntegerType()}; - - auto IndexLiteral = IntegerLiteral::Create( - SemaRef.getASTContext(), IndexVal, SizeT, KernelCallerSrcLoc); - - ExprResult IndexExpr = SemaRef.CreateBuiltinArraySubscriptExpr( - MemberExprBases.back(), KernelCallerSrcLoc, IndexLiteral, - KernelCallerSrcLoc); - - assert(!IndexExpr.isInvalid()); - MemberExprBases.push_back(IndexExpr.get()); + MemberExprBases.push_back( + createArraySubscriptExpr(Index, MemberExprBases.back())); return true; } @@ -3349,8 +3449,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { CollectionInitExprs.pop_back(); ArrayInfos.pop_back(); - // Remove the IndexExpr. - MemberExprBases.pop_back(); + if (!FD->hasAttr()) + MemberExprBases.pop_back(); + else + ArrayParamBases.pop_back(); // Remove the field access expr as well. removeFieldMemberExpr(FD, ArrayType); diff --git a/clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp b/clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp index 0ba2092b68232..cea40fcfa436c 100644 --- a/clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp +++ b/clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp @@ -33,16 +33,9 @@ int main() { return 0; } -// CHECK: %[[WRAPPER_F1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* } -// CHECK: %[[WRAPPER_F2:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* } -// CHECK: %[[GENERATED_A:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* } -// CHECK: %[[WRAPPER_F4_1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* } -// CHECK: %[[WRAPPER_F4_2:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* } +// CHECK: %[[GENERATED_B:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)*, float addrspace(1)*, %[[GENERATED_A:[a-zA-Z0-9_.]+]], [2 x i32 addrspace(1)*] } +// CHECK: %[[GENERATED_A]] = type { float addrspace(1)* } // CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* } // CHECK: define {{.*}}spir_kernel void @{{.*}}structs -// CHECK-SAME: %[[WRAPPER_F1]]* noundef byval(%[[WRAPPER_F1]]) align 8 %_arg_F1, -// CHECK-SAME: %[[WRAPPER_F2]]* noundef byval(%[[WRAPPER_F2]]) align 8 %_arg_F2, -// CHECK-SAME: %[[GENERATED_A]]* noundef byval(%[[GENERATED_A]]) align 8 %_arg_F3, -// CHECK-SAME: %[[WRAPPER_F4_1]]* noundef byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4 -// CHECK-SAME: %[[WRAPPER_F4_2]]* noundef byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41 +// CHECK-SAME: %[[GENERATED_B]]* noundef byval(%[[GENERATED_B]]) align 8 %_arg_Obj // CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(%[[WRAPPER_LAMBDA_PTR]]* noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Lambda) diff --git a/clang/test/CodeGenSYCL/pointers-in-structs.cpp b/clang/test/CodeGenSYCL/pointers-in-structs.cpp index baedde086e81b..fdd95cf100e0a 100644 --- a/clang/test/CodeGenSYCL/pointers-in-structs.cpp +++ b/clang/test/CodeGenSYCL/pointers-in-structs.cpp @@ -33,16 +33,10 @@ int main() { return 0; } -// CHECK: %[[WRAPPER_F1:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } -// CHECK: %[[WRAPPER_F2:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } -// CHECK: %[[GENERATED_A:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } -// CHECK: %[[WRAPPER_F4_1:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } -// CHECK: %[[WRAPPER_F4_2:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } +// CHECK: %[[GENERATED_B:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1), ptr addrspace(1), %[[GENERATED_A:[a-zA-Z0-9_.]+]], [2 x ptr addrspace(1)] } +// CHECK: [[GENERATED_A]] = type { ptr addrspace(1) } // CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } + // CHECK: define {{.*}}spir_kernel void @{{.*}}structs -// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F1]]) align 8 %_arg_F1, -// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F2]]) align 8 %_arg_F2, -// CHECK-SAME: ptr noundef byval(%[[GENERATED_A]]) align 8 %_arg_F3, -// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4 -// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41 +// CHECK-SAME: ptr noundef byval(%[[GENERATED_B]]) align 8 %_arg_Obj // CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(ptr noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Lambda) diff --git a/clang/test/CodeGenSYCL/pointers-int-header.cpp b/clang/test/CodeGenSYCL/pointers-int-header.cpp index 92d35bb334ee4..9874e0b0b5dea 100644 --- a/clang/test/CodeGenSYCL/pointers-int-header.cpp +++ b/clang/test/CodeGenSYCL/pointers-int-header.cpp @@ -7,7 +7,15 @@ #include "Inputs/sycl.hpp" -struct struct_with_pointer { +struct decomposed_struct_with_pointer { + int data_in_struct; + int *ptr_in_struct; + int *ptr_array_in_struct1[2]; + int *ptr_array_in_struct2[2][3]; + sycl::accessor acc; +}; + +struct non_decomposed_struct_with_pointer { int data_in_struct; int *ptr_in_struct; int *ptr_array_in_struct1[2]; @@ -16,12 +24,14 @@ struct struct_with_pointer { int main() { int *ptr; - struct_with_pointer obj; - obj.data_in_struct = 10; + decomposed_struct_with_pointer obj1; + non_decomposed_struct_with_pointer obj2; + obj1.data_in_struct = 10; + obj2.data_in_struct = 10; sycl::kernel_single_task([=]() { *ptr = 50; - int local = obj.data_in_struct; + int local = obj1.data_in_struct + obj2.data_in_struct; }); } @@ -29,11 +39,7 @@ int main() { // CHECK:{ kernel_param_kind_t::kind_pointer, 8, 0 }, // CHECK:{ kernel_param_kind_t::kind_std_layout, 4, 8 }, // CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 16 }, -// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 24 }, -// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 32 }, -// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 40 }, -// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 48 }, -// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 56 }, -// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 64 }, -// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 72 }, -// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 80 }, +// CHECK:{ kernel_param_kind_t::kind_std_layout, 16, 24 }, +// CHECK:{ kernel_param_kind_t::kind_std_layout, 48, 40 }, +// CHECK:{ kernel_param_kind_t::kind_accessor, 4062, 88 }, +// CHECK:{ kernel_param_kind_t::kind_std_layout, 80, 104 }, diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index 13ad595d30e2c..88ef46be9f925 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -22,6 +22,7 @@ int main() { Accessor ReadWriteAccessor[2]; int Array[2]; int *ArrayOfPointers[2]; + int *ArrayOfPointers_2D[2][3]; struct StructWithAccessors { Accessor member_acc[2]; @@ -29,19 +30,6 @@ int main() { S s; - struct StructWithArrayOfPointers { - int x; - int y; - int *ArrayOfPtrs[2]; - }; - - struct DecomposedStruct { - int a; - StructWithArrayOfPointers SWPtrsMem[2]; - int *Array_2D_Ptrs[2][1]; - int c; - }; - // Not decomposed. struct NonDecomposedStruct { int a; @@ -49,19 +37,13 @@ int main() { int c; }; - struct StructWithSimplePointer { + struct StructWithPointers { int *Ptr; - int a; - }; - - struct StructWithNestedPointer { - StructWithSimplePointer SWPointer[2]; + int *ArrayOfPtrs[2]; }; - DecomposedStruct DecompStructArray[2]; NonDecomposedStruct NonDecompStructArray[2]; - StructWithSimplePointer StructWithSimplePointerArray[2]; - StructWithNestedPointer StructWithNestedPointerArray[2]; + StructWithPointers StructWithPointersArray[2]; int array_2D[2][3]; @@ -93,13 +75,6 @@ int main() { }); }); - myQueue.submit([&](sycl::handler &h) { - h.single_task( - [=] { - DecomposedStruct local = DecompStructArray[1]; - }); - }); - myQueue.submit([&](sycl::handler &h) { h.single_task( [=] { @@ -122,18 +97,20 @@ int main() { }); myQueue.submit([&](sycl::handler &h) { - h.single_task( + h.single_task( [=] { - StructWithSimplePointer local = StructWithSimplePointerArray[0]; + StructWithPointers local = StructWithPointersArray[0]; }); }); myQueue.submit([&](sycl::handler &h) { - h.single_task( + h.single_task( [=] { - StructWithNestedPointer local = StructWithNestedPointerArray[0]; + int local1 = *ArrayOfPointers_2D[0][0]; + int local2 = *ArrayOfPointers[0]; }); }); + } // Check Kernel_Accessor parameters @@ -173,21 +150,34 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_Array' '__wrapper_class' // Check Kernel_Array_Ptrs parameters -// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs{{.*}} 'void (__global int *, __global int *)' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers '__global int *' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers '__global int *' +// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs{{.*}} 'void (__wrapper_class)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers '__wrapper_class' // Check Kernel_Array_Ptrs inits // CHECK-NEXT: CompoundStmt // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl {{.*}} cinit // CHECK-NEXT: InitListExpr // CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' '__global int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' '__global int *' +// Initializer for ArrayOfPointers[0] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// Initializer for ArrayOfPointers[1] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 // Check Kernel_StructAccArray parameters // CHECK: FunctionDecl {{.*}}Kernel_StructAccArray{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)' @@ -216,140 +206,6 @@ int main() { // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}}__init -// Check Kernel_DecomposedStruct parameters -// CHECK: FunctionDecl {{.*}}Kernel_DecomposedStruct{{.*}} 'void (int, int, int, __wrapper_class, __wrapper_class, int, int, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, int, int, int, int, __wrapper_class, __wrapper_class, int, int, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, int)' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_x 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_y 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPtrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPtrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_x 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_y 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPtrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPtrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_Array_2D_Ptrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_Array_2D_Ptrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_c 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_x 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_y 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPtrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPtrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_x 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_y 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPtrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPtrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_Array_2D_Ptrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_Array_2D_Ptrs '__wrapper_class' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_c 'int' -// CHECK-NEXT: CompoundStmt -// CHECK-NEXT: DeclStmt -// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit -// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' - -// Initializer for struct array i.e. DecomposedStruct DecompStructArray[2] -// CHECK-NEXT: InitListExpr {{.*}} 'DecomposedStruct[2]' - -// Initializer for first element of DecompStructArray -// CHECK-NEXT: InitListExpr {{.*}} 'DecomposedStruct' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_a' 'int' - -// Initializer for struct array inside DecomposedStruct i.e. StructWithArrayOfPointers SWPtrsMem[2] -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers[2]' -// Initializer for first element of inner struct array -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_x' 'int' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_y' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' -// Initializer for second element of inner struct array -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_x' 'int' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_y' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][1]' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_Array_2D_Ptrs' '__wrapper_class' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_Array_2D_Ptrs' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar {{.*}} '_arg_c' 'int' - -// Initializer for second element of DecompStructArray -// CHECK-NEXT: InitListExpr {{.*}} 'DecomposedStruct' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_a' 'int' - -// Initializer for struct array inside DecomposedStruct i.e. StructWithArrayOfPointers SWPtrsMem[2] -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers[2]' -// Initializer for first element of inner struct array -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_x' 'int' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_y' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' -// Initializer for second element of inner struct array -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_x' 'int' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_y' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][1]' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_Array_2D_Ptrs' '__wrapper_class' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_Array_2D_Ptrs' '__wrapper_class' - // Check Kernel_TemplatedStructArray parameters // CHECK: FunctionDecl {{.*}}Kernel_TemplatedStructArray{{.*}} 'void (S)' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_s 'S':'S' @@ -414,71 +270,158 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_NonDecompStructArray' '__wrapper_class' // CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned -// Check Kernel_StructWithSimplePointer parameters. -// CHECK: FunctionDecl {{.*}}Kernel_StructWithSimplePointer{{.*}} 'void (__generated_StructWithSimplePointer, __generated_StructWithSimplePointer)' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_StructWithSimplePointerArray '__generated_StructWithSimplePointer' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_StructWithSimplePointerArray '__generated_StructWithSimplePointer' -// Check Kernel_StructWithSimplePointer inits +// Check Kernel_StructWithPointers parameters. +// CHECK: FunctionDecl {{.*}}Kernel_StructWithPointers{{.*}} 'void (__wrapper_class)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_StructWithPointersArray '__wrapper_class' +// Check Kernel_StructWithPointers inits // CHECK-NEXT: CompoundStmt // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl {{.*}} cinit // CHECK-NEXT: InitListExpr -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithSimplePointer[2]' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void (const StructWithSimplePointer &) noexcept' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithSimplePointer' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithSimplePointer *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_StructWithSimplePointerArray' '__generated_StructWithSimplePointer' - -// CHECK: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void (const StructWithSimplePointer &) noexcept' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithSimplePointer' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithSimplePointer *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_StructWithSimplePointerArray' '__generated_StructWithSimplePointer' - -// Check Kernel_StructWithNestedPointer parameters. -// CHECK: FunctionDecl {{.*}}Kernel_StructWithNestedPointer{{.*}} 'void (__generated_StructWithSimplePointer, __generated_StructWithSimplePointer, __generated_StructWithSimplePointer, __generated_StructWithSimplePointer)' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_SWPointer '__generated_StructWithSimplePointer' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_SWPointer '__generated_StructWithSimplePointer' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_SWPointer '__generated_StructWithSimplePointer' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_SWPointer '__generated_StructWithSimplePointer' -// Check Kernel_StructWithNestedPointer inits +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithPointers[2]' +// Initializer for StructWithPointersArray[0] +// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithPointers':'StructWithPointers' 'void (const StructWithPointers &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithPointers' lvalue +// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithPointers':'StructWithPointers' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithPointers *' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithPointers *' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__generated_StructWithPointers' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__generated_StructWithPointers *' +// CHECK-NEXT: MemberExpr {{.*}} '__generated_StructWithPointers[2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// Initializer for StructWithPointersArray[1] +// CHECK: CXXConstructExpr {{.*}} 'StructWithPointers':'StructWithPointers' 'void (const StructWithPointers &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithPointers' lvalue +// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithPointers':'StructWithPointers' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithPointers *' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithPointers *' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__generated_StructWithPointers' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__generated_StructWithPointers *' +// CHECK-NEXT: MemberExpr {{.*}} '__generated_StructWithPointers[2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 + +// Check Kernel_Array_Ptrs_2D parameters +// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs_2D 'void (__wrapper_class, __wrapper_class)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers_2D '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers '__wrapper_class' + +// Check Kernel_Array_Ptrs_2D inits // CHECK-NEXT: CompoundStmt // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl {{.*}} cinit // CHECK-NEXT: InitListExpr -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithNestedPointer[2]' -// InitListExpr for first array element of StructWithNestedPointerArray -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithNestedPointer' -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithSimplePointer[2]' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void (const StructWithSimplePointer &) noexcept' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithSimplePointer' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithSimplePointer *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_SWPointer' '__generated_StructWithSimplePointer' - -// CHECK: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void (const StructWithSimplePointer &) noexcept' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithSimplePointer' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithSimplePointer *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_SWPointer' '__generated_StructWithSimplePointer' - -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithNestedPointer' -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithSimplePointer[2]' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void (const StructWithSimplePointer &) noexcept' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithSimplePointer' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithSimplePointer *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_SWPointer' '__generated_StructWithSimplePointer' - -// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void (const StructWithSimplePointer &) noexcept' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithSimplePointer' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithSimplePointer *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_SWPointer' '__generated_StructWithSimplePointer' + +// Initializer for ArrayOfPointers_2D +// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][3]' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]' +// Initializer for ArrayOfPointers_2D[0][0] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 + +// Initializer for ArrayOfPointers_2D[0][1] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 + +// Initializer for ArrayOfPointers_2D[0][2] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 2 + +// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]' + +// Initializer for ArrayOfPointers_2D[1][0] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 + +// Initializer for ArrayOfPointers_2D[1][1] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 + +// Initializer for ArrayOfPointers_2D[1][2] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 2 + +// Initializer for ArrayOfPointers +// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' +// Initializer for ArrayOfPointers[0] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 + +// Initializer for ArrayOfPointers[1] +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp index d1df3c4467efe..e6093f7065e45 100644 --- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp +++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp @@ -79,72 +79,24 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_data' 'int' // Check kernel parameters -// CHECK: {{.*}}kernel_struct{{.*}} 'void (int, __wrapper_class, __wrapper_class, __wrapper_class -// CHECK-SAME: __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class)' -// CHECK: ParmVarDecl {{.*}} used _arg_data 'int' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr '__wrapper_class' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array1 '__wrapper_class' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array1 '__wrapper_class' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class' +// CHECK: {{.*}}kernel_struct{{.*}} 'void (__generated_test_struct)' +// CHECK: ParmVarDecl {{.*}} used _arg_s '__generated_test_struct' // Check that lambda field of struct type is initialized // CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' // CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' -// CHECK-NEXT: InitListExpr {{.*}}'test_struct'{{.*}} -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_data' 'int' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr' '__wrapper_class' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array1' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array1' '__wrapper_class' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][3]' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'test_struct':'test_struct' 'void (const test_struct &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const test_struct' lvalue +// CHECK-NEXT: UnaryOperator {{.*}} 'test_struct':'test_struct' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'test_struct *' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_test_struct *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_test_struct' lvalue ParmVar {{.*}} '_arg_s' // Check kernel parameters -// CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *, __global int *, __global int *)' +// CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *, __wrapper_class)' // CHECK: ParmVarDecl {{.*}} used _arg_new_data_addr '__global int *' // CHECK: ParmVarDecl {{.*}} used _arg_data_addr '__global int *' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array '__global int *' -// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array '__global int *' +// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array '__wrapper_class' // CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' // Check that lambda fields of pointer types are initialized @@ -156,9 +108,22 @@ int main() { // CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' // CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_data_addr' '__global int *' // CHECK: InitListExpr {{.*}} 'int *[2]' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_ptr_array' '__global int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_ptr_array' '__global int *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow +// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast +// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 + diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp index 808f9bcb79a8b..60b1926d0e57b 100644 --- a/clang/test/SemaSYCL/decomposition.cpp +++ b/clang/test/SemaSYCL/decomposition.cpp @@ -50,6 +50,7 @@ struct StructWithNonDecomposedStruct : StructNonDecomposed { struct StructWithPtr { StructNonDecomposed member; int *ptr; + int *ptrArr[2]; int i; }; @@ -179,12 +180,11 @@ int main() { }); // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (__generated_StructWithPtr)' - // FIXME: Stop decomposition of arrays with pointers StructWithArray t1; myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return t1.i; }); }); - // CHECK: FunctionDecl {{.*}}NestedArrayOfStructWithPointer{{.*}} 'void (__generated_StructWithPtr, __generated_StructWithPtr, __generated_StructWithPtr, StructNonDecomposed, int)' + // CHECK: FunctionDecl {{.*}}NestedArrayOfStructWithPointer{{.*}} 'void (__generated_StructWithArray)' DerivedStruct t2; myQueue.submit([&](sycl::handler &h) { @@ -194,8 +194,6 @@ int main() { } { - // FIXME: Stop decomposition for non-trivial types with pointers. - NonTrivialType NonTrivialStructWithPtr(10); myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return NonTrivialStructWithPtr.i;}); @@ -206,7 +204,7 @@ int main() { myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return NonTrivialTypeArray[0].i;}); }); - // CHECK: FunctionDecl {{.*}}ArrayOfNonTrivialStruct{{.*}} 'void (__generated_NonTrivialType, __generated_NonTrivialType)' + // CHECK: FunctionDecl {{.*}}ArrayOfNonTrivialStruct{{.*}} 'void (__wrapper_class)' NonTrivialDerived NonTrivialDerivedStructWithPtr(10); myQueue.submit([&](sycl::handler &h) { From 2dfe874bcecb17b58041f2fa3c7a2c59d72cb131 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Tue, 11 Oct 2022 11:47:32 -0700 Subject: [PATCH 2/4] Apply review comments Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 14 ++++++-------- 1 file changed, 6 insertions(+), 8 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 5c4c0425af06e..e89fa48f31650 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1498,7 +1498,7 @@ void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, visitComplexArray(Owner, Field, ArrayTy, Handlers...); } else if (AnyTrue:: Value) { - // We are currently in PointerHandler visitor + // We are currently in PointerHandler visitor. if (Field->hasAttr()) { // This is an array of pointers, or an array of a type containing // pointers. @@ -1877,7 +1877,7 @@ static QualType ModifyAddressSpace(Sema &SemaRef, QualType Ty) { Qualifiers Quals = PointeeTy.getQualifiers(); auto AS = Quals.getAddressSpace(); // Leave global_device and global_host address spaces as is to help FPGA - // device in memory allocations + // device in memory allocations. if (!PointeeTy->isFunctionType() && AS != LangAS::sycl_global_device && AS != LangAS::sycl_global_host) Quals.setAddressSpace(LangAS::sycl_global); @@ -1975,9 +1975,10 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { } bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - CXXRecordDecl *ModifiedRD = getGeneratedNewRecord(Ty->getAsCXXRecordDecl()); + // Add this record as a field of it's parent record if it is not an + // array element. if (!isArrayElement(FD, Ty)) addField(FD, QualType(ModifiedRD->getTypeForDecl(), 0)); else @@ -2019,7 +2020,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { // Add this array field as a field of it's parent record. addField(FD, ModifiedArray); } else { - // Multi-dimensional array element + // Multi-dimensional array element. ModifiedArrayElementsOrArray.push_back(ModifiedArray); } @@ -3202,17 +3203,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { Expr *createArraySubscriptExpr(uint64_t Index, Expr *ArrayRef) { QualType SizeT = SemaRef.getASTContext().getSizeType(); - llvm::APInt IndexVal{ static_cast(SemaRef.getASTContext().getTypeSize(SizeT)), Index, SizeT->isSignedIntegerType()}; - auto IndexLiteral = IntegerLiteral::Create( SemaRef.getASTContext(), IndexVal, SizeT, KernelCallerSrcLoc); - ExprResult IndexExpr = SemaRef.CreateBuiltinArraySubscriptExpr( ArrayRef, KernelCallerSrcLoc, IndexLiteral, KernelCallerSrcLoc); - assert(!IndexExpr.isInvalid()); return IndexExpr.get(); } @@ -3449,6 +3446,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { CollectionInitExprs.pop_back(); ArrayInfos.pop_back(); + // Remove the IndexExpr. if (!FD->hasAttr()) MemberExprBases.pop_back(); else From 28a35dad33a214bc8503c906732000dbbb198b82 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 12 Oct 2022 09:54:58 -0700 Subject: [PATCH 3/4] Apply review comments Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 59 ++++++++++++++------------ clang/test/CodeGenSYCL/inheritance.cpp | 7 +-- 2 files changed, 36 insertions(+), 30 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e89fa48f31650..45e09b22cc550 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1383,6 +1383,14 @@ class SyclKernelFieldHandler : public SyclKernelFieldHandlerBase { protected: Sema &SemaRef; SyclKernelFieldHandler(Sema &S) : SemaRef(S) {} + + // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) + // is an element of an array. FD will always be the array field. When + // traversing the array field, Ty will be the type of the array field or the + // type of array element (or some decomposed type from array). + bool isArrayElement(const FieldDecl *FD, QualType Ty) const { + return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); + } }; // A class to represent the 'do nothing' case for filtering purposes. @@ -1955,10 +1963,6 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return ModifiedRD; } - bool isArrayElement(const FieldDecl *FD, QualType Ty) const { - return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); - } - public: static constexpr const bool VisitInsideSimpleContainersWithPointer = true; static constexpr const bool VisitNthArrayElement = false; @@ -2065,7 +2069,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { } public: - QualType getNewType() { + QualType getNewRecordType() { CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); ModifiedRD->completeDefinition(); @@ -2270,14 +2274,14 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // defines a PointerHandler visitor which visits this record // recursively and modifies the address spaces of any pointer // found as required, thereby generating a new record with all - // pointers in 'right' address space. PointerHandler.getNewType() + // pointers in 'right' address space. PointerHandler.getNewRecordType() // returns this generated type. - QualType GenerateNewType(const CXXRecordDecl *RD) { + QualType GenerateNewRecordType(const CXXRecordDecl *RD) { SyclKernelPointerHandler PointerHandler(SemaRef, RD); KernelObjVisitor Visitor{SemaRef}; Visitor.VisitRecordBases(RD, PointerHandler); Visitor.VisitRecordFields(RD, PointerHandler); - return PointerHandler.getNewType(); + return PointerHandler.getNewRecordType(); } // If the array has been marked with SYCLGenerateNewTypeAttr, @@ -2285,7 +2289,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // of a type which contains pointers. This function generates // a new array with all pointers in the required address space. QualType GenerateNewArrayType(FieldDecl *FD, QualType FieldTy) { - const CXXRecordDecl *Owner = dyn_cast(FD->getParent()); + const auto *Owner = dyn_cast(FD->getParent()); SyclKernelPointerHandler PointerHandler(SemaRef); KernelObjVisitor Visitor{SemaRef}; Visitor.visitArray(Owner, FD, FieldTy, PointerHandler); @@ -2441,7 +2445,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // Check if we need to generate a new type for this record, // i.e. this record contains pointers. if (FieldRecordDecl->hasAttr()) - addParam(FD, GenerateNewType(FieldRecordDecl)); + addParam(FD, GenerateNewRecordType(FieldRecordDecl)); else addParam(FD, Ty); return true; @@ -2455,7 +2459,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // Check if we need to generate a new type for this record, // i.e. this record contains pointers. if (BaseRecordDecl->hasAttr()) - addParam(BS, GenerateNewType(BaseRecordDecl)); + addParam(BS, GenerateNewRecordType(BaseRecordDecl)); else addParam(BS, Ty); return true; @@ -2897,14 +2901,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return buildMemberExpr(DRE, ArrayField); } - // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) - // is an element of an array. This will determine whether we do - // MemberExprBases in some cases or not, AND determines how we initialize - // values. - bool isArrayElement(const FieldDecl *FD, QualType Ty) const { - return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); - } - // Creates an initialized entity for a field/item. In the case where this is a // field, returns a normal member initializer, if we're in a sub-array of a MD // array, returns an element initializer. @@ -3232,6 +3228,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, T, Initializer); } + // This function is recursive in order to handle + // multi-dimensional arrays. If the array element is + // an array, it implies that the array is multi-dimensional. + // We continue recursion till we reach a non-array element to + // generate required array subscript expressions. void createArrayInit(FieldDecl *FD, QualType T) { const ConstantArrayType *CAT = SemaRef.getASTContext().getAsConstantArrayType(T); @@ -3256,6 +3257,18 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { leaveArray(FD, T, ET); } + // This function is used to create initializers for a top + // level array which contains pointers. The openCl kernel + // parameter for this array will be a wrapper class + // which contains the generated type. This function generates + // code equivalent to: + // void ocl_kernel(__wrapper_class WrappedGT) { + // Kernel KernelObjClone { + // *reinterpret_cast(&WrappedGT.GeneratedArr[0]), + // *reinterpret_cast(&WrappedGT.GeneratedArr[1]), + // *reinterpret_cast(&WrappedGT.GeneratedArr[2]) + // }; + // } void handleGeneratedArrayType(FieldDecl *FD, QualType FieldTy) { ArrayParamBases.push_back(createSimpleArrayParamReferenceExpr(FieldTy)); createArrayInit(FD, FieldTy); @@ -3500,14 +3513,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { static_cast(CurOffset + OffsetAdj)); } - // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) - // is an element of an array. This will determine whether we do - // MemberExprBases in some cases or not, AND determines how we initialize - // values. - bool isArrayElement(const FieldDecl *FD, QualType Ty) const { - return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); - } - public: static constexpr const bool VisitInsideSimpleContainers = false; SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H, diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp index ddf1b78596000..2806779730e9a 100644 --- a/clang/test/CodeGenSYCL/inheritance.cpp +++ b/clang/test/CodeGenSYCL/inheritance.cpp @@ -5,6 +5,7 @@ class second_base { public: int *e; + int *arr[2]; second_base(int *E) : e(E) {} }; @@ -43,9 +44,9 @@ int main() { // CHECK: %struct.base = type { i32, %class.InnerField } // CHECK: %class.InnerField = type { %class.InnerFieldBase, i32 } // CHECK: %class.InnerFieldBase = type { i32 } -// CHECK: %class.__generated_second_base = type { ptr addrspace(1) } +// CHECK: %class.__generated_second_base = type { ptr addrspace(1), [2 x ptr addrspace(1)] } // CHECK: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }> -// CHECK: %class.second_base = type { ptr addrspace(4) } +// CHECK: %class.second_base = type { ptr addrspace(4), [2 x ptr addrspace(4)] } // Check kernel paramters // CHECK: define {{.*}}spir_kernel void @{{.*}}derived @@ -69,7 +70,7 @@ int main() { // First, derived-to-base cast with offset: // CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %[[LOCAL_OBJECT]], i64 16 // Initialize 'second_base' -// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[OFFSET_CALC]], ptr addrspace(4) align 8 %[[ARG_BASE1]], i64 8, i1 false) +// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[OFFSET_CALC]], ptr addrspace(4) align 8 %[[ARG_BASE1]], i64 24, i1 false) // Initialize field 'a' // CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 3 From d60780ab00ef87617f7e8274ad04c3797c9d65a5 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 12 Oct 2022 09:57:32 -0700 Subject: [PATCH 4/4] Apply review comment Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 45e09b22cc550..6570a50199430 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1883,7 +1883,7 @@ static QualType ModifyAddressSpace(Sema &SemaRef, QualType Ty) { // CodeGen. QualType PointeeTy = Ty->getPointeeType(); Qualifiers Quals = PointeeTy.getQualifiers(); - auto AS = Quals.getAddressSpace(); + LangAS AS = Quals.getAddressSpace(); // Leave global_device and global_host address spaces as is to help FPGA // device in memory allocations. if (!PointeeTy->isFunctionType() && AS != LangAS::sycl_global_device &&