diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b3ecda585cf3e..7dd4b7fa3dde6 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -628,6 +628,16 @@ static ParamDesc makeParamDesc(const FieldDecl *Src, QualType Ty) { Ctx.getTrivialTypeSourceInfo(Ty)); } +// Creates a parameter descriptor for kernel object +static ParamDesc makeParamDesc(const CXXRecordDecl *Src, QualType Ty) { + ASTContext &Ctx = Src->getASTContext(); + // There is no name available for lambda object. Name for all + // kernel types (lambda and functor) is set as _arg_kernelObject. + std::string Name = "_arg_kernelObject"; + return std::make_tuple(Ty, &Ctx.Idents.get(Name), + Ctx.getTrivialTypeSourceInfo(Ty)); +} + static ParamDesc makeParamDesc(ASTContext &Ctx, const CXXBaseSpecifier &Src, QualType Ty) { // TODO: There is no name for the base available, but duplicate names are @@ -721,12 +731,17 @@ static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent, // A visitor function that dispatches to functions as defined in // SyclKernelFieldHandler for the purposes of kernel generation. template -static void VisitRecordFields(RecordDecl::field_range Fields, +static void VisitRecordFields(CXXRecordDecl *KernelObject, Handlers &... handlers) { + + QualType KernelType = QualType(KernelObject->getTypeForDecl(), 0); + (void)std::initializer_list{ + (handlers.handleKernelObject(KernelObject, KernelType), 0)...}; + #define KF_FOR_EACH(FUNC) \ (void)std::initializer_list { (handlers.FUNC(Field, FieldTy), 0)... } - for (const auto &Field : Fields) { + for (const auto &Field : KernelObject->fields()) { QualType FieldTy = Field->getType(); if (Util::isSyclAccessorType(FieldTy)) @@ -781,6 +796,7 @@ template class SyclKernelFieldHandler { virtual void handlePointerType(FieldDecl *, QualType) {} virtual void handleArrayType(FieldDecl *, QualType) {} virtual void handleScalarType(FieldDecl *, QualType) {} + virtual void handleKernelObject(CXXRecordDecl *, QualType) {} // Most handlers shouldn't be handling this, just the field checker. virtual void handleOtherType(FieldDecl *, QualType) {} @@ -830,6 +846,16 @@ class SyclKernelFieldChecker << 1 << FieldTy; } } + void handleKernelObject(CXXRecordDecl *KernelObject, + QualType KernelType) final { + // TODO: Is this check correct? SYCL spec only talks about kernel defined as + // named function objects. What about lambda functions? + /*if (!KernelObject->isStandardLayoutType()) + IsInvalid = + Diag.Report(KernelObject->getLocation(), + diag::err_sycl_non_std_layout_type) + << KernelType;*/ + } // We should be able to handle this, so we made it part of the visitor, but // this is 'to be implemented'. @@ -860,6 +886,11 @@ class SyclKernelDeclCreator addParam(newParamDesc, FieldTy); } + void addParam(const CXXRecordDecl *KernelObject, QualType KernelType) { + ParamDesc newParamDesc = makeParamDesc(KernelObject, KernelType); + addParam(newParamDesc, KernelType); + } + void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), BS, FieldTy); @@ -985,12 +1016,9 @@ class SyclKernelDeclCreator addParam(FD, ModTy); } - void handleScalarType(FieldDecl *FD, QualType FieldTy) final { - addParam(FD, FieldTy); - } - + // TODO: Accessors in structs void handleStructType(FieldDecl *FD, QualType FieldTy) final { - addParam(FD, FieldTy); + // addParam(FD, FieldTy); } void handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { @@ -1002,6 +1030,11 @@ class SyclKernelDeclCreator // See https://github.com/intel/llvm/issues/1552 } + void handleKernelObject(CXXRecordDecl *KernelObject, + QualType KernelType) final { + addParam(KernelObject, KernelType); + } + void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } FunctionDecl *getKernelDecl() { return KernelDecl; } @@ -1010,6 +1043,8 @@ class SyclKernelDeclCreator return ArrayRef(std::begin(Params) + LastParamIndex, std::end(Params)); } + + ParmVarDecl *getKernelObjectParam() { return Params.front(); } }; class SyclKernelBodyCreator @@ -1017,9 +1052,6 @@ class SyclKernelBodyCreator SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; llvm::SmallVector FinalizeStmts; - llvm::SmallVector InitExprs; - VarDecl *KernelObjClone; - InitializedEntity VarEntity; CXXRecordDecl *KernelObj; llvm::SmallVector MemberExprBases; FunctionDecl *KernelCallerFunc; @@ -1029,22 +1061,24 @@ class SyclKernelBodyCreator // statements in advance to allocate it, so we cannot do this as we go along. CompoundStmt *createKernelBody() { - Expr *ILE = new (SemaRef.getASTContext()) InitListExpr( - SemaRef.getASTContext(), SourceLocation(), InitExprs, SourceLocation()); - ILE->setType(QualType(KernelObj->getTypeForDecl(), 0)); - KernelObjClone->setInit(ILE); Stmt *FunctionBody = KernelCallerFunc->getBody(); - ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); + // Kernel object parameter from kernel caller function + ParmVarDecl *KernelCallerObjParam = *(KernelCallerFunc->param_begin()); + // Kernel object parameter from generated kernel. + ParmVarDecl *KernelObjParam = DeclCreator.getKernelObjectParam(); // DeclRefExpr with valid source location but with decl which is not marked // as used is invalid. - KernelObjClone->setIsUsed(); + KernelObjParam->setIsUsed(); std::pair MappingPair = - std::make_pair(KernelObjParam, KernelObjClone); + std::make_pair(KernelCallerObjParam, KernelObjParam); // Push the Kernel function scope to ensure the scope isn't empty SemaRef.PushFunctionScope(); + + // Replacing all references to kernel caller function parameter in kernel + // body with references to kernel object parameter in generated kernel. KernelBodyTransform KBT(MappingPair, SemaRef); Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); BodyStmts.push_back(NewBody); @@ -1094,26 +1128,13 @@ class SyclKernelBodyCreator return Result; } + // TODO: Correct Stream + Accessors void createExprForStructOrScalar(FieldDecl *FD) { ParmVarDecl *KernelParameter = DeclCreator.getParamVarDeclsForCurrentField()[0]; - InitializedEntity Entity = - InitializedEntity::InitializeMember(FD, &VarEntity); QualType ParamType = KernelParameter->getOriginalType(); Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, SourceLocation()); - if (FD->getType()->isPointerType() && - FD->getType()->getPointeeType().getAddressSpace() != - ParamType->getPointeeType().getAddressSpace()) - DRE = ImplicitCastExpr::Create(SemaRef.Context, FD->getType(), - CK_AddressSpaceConversion, DRE, nullptr, - VK_RValue); - InitializationKind InitKind = - InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE); - - ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, DRE); - InitExprs.push_back(MemberInit.get()); } void createSpecialMethodCall(const CXXRecordDecl *SpecialClass, Expr *Base, @@ -1153,32 +1174,8 @@ class SyclKernelBodyCreator BodyStmts.push_back(Call); } - // FIXME Avoid creation of kernel obj clone. - // See https://github.com/intel/llvm/issues/1544 for details. - static VarDecl *createKernelObjClone(ASTContext &Ctx, DeclContext *DC, - CXXRecordDecl *KernelObj) { - TypeSourceInfo *TSInfo = - KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr; - VarDecl *VD = VarDecl::Create( - Ctx, DC, SourceLocation(), SourceLocation(), KernelObj->getIdentifier(), - QualType(KernelObj->getTypeForDecl(), 0), TSInfo, SC_None); - - return VD; - } - void handleSpecialType(FieldDecl *FD, QualType Ty) { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); - // Perform initialization only if it is field of kernel object - if (MemberExprBases.size() == 1) { - InitializedEntity Entity = - InitializedEntity::InitializeMember(FD, &VarEntity); - // Initialize with the default constructor. - InitializationKind InitKind = - InitializationKind::CreateDefault(SourceLocation()); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); - ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); - InitExprs.push_back(MemberInit.get()); - } createSpecialMethodCall(RecordDecl, MemberExprBases.back(), InitMethodName, FD); } @@ -1188,20 +1185,8 @@ class SyclKernelBodyCreator CXXRecordDecl *KernelObj, FunctionDecl *KernelCallerFunc) : SyclKernelFieldHandler(S), DeclCreator(DC), - KernelObjClone(createKernelObjClone(S.getASTContext(), - DC.getKernelDecl(), KernelObj)), - VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) { markParallelWorkItemCalls(); - - Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), - SourceLocation(), SourceLocation()); - BodyStmts.push_back(DS); - DeclRefExpr *KernelObjCloneRef = DeclRefExpr::Create( - S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone, - false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0), - VK_LValue); - MemberExprBases.push_back(KernelObjCloneRef); } ~SyclKernelBodyCreator() { @@ -1237,15 +1222,40 @@ class SyclKernelBodyCreator } void handlePointerType(FieldDecl *FD, QualType FieldTy) final { - createExprForStructOrScalar(FD); + ParmVarDecl *KernelParameter = + DeclCreator.getParamVarDeclsForCurrentField()[0]; + QualType ParamType = KernelParameter->getOriginalType(); + Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, + SourceLocation()); + if (FD->getType()->isPointerType() && + FD->getType()->getPointeeType().getAddressSpace() != + ParamType->getPointeeType().getAddressSpace()) + DRE = ImplicitCastExpr::Create(SemaRef.Context, FD->getType(), + CK_AddressSpaceConversion, DRE, nullptr, + VK_RValue); + + MemberExpr *KernelObjectPointerField = + BuildMemberExpr(MemberExprBases.back(), FD); + Expr *AssignPointerParameter = BinaryOperator::Create( + SemaRef.getASTContext(), KernelObjectPointerField, DRE, BO_Assign, + FieldTy, VK_LValue, OK_Ordinary, SourceLocation(), + FPOptions(SemaRef.getASTContext().getLangOpts())); + + BodyStmts.push_back(AssignPointerParameter); } + // TODO: Accessors in structs void handleStructType(FieldDecl *FD, QualType FieldTy) final { - createExprForStructOrScalar(FD); + // createExprForStructOrScalar(FD); } - void handleScalarType(FieldDecl *FD, QualType FieldTy) final { - createExprForStructOrScalar(FD); + void handleKernelObject(CXXRecordDecl *KernelObject, QualType KernelType) { + ParmVarDecl *KernelParameter = + DeclCreator.getParamVarDeclsForCurrentField()[0]; + QualType ParamType = KernelParameter->getOriginalType(); + Expr *KernelObjRef = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, + VK_LValue, SourceLocation()); + MemberExprBases.push_back(KernelObjRef); } void enterStruct(const CXXRecordDecl *, FieldDecl *FD) final { @@ -1376,6 +1386,14 @@ class SyclKernelIntHeaderCreator CurStruct = FD->getType()->getAsCXXRecordDecl(); CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8; } + void handleKernelObject(CXXRecordDecl *KernelObject, + QualType KernelType) final { + uint64_t Size = + SemaRef.getASTContext().getTypeSizeInChars(KernelType).getQuantity(); + // Offset for kernel object is 0 + Header.addParamDesc(SYCLIntegrationHeader::kind_std_layout, + static_cast(Size), 0); + } void leaveStruct(const CXXRecordDecl *RD, FieldDecl *FD) final { CurStruct = RD; @@ -1447,7 +1465,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, StableName); ConstructingOpenCLKernel = true; - VisitRecordFields(KernelLambda->fields(), checker, kernel_decl, kernel_body, + VisitRecordFields(KernelLambda, checker, kernel_decl, kernel_body, int_header); ConstructingOpenCLKernel = false; } diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp index d52ba4c13a7f7..8b82d82717a03 100644 --- a/clang/test/CodeGenSYCL/device-functions.cpp +++ b/clang/test/CodeGenSYCL/device-functions.cpp @@ -21,7 +21,7 @@ int main() { kernel_single_task([]() { foo(); }); return 0; } -// CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel() +// CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) // CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %this) // CHECK: define spir_func void @_Z3foov() // CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg) diff --git a/clang/test/CodeGenSYCL/image_accessor.cpp b/clang/test/CodeGenSYCL/image_accessor.cpp index 8ad7992b56a0a..d1b8a1368a097 100644 --- a/clang/test/CodeGenSYCL/image_accessor.cpp +++ b/clang/test/CodeGenSYCL/image_accessor.cpp @@ -7,27 +7,27 @@ // RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DWO // // CHECK-1DRO: %opencl.image1d_ro_t = type opaque -// CHECK-1DRO: define spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-1DRO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 8 %_arg_kernelObject, %opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) // CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}}) // // CHECK-2DRO: %opencl.image2d_ro_t = type opaque -// CHECK-2DRO: define spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-2DRO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.0"* byval(%"class.{{.*}}.anon.0") align 8 %_arg_kernelObject, %opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) // CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}}) // // CHECK-3DRO: %opencl.image3d_ro_t = type opaque -// CHECK-3DRO: define spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-3DRO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.1"* byval(%"class.{{.*}}.anon.1") align 8 %_arg_kernelObject, %opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) // CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}}) // // CHECK-1DWO: %opencl.image1d_wo_t = type opaque -// CHECK-1DWO: define spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-1DWO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.2"* byval(%"class.{{.*}}.anon.2") align 8 %_arg_kernelObject, %opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) // CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}}) // // CHECK-2DWO: %opencl.image2d_wo_t = type opaque -// CHECK-2DWO: define spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-2DWO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.3"* byval(%"class.{{.*}}.anon.3") align 8 %_arg_kernelObject, %opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) // CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}}) // // CHECK-3DWO: %opencl.image3d_wo_t = type opaque -// CHECK-3DWO: define spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-3DWO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.4"* byval(%"class.{{.*}}.anon.4") align 8 %_arg_kernelObject, %opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) // CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}}) // // TODO: Add tests for the image_array opencl datatype support. diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index 84b35578f48e6..b9969095fcd32 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -1,6 +1,7 @@ // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only // RUN: FileCheck -input-file=%t.h %s -// +// FIXME: Check incorrect header generation for accessor in base classes. +// XFAIL: * // CHECK: #include // // CHECK: class first_kernel; @@ -27,22 +28,26 @@ // CHECK: static constexpr // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- _ZTSZ4mainE12first_kernel +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 40, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 16 }, // CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 32 }, // CHECK-EMPTY: // CHECK-NEXT: //--- _ZTSN16second_namespace13second_kernelIcEE +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 16 }, // CHECK-EMPTY: // CHECK-NEXT: //--- _ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 16 }, // CHECK-EMPTY: // CHECK-NEXT: //--- _ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 16, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 }, // CHECK-EMPTY: diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp index d1352b190fa94..18e02fc695380 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp @@ -21,8 +21,8 @@ void bar() { []() [[intelfpga::no_global_work_offset(0)]]{}); } -// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !no_global_work_offset ![[NUM5]] -// CHECK: define spir_kernel void @{{.*}}kernel_name3() {{.*}} ![[NUM4:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name1(%class.{{.*}}.Foo* byval(%class.{{.*}}.Foo) align 1 %_arg_kernelObject) {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name2(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) {{.*}} !no_global_work_offset ![[NUM5]] +// CHECK: define spir_kernel void @{{.*}}kernel_name3(%"class.{{.*}}.anon.0"* byval(%"class.{{.*}}.anon.0") align 1 %_arg_kernelObject) {{.*}} ![[NUM4:[0-9]+]] // CHECK-NOT: ![[NUM4]] = !{i32 0} // CHECK: ![[NUM5]] = !{} diff --git a/clang/test/CodeGenSYCL/intel-fpga-reg.cpp b/clang/test/CodeGenSYCL/intel-fpga-reg.cpp index 9428243813f40..d3e6365a1d956 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-reg.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-reg.cpp @@ -33,113 +33,113 @@ void foo() { int a=123; myInt myA = 321; int b = __builtin_intel_fpga_reg(a); -// CHECK: %[[V_A1:[0-9]+]] = load i32, i32* %a, align 4, !tbaa !9 -// CHECK-NEXT: %[[V_A2:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_A1]], [[BIFR_STR:i8\* getelementptr inbounds \(\[25 x i8\], \[25 x i8\]\* @.str, i32 0, i32 0\),]] -// CHECK-NEXT: store i32 %[[V_A2]], i32* %b, align 4, !tbaa !9 + // CHECK: %[[V_A1:[0-9]+]] = load i32, i32* %a, align 4, !tbaa [[ONE:![0-9]*]] + // CHECK-NEXT: %[[V_A2:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_A1]], [[BIFR_STR:i8\* getelementptr inbounds \(\[25 x i8\], \[25 x i8\]\* @.str, i32 0, i32 0\),]] + // CHECK-NEXT: store i32 %[[V_A2]], i32* %b, align 4, !tbaa [[ONE]] int myB = __builtin_intel_fpga_reg(myA); -// CHECK: %[[V_MYA1:[0-9]+]] = load i32, i32* %myA -// CHECK-NEXT: %[[V_MYA2:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_MYA1]], [[BIFR_STR]] -// CHECK-NEXT: store i32 %[[V_MYA2]], i32* %myB, align 4, !tbaa !9 + // CHECK: %[[V_MYA1:[0-9]+]] = load i32, i32* %myA + // CHECK-NEXT: %[[V_MYA2:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_MYA1]], [[BIFR_STR]] + // CHECK-NEXT: store i32 %[[V_MYA2]], i32* %myB, align 4, !tbaa [[ONE]] int c = __builtin_intel_fpga_reg(2.0f); -// CHECK: %[[V_CF1:[0-9]+]] = call i32 @llvm.annotation.i32(i32 1073741824, [[BIFR_STR]] -// CHECK-NEXT: %[[V_FBITCAST:[0-9]+]] = bitcast i32 %[[V_CF1]] to float -// CHECK-NEXT: %[[V_CF2:conv]] = fptosi float %[[V_FBITCAST]] to i32 -// CHECK-NEXT: store i32 %[[V_CF2]], i32* %c, align 4, !tbaa !9 + // CHECK: %[[V_CF1:[0-9]+]] = call i32 @llvm.annotation.i32(i32 1073741824, [[BIFR_STR]] + // CHECK-NEXT: %[[V_FBITCAST:[0-9]+]] = bitcast i32 %[[V_CF1]] to float + // CHECK-NEXT: %[[V_CF2:conv]] = fptosi float %[[V_FBITCAST]] to i32 + // CHECK-NEXT: store i32 %[[V_CF2]], i32* %c, align 4, !tbaa [[ONE]] int d = __builtin_intel_fpga_reg( __builtin_intel_fpga_reg( b+12 )); -// CHECK: %[[V_B1:[0-9]+]] = load i32, i32* %b -// CHECK-NEXT: %[[V_B2:add]] = add nsw i32 %[[V_B1]], 12 -// CHECK-NEXT: %[[V_B3:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_B2]], [[BIFR_STR]] -// CHECK-NEXT: %[[V_B4:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_B3]], [[BIFR_STR]] -// CHECK-NEXT: store i32 %[[V_B4]], i32* %d, align 4, !tbaa !9 + // CHECK: %[[V_B1:[0-9]+]] = load i32, i32* %b + // CHECK-NEXT: %[[V_B2:add]] = add nsw i32 %[[V_B1]], 12 + // CHECK-NEXT: %[[V_B3:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_B2]], [[BIFR_STR]] + // CHECK-NEXT: %[[V_B4:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_B3]], [[BIFR_STR]] + // CHECK-NEXT: store i32 %[[V_B4]], i32* %d, align 4, !tbaa [[ONE]] int e = __builtin_intel_fpga_reg( __builtin_intel_fpga_reg( a+b )); -// CHECK: %[[V_AB1:[0-9]+]] = load i32, i32* %a -// CHECK-NEXT: %[[V_AB2:[0-9]+]] = load i32, i32* %b -// CHECK-NEXT: %[[V_AB3:add[0-9]+]] = add nsw i32 %[[V_AB1]], %[[V_AB2]] -// CHECK-NEXT: %[[V_AB4:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_AB3]], [[BIFR_STR]] -// CHECK-NEXT: %[[V_AB5:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_AB4]], [[BIFR_STR]] -// CHECK-NEXT: store i32 %[[V_AB5]], i32* %e, align 4, !tbaa !9 + // CHECK: %[[V_AB1:[0-9]+]] = load i32, i32* %a + // CHECK-NEXT: %[[V_AB2:[0-9]+]] = load i32, i32* %b + // CHECK-NEXT: %[[V_AB3:add[0-9]+]] = add nsw i32 %[[V_AB1]], %[[V_AB2]] + // CHECK-NEXT: %[[V_AB4:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_AB3]], [[BIFR_STR]] + // CHECK-NEXT: %[[V_AB5:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_AB4]], [[BIFR_STR]] + // CHECK-NEXT: store i32 %[[V_AB5]], i32* %e, align 4, !tbaa [[ONE]] int f; f = __builtin_intel_fpga_reg(a); -// CHECK: %[[V_F1:[0-9]+]] = load i32, i32* %a -// CHECK-NEXT: %[[V_F2:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_F1]], [[BIFR_STR]] -// CHECK-NEXT: store i32 %[[V_F2]], i32* %f, align 4, !tbaa !9 + // CHECK: %[[V_F1:[0-9]+]] = load i32, i32* %a + // CHECK-NEXT: %[[V_F2:[0-9]+]] = call i32 @llvm.annotation.i32(i32 %[[V_F1]], [[BIFR_STR]] + // CHECK-NEXT: store i32 %[[V_F2]], i32* %f, align 4, !tbaa [[ONE]] struct st i = {1, 5.0f}; struct st i2 = i; struct st ii = __builtin_intel_fpga_reg(i); -// CHECK: %[[V_TI1:[0-9]+]] = bitcast %[[T_ST]]* %agg-temp to i8* -// CHECK-NEXT: %[[V_I:[0-9]+]] = bitcast %[[T_ST]]* %i to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TI1]], i8* align 4 %[[V_I]], i64 8, i1 false), !tbaa.struct !11 -// CHECK-NEXT: %[[V_TI2:[0-9]+]] = bitcast %[[T_ST]]* %agg-temp to i8* -// CHECK-NEXT: %[[V_TI3:[0-9]+]] = call i8* @llvm.ptr.annotation.p0i8(i8* %[[V_TI2]], [[BIFR_STR]] -// CHECK-NEXT: %[[V_TI4:[0-9]+]] = bitcast i8* %[[V_TI3]] to %[[T_ST]]* -// CHECK-NEXT: %[[V_II:[0-9]+]] = bitcast %[[T_ST]]* %ii to i8* -// CHECK-NEXT: %[[V_TI5:[0-9]+]] = bitcast %[[T_ST]]* %[[V_TI4]] to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_II]], i8* align 4 %[[V_TI5]], i64 8, i1 false) + // CHECK: %[[V_TI1:[0-9]+]] = bitcast %[[T_ST]]* %agg-temp to i8* + // CHECK-NEXT: %[[V_I:[0-9]+]] = bitcast %[[T_ST]]* %i to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TI1]], i8* align 4 %[[V_I]], i64 8, i1 false), !tbaa.struct [[TWO:![0-9]*]] + // CHECK-NEXT: %[[V_TI2:[0-9]+]] = bitcast %[[T_ST]]* %agg-temp to i8* + // CHECK-NEXT: %[[V_TI3:[0-9]+]] = call i8* @llvm.ptr.annotation.p0i8(i8* %[[V_TI2]], [[BIFR_STR]] + // CHECK-NEXT: %[[V_TI4:[0-9]+]] = bitcast i8* %[[V_TI3]] to %[[T_ST]]* + // CHECK-NEXT: %[[V_II:[0-9]+]] = bitcast %[[T_ST]]* %ii to i8* + // CHECK-NEXT: %[[V_TI5:[0-9]+]] = bitcast %[[T_ST]]* %[[V_TI4]] to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_II]], i8* align 4 %[[V_TI5]], i64 8, i1 false) struct st iii; iii = __builtin_intel_fpga_reg(ii); -// CHECK: %[[V_TII1:[0-9]+]] = bitcast %[[T_ST]]* %agg-temp2 to i8* -// CHECK-NEXT: %[[V_II:[0-9]+]] = bitcast %[[T_ST]]* %ii to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TII1]], i8* align 4 %[[V_II]], i64 8, i1 false), !tbaa.struct !11 -// CHECK-NEXT: %[[V_TII2:[0-9]+]] = bitcast %[[T_ST]]* %agg-temp2 to i8* -// CHECK-NEXT: %[[V_TII3:[0-9]+]] = call i8* @llvm.ptr.annotation.p0i8(i8* %[[V_TII2]], [[BIFR_STR]] -// CHECK-NEXT: %[[V_TII4:[0-9]+]] = bitcast i8* %[[V_TII3]] to %[[T_ST]]* -// CHECK-NEXT: %[[V_TII5:[0-9]+]] = bitcast %[[T_ST]]* %ref.tmp to i8* -// CHECK-NEXT: %[[V_TII6:[0-9]+]] = bitcast %[[T_ST]]* %[[V_TII4]] to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TII5]], i8* align 4 %[[V_TII6]], i64 8, i1 false) -// CHECK-NEXT: %[[V_TIII:[0-9]+]] = bitcast %[[T_ST]]* %iii to i8* -// CHECK-NEXT: %[[V_TII7:[0-9]+]] = bitcast %[[T_ST]]* %ref.tmp to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TIII]], i8* align 4 %[[V_TII7]], i64 8, i1 false), !tbaa.struct !11 + // CHECK: %[[V_TII1:[0-9]+]] = bitcast %[[T_ST]]* %agg-temp2 to i8* + // CHECK-NEXT: %[[V_II:[0-9]+]] = bitcast %[[T_ST]]* %ii to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TII1]], i8* align 4 %[[V_II]], i64 8, i1 false), !tbaa.struct [[TWO]] + // CHECK-NEXT: %[[V_TII2:[0-9]+]] = bitcast %[[T_ST]]* %agg-temp2 to i8* + // CHECK-NEXT: %[[V_TII3:[0-9]+]] = call i8* @llvm.ptr.annotation.p0i8(i8* %[[V_TII2]], [[BIFR_STR]] + // CHECK-NEXT: %[[V_TII4:[0-9]+]] = bitcast i8* %[[V_TII3]] to %[[T_ST]]* + // CHECK-NEXT: %[[V_TII5:[0-9]+]] = bitcast %[[T_ST]]* %ref.tmp to i8* + // CHECK-NEXT: %[[V_TII6:[0-9]+]] = bitcast %[[T_ST]]* %[[V_TII4]] to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TII5]], i8* align 4 %[[V_TII6]], i64 8, i1 false) + // CHECK-NEXT: %[[V_TIII:[0-9]+]] = bitcast %[[T_ST]]* %iii to i8* + // CHECK-NEXT: %[[V_TII7:[0-9]+]] = bitcast %[[T_ST]]* %ref.tmp to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TIII]], i8* align 4 %[[V_TII7]], i64 8, i1 false), !tbaa.struct [[TWO]] struct st *iiii = __builtin_intel_fpga_reg(&iii); -// CHECK: %[[V_T3I0:[0-9]+]] = ptrtoint %[[T_ST]]* %iii to i64 -// CHECK-NEXT: %[[V_T3I1:[0-9]+]] = call i64 @llvm.annotation.i64(i64 %[[V_T3I0]], [[BIFR_STR]] -// CHECK-NEXT: %[[V_T3I2:[0-9]+]] = inttoptr i64 %[[V_T3I1]] to %[[T_ST]]* -// CHECK-NEXT: %[[V_T3I3:[0-9]+]] = addrspacecast %[[T_ST]]* %[[V_T3I2]] to %[[T_ST]] addrspace(4)* -// CHECK-NEXT: store %[[T_ST]] addrspace(4)* %[[V_T3I3]], %[[T_ST]] addrspace(4)** %iiii, align 8, !tbaa !5 + // CHECK: %[[V_T3I0:[0-9]+]] = ptrtoint %[[T_ST]]* %iii to i64 + // CHECK-NEXT: %[[V_T3I1:[0-9]+]] = call i64 @llvm.annotation.i64(i64 %[[V_T3I0]], [[BIFR_STR]] + // CHECK-NEXT: %[[V_T3I2:[0-9]+]] = inttoptr i64 %[[V_T3I1]] to %[[T_ST]]* + // CHECK-NEXT: %[[V_T3I3:[0-9]+]] = addrspacecast %[[T_ST]]* %[[V_T3I2]] to %[[T_ST]] addrspace(4)* + // CHECK-NEXT: store %[[T_ST]] addrspace(4)* %[[V_T3I3]], %[[T_ST]] addrspace(4)** %iiii, align 8, !tbaa [[THREE:![0-9]*]] union un u1 = {1}; union un u2, *u3; u2 = __builtin_intel_fpga_reg(u1); -// CHECK: %[[V_TU1:[0-9]+]] = bitcast %[[T_UN]]* %agg-temp4 to i8* -// CHECK-NEXT: %[[V_TU2:[0-9]+]] = bitcast %[[T_UN]]* %u1 to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TU1]], i8* align 4 %[[V_TU2]], i64 4, i1 false), !tbaa.struct !14 -// CHECK-NEXT: %[[V_TU3:[0-9]+]] = bitcast %[[T_UN]]* %agg-temp4 to i8* -// CHECK-NEXT: %[[V_TU4:[0-9]+]] = call i8* @llvm.ptr.annotation.p0i8(i8* %[[V_TU3]], [[BIFR_STR]] -// CHECK-NEXT: %[[V_TU5:[0-9]+]] = bitcast i8* %[[V_TU4]] to %[[T_UN]]* -// CHECK-NEXT: %[[V_TU6:[0-9]+]] = bitcast %[[T_UN]]* %ref.tmp3 to i8* -// CHECK-NEXT: %[[V_TU7:[0-9]+]] = bitcast %[[T_UN]]* %[[V_TU5]] to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TU6]], i8* align 4 %[[V_TU7]], i64 8, i1 false) -// CHECK-NEXT: %[[V_TU8:[0-9]+]] = bitcast %[[T_UN]]* %u2 to i8* -// CHECK-NEXT: %[[V_TU9:[0-9]+]] = bitcast %[[T_UN]]* %ref.tmp3 to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TU8]], i8* align 4 %[[V_TU9]], i64 4, i1 false), !tbaa.struct !14 + // CHECK: %[[V_TU1:[0-9]+]] = bitcast %[[T_UN]]* %agg-temp4 to i8* + // CHECK-NEXT: %[[V_TU2:[0-9]+]] = bitcast %[[T_UN]]* %u1 to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TU1]], i8* align 4 %[[V_TU2]], i64 4, i1 false), !tbaa.struct [[FOUR:![0-9]*]] + // CHECK-NEXT: %[[V_TU3:[0-9]+]] = bitcast %[[T_UN]]* %agg-temp4 to i8* + // CHECK-NEXT: %[[V_TU4:[0-9]+]] = call i8* @llvm.ptr.annotation.p0i8(i8* %[[V_TU3]], [[BIFR_STR]] + // CHECK-NEXT: %[[V_TU5:[0-9]+]] = bitcast i8* %[[V_TU4]] to %[[T_UN]]* + // CHECK-NEXT: %[[V_TU6:[0-9]+]] = bitcast %[[T_UN]]* %ref.tmp3 to i8* + // CHECK-NEXT: %[[V_TU7:[0-9]+]] = bitcast %[[T_UN]]* %[[V_TU5]] to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TU6]], i8* align 4 %[[V_TU7]], i64 8, i1 false) + // CHECK-NEXT: %[[V_TU8:[0-9]+]] = bitcast %[[T_UN]]* %u2 to i8* + // CHECK-NEXT: %[[V_TU9:[0-9]+]] = bitcast %[[T_UN]]* %ref.tmp3 to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TU8]], i8* align 4 %[[V_TU9]], i64 4, i1 false), !tbaa.struct [[FOUR]] u3 = __builtin_intel_fpga_reg(&u2); -// CHECK: %[[V_TPU1:[0-9]+]] = ptrtoint %[[T_UN]]* %u2 to i64 -// CHECK-NEXT: %[[V_TPU2:[0-9]+]] = call i64 @llvm.annotation.i64(i64 %[[V_TPU1]], [[BIFR_STR]] -// CHECK-NEXT: %[[V_TPU3:[0-9]+]] = inttoptr i64 %[[V_TPU2]] to %[[T_UN]]* -// CHECK-NEXT: %[[V_TPU4:[0-9]+]] = addrspacecast %[[T_UN]]* %[[V_TPU3]] to %[[T_UN]] addrspace(4)* -// CHECK-NEXT: store %[[T_UN]] addrspace(4)* %[[V_TPU4]], %[[T_UN]] addrspace(4)** %u3, align 8, !tbaa !5 + // CHECK: %[[V_TPU1:[0-9]+]] = ptrtoint %[[T_UN]]* %u2 to i64 + // CHECK-NEXT: %[[V_TPU2:[0-9]+]] = call i64 @llvm.annotation.i64(i64 %[[V_TPU1]], [[BIFR_STR]] + // CHECK-NEXT: %[[V_TPU3:[0-9]+]] = inttoptr i64 %[[V_TPU2]] to %[[T_UN]]* + // CHECK-NEXT: %[[V_TPU4:[0-9]+]] = addrspacecast %[[T_UN]]* %[[V_TPU3]] to %[[T_UN]] addrspace(4)* + // CHECK-NEXT: store %[[T_UN]] addrspace(4)* %[[V_TPU4]], %[[T_UN]] addrspace(4)** %u3, align 8, !tbaa [[THREE]] A ca(213); A cb = __builtin_intel_fpga_reg(ca); -// CHECK: %[[V_TCA1:[0-9]+]] = bitcast %[[T_CL]]* %agg-temp5 to i8* -// CHECK-NEXT: %[[V_CA:[0-9]+]] = bitcast %[[T_CL]]* %ca to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TCA1]], i8* align 4 %[[V_CA]], i64 4, i1 false), !tbaa.struct !16 -// CHECK-NEXT: %[[V_TCA2:[0-9]+]] = bitcast %[[T_CL]]* %agg-temp5 to i8* -// CHECK-NEXT: %[[V_TCA3:[0-9]+]] = call i8* @llvm.ptr.annotation.p0i8(i8* %[[V_TCA2]], [[BIFR_STR]] -// CHECK-NEXT: %[[V_TCA4:[0-9]+]] = bitcast i8* %[[V_TCA3]] to %[[T_CL]]* -// CHECK-NEXT: %[[V_CB:[0-9]+]] = bitcast %[[T_CL]]* %cb to i8* -// CHECK-NEXT: %[[V_TCA5:[0-9]+]] = bitcast %[[T_CL]]* %[[V_TCA4]] to i8* -// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_CB]], i8* align 4 %[[V_TCA5]], i64 8, i1 false) + // CHECK: %[[V_TCA1:[0-9]+]] = bitcast %[[T_CL]]* %agg-temp5 to i8* + // CHECK-NEXT: %[[V_CA:[0-9]+]] = bitcast %[[T_CL]]* %ca to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_TCA1]], i8* align 4 %[[V_CA]], i64 4, i1 false), !tbaa.struct [[FIVE:![0-9]*]] + // CHECK-NEXT: %[[V_TCA2:[0-9]+]] = bitcast %[[T_CL]]* %agg-temp5 to i8* + // CHECK-NEXT: %[[V_TCA3:[0-9]+]] = call i8* @llvm.ptr.annotation.p0i8(i8* %[[V_TCA2]], [[BIFR_STR]] + // CHECK-NEXT: %[[V_TCA4:[0-9]+]] = bitcast i8* %[[V_TCA3]] to %[[T_CL]]* + // CHECK-NEXT: %[[V_CB:[0-9]+]] = bitcast %[[T_CL]]* %cb to i8* + // CHECK-NEXT: %[[V_TCA5:[0-9]+]] = bitcast %[[T_CL]]* %[[V_TCA4]] to i8* + // CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[V_CB]], i8* align 4 %[[V_TCA5]], i64 8, i1 false) int *ap = &a; int *bp = __builtin_intel_fpga_reg(ap); -// CHECK: %[[V_AP0:[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %ap, align 8, !tbaa !5 -// CHECK-NEXT: %[[V_AP1:[0-9]+]] = ptrtoint i32 addrspace(4)* %[[V_AP0]] to i64 -// CHECK-NEXT: %[[V_AP2:[0-9]+]] = call i64 @llvm.annotation.i64(i64 %[[V_AP1]], [[BIFR_STR]] -// CHECK-NEXT: %[[V_AP3:[0-9]+]] = inttoptr i64 %[[V_AP2]] to i32 addrspace(4)* -// CHECK-NEXT: store i32 addrspace(4)* %[[V_AP3]], i32 addrspace(4)** %bp, align 8, !tbaa !5 + // CHECK: %[[V_AP0:[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %ap, align 8, !tbaa [[THREE]] + // CHECK-NEXT: %[[V_AP1:[0-9]+]] = ptrtoint i32 addrspace(4)* %[[V_AP0]] to i64 + // CHECK-NEXT: %[[V_AP2:[0-9]+]] = call i64 @llvm.annotation.i64(i64 %[[V_AP1]], [[BIFR_STR]] + // CHECK-NEXT: %[[V_AP3:[0-9]+]] = inttoptr i64 %[[V_AP2]] to i32 addrspace(4)* + // CHECK-NEXT: store i32 addrspace(4)* %[[V_AP3]], i32 addrspace(4)** %bp, align 8, !tbaa [[THREE]] } template diff --git a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp index 5208db6ec3908..e7c50ce3652bc 100644 --- a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp +++ b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp @@ -18,7 +18,7 @@ void bar() { []() [[intelfpga::max_global_work_dim(2)]] {}); } -// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !max_global_work_dim ![[NUM1:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !max_global_work_dim ![[NUM8:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name1(%class.{{.*}}.Foo* byval(%class.{{.*}}.Foo) align 1 %_arg_kernelObject) {{.*}} !max_global_work_dim ![[NUM1:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name2(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) {{.*}} !max_global_work_dim ![[NUM8:[0-9]+]] // CHECK: ![[NUM1]] = !{i32 1} // CHECK: ![[NUM8]] = !{i32 2} diff --git a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp index 13bbb54f34198..1c45fa07544ca 100644 --- a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp @@ -18,7 +18,7 @@ void bar() { []() [[intelfpga::max_work_group_size(8, 8, 8)]] {}); } -// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !max_work_group_size ![[NUM1:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !max_work_group_size ![[NUM8:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name1(%class.{{.*}}.Foo* byval(%class.{{.*}}.Foo) align 1 %_arg_kernelObject) {{.*}} !max_work_group_size ![[NUM1:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name2(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) {{.*}} !max_work_group_size ![[NUM8:[0-9]+]] // CHECK: ![[NUM1]] = !{i32 1, i32 1, i32 1} // CHECK: ![[NUM8]] = !{i32 8, i32 8, i32 8} diff --git a/clang/test/CodeGenSYCL/intel-restrict.cpp b/clang/test/CodeGenSYCL/intel-restrict.cpp index 06d9d7ef4d59f..bdaeb76ee0ea2 100644 --- a/clang/test/CodeGenSYCL/intel-restrict.cpp +++ b/clang/test/CodeGenSYCL/intel-restrict.cpp @@ -11,7 +11,7 @@ int main() { int *c; kernel( [a,b,c]() [[intel::kernel_args_restrict]] { c[0] = a[0] + b[0];}); -// CHECK: define spir_kernel {{.*}}kernel_restrict(i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}}) + // CHECK: define spir_kernel {{.*}}kernel_restrict(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 8 %_arg_kernelObject, i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}}) int *d; int *e; @@ -19,10 +19,10 @@ int main() { kernel( [d,e,f]() { f[0] = d[0] + e[0];}); -// CHECK: define spir_kernel {{.*}}kernel_norestrict(i32 addrspace(1)* %{{.*}}, i32 addrspace(1)* %{{.*}}, i32 addrspace(1)* %{{.*}}) + // CHECK: define spir_kernel {{.*}}kernel_norestrict(%"class.{{.*}}.anon.0"* byval(%"class.{{.*}}.anon.0") align 8 %_arg_kernelObject, i32 addrspace(1)* %{{.*}}, i32 addrspace(1)* %{{.*}}, i32 addrspace(1)* %{{.*}}) int g = 42; kernel( [a,b,c,g]() [[intel::kernel_args_restrict]] { c[0] = a[0] + b[0] + g;}); -// CHECK: define spir_kernel {{.*}}kernel_restrict_other_types(i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}}, i32 %{{.*}}) + // CHECK: define spir_kernel {{.*}}kernel_restrict_other_types(%"class.{{.*}}.anon.1"* byval(%"class.{{.*}}.anon.1") align 8 %_arg_kernelObject, i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}}, i32 addrspace(1)* noalias %{{.*}}) } diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp index 7e07220663868..5bddd50aa3c13 100644 --- a/clang/test/CodeGenSYCL/kernel-metadata.cpp +++ b/clang/test/CodeGenSYCL/kernel-metadata.cpp @@ -1,4 +1,6 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s +// FIXME: Confirm metadata change +// XFAIL: * // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function() {{[^{]+}} !kernel_arg_addr_space ![[MD:[0-9]+]] !kernel_arg_access_qual ![[MD]] !kernel_arg_type ![[MD]] !kernel_arg_base_type ![[MD]] !kernel_arg_type_qual ![[MD]] // CHECK: ![[MD]] = !{} diff --git a/clang/test/CodeGenSYCL/module-id.cpp b/clang/test/CodeGenSYCL/module-id.cpp index d120ee295c288..9ac45a21110e6 100644 --- a/clang/test/CodeGenSYCL/module-id.cpp +++ b/clang/test/CodeGenSYCL/module-id.cpp @@ -9,6 +9,6 @@ int main() { kernel_single_task([]() {}); return 0; } -// CHECK: define spir_kernel void @{{.*}}kernel{{.*}}() #[[KERN_ATTR:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel{{.*}}(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) #[[KERN_ATTR:[0-9]+]] // CHECK: #[[KERN_ATTR]] = { {{.*}}"sycl-module-id"="{{.*}}module-id.cpp"{{.*}} } diff --git a/clang/test/CodeGenSYCL/num-simd-work-items.cpp b/clang/test/CodeGenSYCL/num-simd-work-items.cpp index 8b8b8ba22d0da..03121a0061d8c 100644 --- a/clang/test/CodeGenSYCL/num-simd-work-items.cpp +++ b/clang/test/CodeGenSYCL/num-simd-work-items.cpp @@ -19,8 +19,7 @@ void bar() { } -// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !num_simd_work_items ![[NUM1:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !num_simd_work_items ![[NUM42:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name1(%class.{{.*}}.Foo* byval(%class.{{.*}}.Foo) align 1 %_arg_kernelObject) {{.*}} !num_simd_work_items ![[NUM1:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name2(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) {{.*}} !num_simd_work_items ![[NUM42:[0-9]+]] // CHECK: ![[NUM1]] = !{i32 1} // CHECK: ![[NUM42]] = !{i32 42} - diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp index f290ca4757c1d..3f2bc52889866 100644 --- a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp @@ -30,10 +30,9 @@ void bar() { []() [[cl::intel_reqd_sub_group_size(4)]] {}); } -// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !intel_reqd_sub_group_size ![[SGSIZE16:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !intel_reqd_sub_group_size ![[SGSIZE8:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name3() {{.*}} !intel_reqd_sub_group_size ![[SGSIZE4:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name1(%class.{{.*}}.Functor16* byval(%class.{{.*}}.Functor16) align 1 %_arg_kernelObject) {{.*}} !intel_reqd_sub_group_size ![[SGSIZE16:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name2(%class.{{.*}}.Functor* byval(%class.{{.*}}.Functor) align 1 %_arg_kernelObject) {{.*}} !intel_reqd_sub_group_size ![[SGSIZE8:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name3(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) {{.*}} !intel_reqd_sub_group_size ![[SGSIZE4:[0-9]+]] // CHECK: ![[SGSIZE16]] = !{i32 16} // CHECK: ![[SGSIZE8]] = !{i32 8} // CHECK: ![[SGSIZE4]] = !{i32 4} - diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp index bfb08c7ce6c2d..18a3a332d87dc 100644 --- a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp @@ -30,9 +30,9 @@ void bar() { []() [[cl::reqd_work_group_size(8, 8, 8)]]{}); } -// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !reqd_work_group_size ![[WGSIZE32:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !reqd_work_group_size ![[WGSIZE8:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name3() {{.*}} !reqd_work_group_size ![[WGSIZE88:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name1(%class.{{.*}}.Functor32x16x16* byval(%class.{{.*}}.Functor32x16x16) align 1 %_arg_kernelObject) {{.*}} !reqd_work_group_size ![[WGSIZE32:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name2(%class.{{.*}}.Functor* byval(%class.{{.*}}.Functor) align 1 %_arg_kernelObject) {{.*}} !reqd_work_group_size ![[WGSIZE8:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name3(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) {{.*}} !reqd_work_group_size ![[WGSIZE88:[0-9]+]] // CHECK: ![[WGSIZE32]] = !{i32 16, i32 16, i32 32} // CHECK: ![[WGSIZE8]] = !{i32 1, i32 1, i32 8} // CHECK: ![[WGSIZE88]] = !{i32 8, i32 8, i32 8} diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp index 947a650afea12..bd56dcb33a258 100644 --- a/clang/test/CodeGenSYCL/sampler.cpp +++ b/clang/test/CodeGenSYCL/sampler.cpp @@ -1,5 +1,5 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -I %S/Inputs -disable-llvm-passes -emit-llvm %s -o - | FileCheck --enable-var-scope %s -// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]]) +// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 8 %_arg_kernelObject, %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]]) // CHECK-NEXT: entry: // CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8 // CHECK-NEXT: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 8 diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp index bed31dcb96e48..47e06851714e9 100644 --- a/clang/test/CodeGenSYCL/spir-calling-conv.cpp +++ b/clang/test/CodeGenSYCL/spir-calling-conv.cpp @@ -7,9 +7,11 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { int main() { - // CHECK: define spir_kernel void @_ZTSZ4mainE15kernel_function() + // CHECK: define spir_kernel void @_ZTSZ4mainE15kernel_function(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) - // CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %2) + // CHECK: [[CAST:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon"* %_arg_kernelObject to %"class.{{.*}}.anon" addrspace(4)* + + // CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* [[CAST]]) // CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon" addrspace(4)* %this) diff --git a/clang/test/CodeGenSYCL/spir-enum.cpp b/clang/test/CodeGenSYCL/spir-enum.cpp index 738d1337e02ab..7e21dde3be247 100644 --- a/clang/test/CodeGenSYCL/spir-enum.cpp +++ b/clang/test/CodeGenSYCL/spir-enum.cpp @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// FIXME: What is this test checking? Is it required now that we're passing kernel object directly? +// XFAIL: * template __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { kernelFunc(); @@ -20,7 +22,7 @@ void test(enum_type val) int main() { - // CHECK: define spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 %_arg_) + // CHECK: define spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 4 %_arg_kernelObject, i32 %_arg_) // CHECK: getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* // CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %4) diff --git a/clang/test/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp index 8e6fbcec309dd..31d0909f2b01f 100644 --- a/clang/test/CodeGenSYCL/struct_kernel_param.cpp +++ b/clang/test/CodeGenSYCL/struct_kernel_param.cpp @@ -3,12 +3,12 @@ // CHECK: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 36, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 12 }, // CHECK-EMPTY: // CHECK-NEXT:}; - // This test checks if compiler accepts structures as kernel parameters. #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp index 3df365b3fef10..b72b19539e275 100644 --- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp +++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp @@ -16,6 +16,6 @@ void bar() { kernel(foo); } -// CHECK: define spir_kernel void @{{.*}}kernel_name() {{.*}} !reqd_work_group_size ![[WGSIZE:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name(%class.{{.*}}.Functor* byval(%class.{{.*}}.Functor) align 1 %_arg_kernelObject) {{.*}} !reqd_work_group_size ![[WGSIZE:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE:[0-9]+]] // CHECK: ![[WGSIZE]] = !{i32 16, i32 16, i32 32} // CHECK: ![[SGSIZE]] = !{i32 4} diff --git a/clang/test/CodeGenSYCL/usm-int-header.cpp b/clang/test/CodeGenSYCL/usm-int-header.cpp index e3cedd5302b78..a7df2be5bf4a5 100644 --- a/clang/test/CodeGenSYCL/usm-int-header.cpp +++ b/clang/test/CodeGenSYCL/usm-int-header.cpp @@ -31,7 +31,7 @@ int main() { }); } -// CHECK: FunctionDecl {{.*}}usm_test 'void (__global int *, __global float *)' +// CHECK: FunctionDecl {{.*}}usm_test 'void ((lambda at {{.*}}usm-int-header.cpp{{.*}}), __global int *, __global float *)' // TODO: SYCL specific fail - analyze and enable // XFAIL: windows-msvc diff --git a/clang/test/CodeGenSYCL/wrapped-accessor.cpp b/clang/test/CodeGenSYCL/wrapped-accessor.cpp index 0cd651efc58f5..920b4187dfc38 100644 --- a/clang/test/CodeGenSYCL/wrapped-accessor.cpp +++ b/clang/test/CodeGenSYCL/wrapped-accessor.cpp @@ -18,6 +18,7 @@ // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- _ZTSZ4mainE14wrapped_access // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 12, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 12, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 }, // CHECK-EMPTY: // CHECK-NEXT: }; diff --git a/clang/test/SemaSYCL/accessors-targets-image.cpp b/clang/test/SemaSYCL/accessors-targets-image.cpp index 18fac9940cb1f..938156c38d006 100644 --- a/clang/test/SemaSYCL/accessors-targets-image.cpp +++ b/clang/test/SemaSYCL/accessors-targets-image.cpp @@ -63,12 +63,12 @@ int main() { }); } -// CHECK: {{.*}}use_image1d_r 'void (__read_only image1d_t)' -// CHECK: {{.*}}use_image2d_r 'void (__read_only image2d_t)' -// CHECK: {{.*}}use_image3d_r 'void (__read_only image3d_t)' -// CHECK: {{.*}}use_image1d_w 'void (__write_only image1d_t)' -// CHECK: {{.*}}use_image2d_w 'void (__write_only image2d_t)' -// CHECK: {{.*}}use_image3d_w 'void (__write_only image3d_t)' +// CHECK: {{.*}}use_image1d_r 'void ((lambda at {{.*}}accessors-targets-image.cpp{{.*}}), __read_only image1d_t)' +// CHECK: {{.*}}use_image2d_r 'void ((lambda at {{.*}}accessors-targets-image.cpp{{.*}}), __read_only image2d_t)' +// CHECK: {{.*}}use_image3d_r 'void ((lambda at {{.*}}accessors-targets-image.cpp{{.*}}), __read_only image3d_t)' +// CHECK: {{.*}}use_image1d_w 'void ((lambda at {{.*}}accessors-targets-image.cpp{{.*}}), __write_only image1d_t)' +// CHECK: {{.*}}use_image2d_w 'void ((lambda at {{.*}}accessors-targets-image.cpp{{.*}}), __write_only image2d_t)' +// CHECK: {{.*}}use_image3d_w 'void ((lambda at {{.*}}accessors-targets-image.cpp{{.*}}), __write_only image3d_t)' // TODO: SYCL specific fail - analyze and enable // XFAIL: windows-msvc diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp index dbaab2664e95c..0869f0a884b73 100644 --- a/clang/test/SemaSYCL/accessors-targets.cpp +++ b/clang/test/SemaSYCL/accessors-targets.cpp @@ -36,6 +36,6 @@ int main() { constant_acc.use(); }); } -// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' -// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' -// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_local{{.*}} 'void ((lambda at {{.*}}accessors-targets.cpp{{.*}}), __local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_global{{.*}} 'void ((lambda at {{.*}}accessors-targets.cpp{{.*}}), __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_constant{{.*}} 'void ((lambda at {{.*}}accessors-targets.cpp{{.*}}), __constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp index 1f500eff0a888..56864d1939c55 100644 --- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp @@ -23,10 +23,11 @@ int main() { // Check declaration of the kernel -// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void ((lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}}), __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // Check parameters of the kernel +// CHECK: ParmVarDecl {{.*}} used [[_arg_KernelObject:[0-9a-zA-Z_]+]] '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' // CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global int *' // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp index 82cd21bf01552..4ffbe632971f1 100644 --- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp +++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp @@ -45,47 +45,32 @@ int main() { return 0; } // Check kernel parameters -// CHECK: FunctionDecl {{.*}}kernel_const{{.*}} 'void (const int)' -// CHECK: ParmVarDecl {{.*}} used _arg_ 'const int' - -// Check that lambda field of const built-in type is initialized -// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' -// CHECK-NEXT: InitListExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'const int' lvalue ParmVar {{.*}} '_arg_' 'const int' +// CHECK: FunctionDecl {{.*}}kernel_const{{.*}} 'void ((lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}}))' +// CHECK: ParmVarDecl {{.*}} used _arg_kernelObject '(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}}) // Check kernel parameters -// CHECK: {{.*}}kernel_int{{.*}} 'void (int)' -// CHECK: ParmVarDecl {{.*}} used _arg_ 'int' - -// Check that lambda field of built-in type is initialized -// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' -// CHECK-NEXT: InitListExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' +// CHECK: {{.*}}kernel_int{{.*}} 'void ((lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}}))' +// CHECK: ParmVarDecl {{.*}} used _arg_kernelObject '(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}}) // Check kernel parameters -// CHECK: {{.*}}kernel_struct{{.*}} 'void (test_struct)' -// CHECK: ParmVarDecl {{.*}} used _arg_ 'test_struct' - -// Check that lambda field of struct type is initialized -// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' -// CHECK-NEXT: InitListExpr -// CHECK-NEXT: CXXConstructExpr {{.*}}'test_struct'{{.*}}void (const test_struct &) -// CHECK-NEXT: ImplicitCastExpr {{.*}}'const test_struct' lvalue -// CHECK-NEXT: DeclRefExpr {{.*}} 'test_struct' lvalue ParmVar {{.*}} '_arg_' 'test_struct' +// CHECK: {{.*}}kernel_struct{{.*}} 'void ((lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}}))' +// CHECK: ParmVarDecl {{.*}} used _arg_kernelObject '(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}}) // Check kernel parameters -// CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *)' +// CHECK: {{.*}}kernel_pointer{{.*}} 'void ((lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}}), __global int *, __global int *)' +// CHECK: ParmVarDecl {{.*}} used _arg_kernelObject '(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}}) // CHECK: ParmVarDecl {{.*}} used _arg_ '__global int *' // CHECK: ParmVarDecl {{.*}} used _arg_ '__global int *' -// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' -// Check that lambda fields of pointer types are initialized -// CHECK: InitListExpr +// Check that lambda fields of pointer types are assigned with kernel pointer parameters. +// CHECK: BinaryOperator {{.*}} '=' +// CHECK-NEXT: MemberExpr {{.*}} 'int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' lvalue ParmVar {{.*}} '_arg_kernelObject' '(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' // CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *' + +// CHECK: BinaryOperator {{.*}} '=' +// CHECK-NEXT: MemberExpr {{.*}} 'int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' lvalue ParmVar {{.*}} '_arg_kernelObject' '(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' // CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *' - -// Check kernel parameters diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp index 24d36a6ba54b6..0ac434172e165 100644 --- a/clang/test/SemaSYCL/fake-accessors.cpp +++ b/clang/test/SemaSYCL/fake-accessors.cpp @@ -51,6 +51,6 @@ int main() { }); return 0; } -// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: fake_accessors{{.*}} 'void ((lambda at {{.*}}fake-accessors.cpp{{.*}}), __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_typedef{{.*}} 'void ((lambda at {{.*}}fake-accessors.cpp{{.*}}), __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_alias{{.*}} 'void ((lambda at {{.*}}fake-accessors.cpp{{.*}}), __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) diff --git a/clang/test/SemaSYCL/sampler.cpp b/clang/test/SemaSYCL/sampler.cpp index c9f4a5bbfdfcc..c558768e4c2f8 100644 --- a/clang/test/SemaSYCL/sampler.cpp +++ b/clang/test/SemaSYCL/sampler.cpp @@ -16,9 +16,10 @@ int main() { } // Check declaration of the test kernel -// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (sampler_t)' +// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void ((lambda at {{.*}}sampler.cpp{{.*}}), sampler_t)' // // Check parameters of the test kernel +// CHECK: ParmVarDecl {{.*}} used _arg_kernelObject '(lambda at {{.*}}sampler.cpp{{.*}})' // CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] 'sampler_t' // // Check that sampler field of the test kernel object is initialized using __init method diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index 83bb3ff2448fb..62869936dc430 100644 --- a/clang/test/SemaSYCL/wrapped-accessor.cpp +++ b/clang/test/SemaSYCL/wrapped-accessor.cpp @@ -23,9 +23,10 @@ int main() { } // Check declaration of the kernel -// CHECK: wrapped_access{{.*}} 'void (AccWrapper>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: wrapped_access{{.*}} 'void ((lambda at {{.*}}wrapped-accessor.cpp{{.*}}), AccWrapper>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // Check parameters of the kernel +// CHECK: ParmVarDecl {{.*}} used _arg_kernelObject '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' // CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper>':'AccWrapper>' // CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *' // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'