diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index 3b9000133438e..5d7f0976bcd86 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -114,6 +114,16 @@ class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo { return CC_SpirFunction; } + llvm::Optional getConstantAddressSpace() const override { + // If we assign "opencl_constant" address space the following code becomes + // illegal, because it can't be cast to any other address space: + // + // const char *getLiteral() { + // return "AB"; + // } + return LangAS::opencl_global; + } + void setSupportedOpenCLOpts() override { // Assume all OpenCL extensions and optional core features are supported // for SPIR since it is a generic target. diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 0bdf127a4b156..8e758730d3309 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -4771,17 +4771,6 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, V->getType()->isIntegerTy()) V = Builder.CreateZExt(V, ArgInfo.getCoerceToType()); - if (FirstIRArg < IRFuncTy->getNumParams()) { - const auto *LHSPtrTy = - dyn_cast_or_null(V->getType()); - const auto *RHSPtrTy = dyn_cast_or_null( - IRFuncTy->getParamType(FirstIRArg)); - if (LHSPtrTy && RHSPtrTy && - LHSPtrTy->getAddressSpace() != RHSPtrTy->getAddressSpace()) - V = Builder.CreateAddrSpaceCast(V, - IRFuncTy->getParamType(FirstIRArg)); - } - // If the argument doesn't match, perform a bitcast to coerce it. This // can happen due to trivial type mismatches. if (FirstIRArg < IRFuncTy->getNumParams() && @@ -5002,20 +4991,6 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo, if (!CallArgs.getCleanupsToDeactivate().empty()) deactivateArgCleanupsBeforeCall(*this, CallArgs); - // Addrspace cast to generic if necessary - for (unsigned i = 0; i < IRFuncTy->getNumParams(); ++i) { - if (auto *PtrTy = dyn_cast(IRCallArgs[i]->getType())) { - auto *ExpectedPtrType = - cast(IRFuncTy->getParamType(i)); - unsigned ValueAS = PtrTy->getAddressSpace(); - unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); - if (ValueAS != ExpectedAS) { - IRCallArgs[i] = Builder.CreatePointerBitCastOrAddrSpaceCast( - IRCallArgs[i], ExpectedPtrType); - } - } - } - // Assert that the arguments we computed match up. The IR verifier // will catch this, but this is a common enough source of problems // during IRGen changes that it's way better for debugging to catch diff --git a/clang/lib/CodeGen/CGClass.cpp b/clang/lib/CodeGen/CGClass.cpp index 21fe47f600afa..ba221dbbc83b7 100644 --- a/clang/lib/CodeGen/CGClass.cpp +++ b/clang/lib/CodeGen/CGClass.cpp @@ -342,7 +342,7 @@ Address CodeGenFunction::GetAddressOfBaseClass( EmitTypeCheck(TCK_Upcast, Loc, Value.getPointer(), DerivedTy, DerivedAlign, SkippedChecks); } - return Builder.CreatePointerBitCastOrAddrSpaceCast(Value, BasePtrTy); + return Builder.CreateBitCast(Value, BasePtrTy); } llvm::BasicBlock *origBB = nullptr; diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp index 683be93b4b873..bd34195a5d7da 100644 --- a/clang/lib/CodeGen/CGDecl.cpp +++ b/clang/lib/CodeGen/CGDecl.cpp @@ -1642,8 +1642,12 @@ CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) { CGM.generateIntelFPGAAnnotation(&D, AnnotStr); if (!AnnotStr.empty()) { llvm::Value *V = address.getPointer(); - EmitAnnotationCall(CGM.getIntrinsic(llvm::Intrinsic::var_annotation), - Builder.CreateBitCast(V, CGM.Int8PtrTy, V->getName()), + llvm::Type *DestPtrTy = llvm::PointerType::getInt8PtrTy( + CGM.getLLVMContext(), address.getAddressSpace()); + llvm::Value *Arg = Builder.CreateBitCast(V, DestPtrTy, V->getName()); + if (address.getAddressSpace() != 0) + Arg = Builder.CreateAddrSpaceCast(Arg, CGM.Int8PtrTy, V->getName()); + EmitAnnotationCall(CGM.getIntrinsic(llvm::Intrinsic::var_annotation), Arg, AnnotStr, D.getLocation()); } } diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index af1b79dc6a15f..e752aad5aa97e 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -1125,8 +1125,10 @@ Address CodeGenFunction::EmitPointerWithAlignment(const Expr *E, CodeGenFunction::CFITCK_UnrelatedCast, CE->getBeginLoc()); } - return Builder.CreatePointerBitCastOrAddrSpaceCast( - Addr, ConvertType(E->getType())); + return CE->getCastKind() != CK_AddressSpaceConversion + ? Builder.CreateBitCast(Addr, ConvertType(E->getType())) + : Builder.CreateAddrSpaceCast(Addr, + ConvertType(E->getType())); } break; @@ -1855,16 +1857,6 @@ void CodeGenFunction::EmitStoreOfScalar(llvm::Value *Value, Address Addr, return; } - if (auto *PtrTy = dyn_cast(Value->getType())) { - auto *ExpectedPtrType = - cast(Addr.getType()->getElementType()); - unsigned ValueAS = PtrTy->getAddressSpace(); - unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); - if (ValueAS != ExpectedAS) { - Value = - Builder.CreatePointerBitCastOrAddrSpaceCast(Value, ExpectedPtrType); - } - } llvm::StoreInst *Store = Builder.CreateStore(Value, Addr, Volatile); if (isNontemporal) { llvm::MDNode *Node = @@ -4614,39 +4606,14 @@ EmitConditionalOperatorLValue(const AbstractConditionalOperator *expr) { EmitBlock(contBlock); if (lhs && rhs) { - llvm::Value *lhsPtr = lhs->getPointer(*this); - llvm::Value *rhsPtr = rhs->getPointer(*this); - if (rhsPtr->getType() != lhsPtr->getType()) { - if (!getLangOpts().SYCLIsDevice) - llvm_unreachable( - "Unable to find a common address space for two pointers."); - - auto CastToAS = [](llvm::Value *V, llvm::BasicBlock *BB, unsigned AS) { - auto *Ty = cast(V->getType()); - if (Ty->getAddressSpace() == AS) - return V; - llvm::IRBuilder<> Builder(BB->getTerminator()); - auto *TyAS = llvm::PointerType::get(Ty->getElementType(), AS); - return Builder.CreatePointerBitCastOrAddrSpaceCast(V, TyAS); - }; - - // Language rules define if it is legal to cast from one address space - // to another, and which address space we should use as a "common - // denominator". In SYCL, generic address space overlaps with all other - // address spaces. - unsigned GenericAS = - getContext().getTargetAddressSpace(LangAS::opencl_generic); - - lhsPtr = CastToAS(lhsPtr, lhsBlock, GenericAS); - rhsPtr = CastToAS(rhsPtr, rhsBlock, GenericAS); - } - llvm::PHINode *phi = Builder.CreatePHI(lhsPtr->getType(), 2, "cond-lvalue"); - phi->addIncoming(lhsPtr, lhsBlock); - phi->addIncoming(rhsPtr, rhsBlock); + llvm::PHINode *phi = + Builder.CreatePHI(lhs->getPointer(*this)->getType(), 2, "cond-lvalue"); + phi->addIncoming(lhs->getPointer(*this), lhsBlock); + phi->addIncoming(rhs->getPointer(*this), rhsBlock); Address result(phi, std::min(lhs->getAlignment(), rhs->getAlignment())); AlignmentSource alignSource = - std::max(lhs->getBaseInfo().getAlignmentSource(), - rhs->getBaseInfo().getAlignmentSource()); + std::max(lhs->getBaseInfo().getAlignmentSource(), + rhs->getBaseInfo().getAlignmentSource()); TBAAAccessInfo TBAAInfo = CGM.mergeTBAAInfoForConditionalOperator( lhs->getTBAAInfo(), rhs->getTBAAInfo()); return MakeAddrLValue(result, expr->getType(), LValueBaseInfo(alignSource), diff --git a/clang/lib/CodeGen/CGExprAgg.cpp b/clang/lib/CodeGen/CGExprAgg.cpp index a8bc1a2b93a88..2e68103b350bf 100644 --- a/clang/lib/CodeGen/CGExprAgg.cpp +++ b/clang/lib/CodeGen/CGExprAgg.cpp @@ -498,13 +498,7 @@ void AggExprEmitter::EmitArrayInit(Address DestPtr, llvm::ArrayType *AType, elementType.isTriviallyCopyableType(CGF.getContext())) { CodeGen::CodeGenModule &CGM = CGF.CGM; ConstantEmitter Emitter(CGF); - LangAS AS = ArrayQTy.getAddressSpace(); - if (CGM.getLangOpts().SYCLIsDevice && AS == LangAS::Default) { - // SYCL's default AS is 'generic', which can't be used to define constant - // initializer data in. It is reasonable to keep it in the same AS - // as string literals. - AS = CGM.getStringLiteralAddressSpace(); - } + LangAS AS = CGM.GetGlobalVarAddressSpace(/*VarDecl= */ nullptr); if (llvm::Constant *C = Emitter.tryEmitForInitializer(E, AS, ArrayQTy)) { auto GV = new llvm::GlobalVariable( CGM.getModule(), C->getType(), diff --git a/clang/lib/CodeGen/CGExprScalar.cpp b/clang/lib/CodeGen/CGExprScalar.cpp index ece7969786196..6f7e8263fa101 100644 --- a/clang/lib/CodeGen/CGExprScalar.cpp +++ b/clang/lib/CodeGen/CGExprScalar.cpp @@ -1964,26 +1964,10 @@ Value *ScalarExprEmitter::VisitCastExpr(CastExpr *CE) { Value *Src = Visit(const_cast(E)); llvm::Type *SrcTy = Src->getType(); llvm::Type *DstTy = ConvertType(DestTy); - bool NeedAddrspaceCast = false; if (SrcTy->isPtrOrPtrVectorTy() && DstTy->isPtrOrPtrVectorTy() && SrcTy->getPointerAddressSpace() != DstTy->getPointerAddressSpace()) { - // If we have the same address space in AST, which is then codegen'ed to - // different address spaces in IR, then an address space cast should be - // valid. - // - // This is the case for SYCL, where both types have Default address space - // in AST, but in IR one of them may be in opencl_private, and another in - // opencl_generic address space: - // - // int arr[5]; // automatic variable, default AS in AST, - // // private AS in IR - // - // char* p = arr; // default AS in AST, generic AS in IR - // - if (E->getType().getAddressSpace() != DestTy.getAddressSpace()) - llvm_unreachable("wrong cast for pointers in different address spaces" - "(must be an address space cast)!"); - NeedAddrspaceCast = true; + llvm_unreachable("wrong cast for pointers in different address spaces" + "(must be an address space cast)!"); } if (CGF.SanOpts.has(SanitizerKind::CFIUnrelatedCast)) { @@ -2022,14 +2006,6 @@ Value *ScalarExprEmitter::VisitCastExpr(CastExpr *CE) { } } - if (NeedAddrspaceCast) { - llvm::Type *SrcPointeeTy = Src->getType()->getPointerElementType(); - llvm::Type *SrcNewAS = llvm::PointerType::get( - SrcPointeeTy, cast(DstTy)->getAddressSpace()); - - Src = Builder.CreateAddrSpaceCast(Src, SrcNewAS); - } - // If Src is a fixed vector and Dst is a scalable vector, and both have the // same element type, use the llvm.experimental.vector.insert intrinsic to // perform the bitcast. @@ -2966,53 +2942,6 @@ Value *ScalarExprEmitter::VisitUnaryImag(const UnaryOperator *E) { // Binary Operators //===----------------------------------------------------------------------===// -static Value *insertAddressSpaceCast(Value *V, unsigned NewAS) { - auto *VTy = cast(V->getType()); - if (VTy->getAddressSpace() == NewAS) - return V; - - llvm::PointerType *VTyNewAS = - llvm::PointerType::get(VTy->getElementType(), NewAS); - - if (auto *Constant = dyn_cast(V)) - return llvm::ConstantExpr::getAddrSpaceCast(Constant, VTyNewAS); - - llvm::Instruction *NewV = - new llvm::AddrSpaceCastInst(V, VTyNewAS, V->getName() + ".ascast"); - NewV->insertAfter(cast(V)); - return NewV; -} - -static void ensureSameAddrSpace(Value *&RHS, Value *&LHS, - bool CanInsertAddrspaceCast, - const LangOptions &Opts, - const ASTContext &Context) { - if (RHS->getType() == LHS->getType()) - return; - - auto *RHSTy = dyn_cast(RHS->getType()); - auto *LHSTy = dyn_cast(LHS->getType()); - if (!RHSTy || !LHSTy || RHSTy->getAddressSpace() == LHSTy->getAddressSpace()) - return; - - if (!CanInsertAddrspaceCast) - // Pointers have different address spaces and we cannot do anything with - // this. - llvm_unreachable("Pointers are expected to have the same address space."); - - // Language rules define if it is legal to cast from one address space to - // another, and which address space we should use as a "common - // denominator". In SYCL, generic address space overlaps with all other - // address spaces. - if (Opts.SYCLIsDevice) { - unsigned GenericAS = Context.getTargetAddressSpace(LangAS::opencl_generic); - RHS = insertAddressSpaceCast(RHS, GenericAS); - LHS = insertAddressSpaceCast(LHS, GenericAS); - } else - llvm_unreachable("Unable to find a common address space for " - "two pointers."); -} - BinOpInfo ScalarExprEmitter::EmitBinOps(const BinaryOperator *E) { TestAndClearIgnoreResultAssign(); BinOpInfo Result; @@ -4134,14 +4063,6 @@ Value *ScalarExprEmitter::EmitCompare(const BinaryOperator *E, RHS = Builder.CreateStripInvariantGroup(RHS); } - // Expression operands may have the same addrspace in AST, but different - // addrspaces in LLVM IR, in which case an addrspacecast should be valid. - bool CanInsertAddrspaceCast = - LHSTy.getAddressSpace() == RHSTy.getAddressSpace(); - - ensureSameAddrSpace(RHS, LHS, CanInsertAddrspaceCast, CGF.getLangOpts(), - CGF.getContext()); - Result = Builder.CreateICmp(UICmpOpc, LHS, RHS, "cmp"); } @@ -4513,6 +4434,7 @@ static bool isCheapEnoughToEvaluateUnconditionally(const Expr *E, // exist in the source-level program. } + Value *ScalarExprEmitter:: VisitAbstractConditionalOperator(const AbstractConditionalOperator *E) { TestAndClearIgnoreResultAssign(); @@ -4621,15 +4543,6 @@ VisitAbstractConditionalOperator(const AbstractConditionalOperator *E) { assert(!RHS && "LHS and RHS types must match"); return nullptr; } - - // Expressions may have the same addrspace in AST, but different address - // space in LLVM IR, in which case an addrspacecast should be valid. - bool CanInsertAddrspaceCast = rhsExpr->getType().getAddressSpace() == - lhsExpr->getType().getAddressSpace(); - - ensureSameAddrSpace(RHS, LHS, CanInsertAddrspaceCast, CGF.getLangOpts(), - CGF.getContext()); - return Builder.CreateSelect(CondV, LHS, RHS, "cond"); } @@ -4664,14 +4577,6 @@ VisitAbstractConditionalOperator(const AbstractConditionalOperator *E) { if (!RHS) return LHS; - // Expressions may have the same addrspace in AST, but different address - // space in LLVM IR, in which case an addrspacecast should be valid. - bool CanInsertAddrspaceCast = rhsExpr->getType().getAddressSpace() == - lhsExpr->getType().getAddressSpace(); - - ensureSameAddrSpace(RHS, LHS, CanInsertAddrspaceCast, CGF.getLangOpts(), - CGF.getContext()); - // Create a PHI node for the real part. llvm::PHINode *PN = Builder.CreatePHI(LHS->getType(), 2, "cond"); PN->addIncoming(LHS, LHSBlock); diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp index f5b16fb842a45..a1a72a9f668d2 100644 --- a/clang/lib/CodeGen/CGStmt.cpp +++ b/clang/lib/CodeGen/CGStmt.cpp @@ -1211,35 +1211,12 @@ void CodeGenFunction::EmitReturnStmt(const ReturnStmt &S) { // If this function returns a reference, take the address of the expression // rather than the value. RValue Result = EmitReferenceBindingToExpr(RV); - llvm::Value *Val = Result.getScalarVal(); - if (auto *PtrTy = dyn_cast(Val->getType())) { - auto *ExpectedPtrType = - cast(ReturnValue.getType()->getElementType()); - unsigned ValueAS = PtrTy->getAddressSpace(); - unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); - if (ValueAS != ExpectedAS) { - Val = Builder.CreatePointerBitCastOrAddrSpaceCast(Val, ExpectedPtrType); - } - } - Builder.CreateStore(Val, ReturnValue); + Builder.CreateStore(Result.getScalarVal(), ReturnValue); } else { switch (getEvaluationKind(RV->getType())) { case TEK_Scalar: - { - llvm::Value *Val = EmitScalarExpr(RV); - if (auto *PtrTy = dyn_cast(Val->getType())) { - auto *ExpectedPtrType = - cast(ReturnValue.getType()->getElementType()); - unsigned ValueAS = PtrTy->getAddressSpace(); - unsigned ExpectedAS = ExpectedPtrType->getAddressSpace(); - if (ValueAS != ExpectedAS) { - Val = - Builder.CreatePointerBitCastOrAddrSpaceCast(Val, ExpectedPtrType); - } - } - Builder.CreateStore(Val, ReturnValue); + Builder.CreateStore(EmitScalarExpr(RV), ReturnValue); break; - } case TEK_Complex: EmitComplexExprIntoLValue(RV, MakeAddrLValue(ReturnValue, RV->getType()), /*isInit*/ true); diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 3c49cfe9e8786..2142596ed4b75 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3991,7 +3991,6 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, : (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default); assert(getContext().getTargetAddressSpace(ExpectedAS) == Ty->getPointerAddressSpace()); - if (AddrSpace != ExpectedAS) return getTargetCodeGenInfo().performAddrSpaceCast(*this, GV, AddrSpace, ExpectedAS, Ty); @@ -4149,14 +4148,10 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { return AddrSpace; } - if (LangOpts.SYCLIsDevice) { - if (!D) - return LangAS::opencl_global; + if (LangOpts.SYCLIsDevice && D) { auto *Scope = D->getAttr(); if (Scope && Scope->isWorkGroup()) return LangAS::opencl_local; - if (D->getType().getAddressSpace() == LangAS::Default) - return LangAS::opencl_global; } if (LangOpts.CUDA && LangOpts.CUDAIsDevice) { @@ -4184,17 +4179,6 @@ LangAS CodeGenModule::getStringLiteralAddressSpace() const { // OpenCL v1.2 s6.5.3: a string literal is in the constant address space. if (LangOpts.OpenCL) return LangAS::opencl_constant; - if (LangOpts.SYCLIsDevice) - // If we keep a literal string in constant address space, the following code - // becomes illegal: - // - // const char *getLiteral() n{ - // return "AB"; - // } - // Use global address space to avoid illegal casts from constant to generic. - // Private address space is not used here because in SPIR-V global values - // cannot have private address space. - return LangAS::opencl_global; if (auto AS = getTarget().getConstantAddressSpace()) return AS.getValue(); return LangAS::Default; @@ -5086,7 +5070,8 @@ void CodeGenModule::EmitAliasDefinition(GlobalDecl GD) { LT = getFunctionLinkage(GD); AS = Aliasee->getType()->getPointerAddressSpace(); } else { - AS = ArgInfoAddressSpace(GetGlobalVarAddressSpace(/*D=*/nullptr)); + const auto *VarD = cast(GD.getDecl()); + AS = ArgInfoAddressSpace(GetGlobalVarAddressSpace(VarD)); Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(), DeclTy->getPointerTo(AS), /*D=*/nullptr); if (const auto *VD = dyn_cast(GD.getDecl())) diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index da450282de421..dd99cf2dac68b 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -9949,6 +9949,15 @@ class SPIRTargetCodeGenInfo : public TargetCodeGenInfo { SPIRTargetCodeGenInfo(CodeGen::CodeGenTypes &CGT) : TargetCodeGenInfo(std::make_unique(CGT)) {} unsigned getOpenCLKernelCallingConv() const override; + + LangAS getASTAllocaAddressSpace() const override { + return getLangASFromTargetAS( + getABIInfo().getDataLayout().getAllocaAddrSpace()); + } + + LangAS getGlobalVarAddressSpace(CodeGenModule &CGM, + const VarDecl *D) const override; + bool shouldEmitStaticExternCAliases() const override; }; } // End anonymous namespace. @@ -9970,6 +9979,34 @@ unsigned SPIRTargetCodeGenInfo::getOpenCLKernelCallingConv() const { return llvm::CallingConv::SPIR_KERNEL; } +bool SPIRTargetCodeGenInfo::shouldEmitStaticExternCAliases() const { + return false; +} + +LangAS SPIRTargetCodeGenInfo::getGlobalVarAddressSpace(CodeGenModule &CGM, + const VarDecl *D) const { + assert(!CGM.getLangOpts().OpenCL && + !(CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice) && + "Address space agnostic languages only"); + LangAS DefaultGlobalAS = getLangASFromTargetAS( + CGM.getContext().getTargetAddressSpace(LangAS::opencl_global)); + if (!D) + return DefaultGlobalAS; + + LangAS AddrSpace = D->getType().getAddressSpace(); + assert(AddrSpace == LangAS::Default || isTargetAddressSpace(AddrSpace) || + // allow applying clang AST address spaces in SYCL mode + (CGM.getLangOpts().SYCL && CGM.getLangOpts().SYCLIsDevice)); + if (AddrSpace != LangAS::Default) + return AddrSpace; + + if (CGM.isTypeConstant(D->getType(), false)) { + if (auto ConstAS = CGM.getTarget().getConstantAddressSpace()) + return ConstAS.getValue(); + } + return DefaultGlobalAS; +} + static bool appendType(SmallStringEnc &Enc, QualType QType, const CodeGen::CodeGenModule &CGM, TypeStringCache &TSC); diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp index f0196d8be0e8e..2a421b1389ab3 100644 --- a/clang/test/CodeGenSYCL/accessor_inheritance.cpp +++ b/clang/test/CodeGenSYCL/accessor_inheritance.cpp @@ -36,59 +36,62 @@ int main() { // CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC2_ID:%[a-zA-Z0-9_]+]], // CHECK-SAME: i32 [[ARG_C:%[a-zA-Z0-9_]+]]) -// Allocas for kernel parameters +// Allocas and addrspacecasts for kernel parameters // CHECK: [[ARG_A]].addr = alloca i32 +// CHECK: [[ARG_A]].addr.ascast = addrspacecast i32* [[ARG_A]].addr to i32 addrspace(4)* // CHECK: [[ARG_B]].addr = alloca i32 +// CHECK: [[ARG_B]].addr.ascast = addrspacecast i32* [[ARG_B]].addr to i32 addrspace(4)* // CHECK: [[ACC1_DATA]].addr = alloca i8 addrspace(1)* +// CHECK: [[ACC1_DATA]].addr.ascast = addrspacecast i8 addrspace(1)** [[ACC1_DATA]].addr to i8 addrspace(1)* addrspace(4)* // CHECK: [[ACC2_DATA]].addr = alloca i8 addrspace(1)* +// CHECK: [[ACC2_DATA]].addr.ascast = addrspacecast i8 addrspace(1)** [[ACC2_DATA]].addr to i8 addrspace(1)* addrspace(4)* // CHECK: [[ARG_C]].addr = alloca i32 +// CHECK: [[ARG_C]].addr.ascast = addrspacecast i32* [[ARG_C]].addr to i32 addrspace(4)* // // Lambda object alloca -// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = alloca %"class.{{.*}}.anon" +// CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %"class{{.*}}.anon" +// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}.anon"* [[KERNEL]] to %"class{{.*}}.anon" addrspace(4)* // // Kernel argument stores -// CHECK: store i32 [[ARG_A]], i32* [[ARG_A]].addr -// CHECK: store i32 [[ARG_B]], i32* [[ARG_B]].addr -// CHECK: store i8 addrspace(1)* [[ACC1_DATA]], i8 addrspace(1)** [[ACC1_DATA]].addr -// CHECK: store i8 addrspace(1)* [[ACC2_DATA]], i8 addrspace(1)** [[ACC2_DATA]].addr -// CHECK: store i32 [[ARG_C]], i32* [[ARG_C]].addr +// CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast +// CHECK: store i32 [[ARG_B]], i32 addrspace(4)* [[ARG_B]].addr.ascast +// CHECK: store i8 addrspace(1)* [[ACC1_DATA]], i8 addrspace(1)* addrspace(4)* [[ACC1_DATA]].addr.ascast +// CHECK: store i8 addrspace(1)* [[ACC2_DATA]], i8 addrspace(1)* addrspace(4)* [[ACC2_DATA]].addr.ascast +// CHECK: store i32 [[ARG_C]], i32 addrspace(4)* [[ARG_C]].addr.ascast // // Check A and B scalar fields initialization -// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0 -// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to %struct{{.*}}Base* -// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 0 -// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_A]].addr -// CHECK: store i32 [[ARG_A_LOAD]], i32* [[FIELD_A]] -// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 1 -// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_B]].addr -// CHECK: store i32 [[ARG_B_LOAD]], i32* [[FIELD_B]] +// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0 +// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to %struct{{.*}}Base addrspace(4)* +// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 0 +// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast +// CHECK: store i32 [[ARG_A_LOAD]], i32 addrspace(4)* [[FIELD_A]] +// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 1 +// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_B]].addr.ascast +// CHECK: store i32 [[ARG_B_LOAD]], i32 addrspace(4)* [[FIELD_B]] // // Check accessors initialization -// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2 -// CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* +// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 2 // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_AS_CAST]]) -// CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8* -// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20 -// CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"* -// CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC_FIELD]]) +// CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to i8 addrspace(4)* +// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8 addrspace(4)* [[BITCAST1]], i64 20 +// CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8 addrspace(4)* [[GEP1]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC2_AS_CAST]]) +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST2]]) // CHECK C field initialization -// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2 -// CHECK: [[ARG_C_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_C]].addr -// CHECK: store i32 [[ARG_C_LOAD]], i32* [[FIELD_C]] +// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured addrspace(4)* [[GEP]], i32 0, i32 2 +// CHECK: [[ARG_C_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_C]].addr.ascast +// CHECK: store i32 [[ARG_C_LOAD]], i32 addrspace(4)* [[FIELD_C]] // // Check __init method calls -// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0 -// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP2]] to %struct{{.*}}Base* -// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST3]], i32 0, i32 2 -// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC1_DATA]].addr -// CHECK: [[ACC1_AS_CAST1:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC1_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_AS_CAST1]], i8 addrspace(1)* [[ACC1_DATA_LOAD]] +// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0 +// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP2]] to %struct{{.*}}Base addrspace(4)* +// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST3]], i32 0, i32 2 +// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC1_DATA]].addr.ascast +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_FIELD]], i8 addrspace(1)* [[ACC1_DATA_LOAD]] // -// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0 -// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC2_DATA]].addr -// CHECK: [[AS_CAST_CAPTURED:%[a-zA-Z0-9_]+]] = addrspacecast %struct{{.*}}Captured* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[AS_CAST_CAPTURED]], i8 addrspace(1)* [[ACC2_DATA_LOAD]] +// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0 +// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC2_DATA]].addr.ascast +// CHECK: [[BITCAST4:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST4]], i8 addrspace(1)* [[ACC2_DATA_LOAD]] diff --git a/clang/test/CodeGenSYCL/address-space-cond-op.cpp b/clang/test/CodeGenSYCL/address-space-cond-op.cpp index b2428dca72dfd..dd65183485d4c 100644 --- a/clang/test/CodeGenSYCL/address-space-cond-op.cpp +++ b/clang/test/CodeGenSYCL/address-space-cond-op.cpp @@ -1,24 +1,36 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py +// RUN: %clang_cc1 -x c++ -triple spir64-unknown-linux-sycldevice -disable-llvm-passes -fsycl -fsycl-is-device -emit-llvm %s -o - | FileCheck %s -// CHECK: [[STYPE:%.+]] = type { i16 } struct S { unsigned short x; }; -S foo(bool cond, S &lhs, S rhs) { -// CHECK-LABEL:@_Z3foobR1SS_ -// CHECK: br i1 {{.+}}, label %[[BTRUE:.+]], label %[[BFALSE:.+]] -// -// CHECK: [[BTRUE]]: -// CHECK: %[[LHS:.+]] = load [[STYPE]] addrspace(4)*, [[STYPE]] addrspace(4)** -// CHECK: br label %[[BEND:.+]] +// CHECK-LABEL: @_Z3foobR1SS_( +// CHECK: entry: +// CHECK-NEXT: [[COND_ADDR:%.*]] = alloca i8, align 1 +// CHECK-NEXT: [[COND_ADDR_ASCAST:%.*]] = addrspacecast i8* [[COND_ADDR]] to i8 addrspace(4)* +// CHECK-NEXT: [[LHS_ADDR:%.*]] = alloca [[STRUCT__ZTS1S_S:%.*]] addrspace(4)*, align 8 +// CHECK-NEXT: [[LHS_ADDR_ASCAST:%.*]] = addrspacecast [[STRUCT__ZTS1S_S]] addrspace(4)** [[LHS_ADDR]] to [[STRUCT__ZTS1S_S]] addrspace(4)* addrspace(4)* +// CHECK-NEXT: [[FROMBOOL:%.*]] = zext i1 [[COND:%.*]] to i8 +// CHECK-NEXT: store i8 [[FROMBOOL]], i8 addrspace(4)* [[COND_ADDR_ASCAST]], align 1, [[TBAA12:!tbaa !.*]] +// CHECK-NEXT: store [[STRUCT__ZTS1S_S]] addrspace(4)* [[LHS:%.*]], [[STRUCT__ZTS1S_S]] addrspace(4)* addrspace(4)* [[LHS_ADDR_ASCAST]], align 8, [[TBAA5:!tbaa !.*]] +// CHECK-NEXT: [[RHS_ASCAST:%.*]] = addrspacecast %struct._ZTS1S.S* [[RHS:%.*]] to [[STRUCT__ZTS1S_S]] addrspace(4)* +// CHECK-NEXT: [[TMP0:%.*]] = load i8, i8 addrspace(4)* [[COND_ADDR_ASCAST]], align 1, [[TBAA12]], [[RNG14:!range !.*]] +// CHECK-NEXT: [[TOBOOL:%.*]] = trunc i8 [[TMP0]] to i1 +// CHECK-NEXT: br i1 [[TOBOOL]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]] +// CHECK: cond.true: +// CHECK-NEXT: [[TMP1:%.*]] = load [[STRUCT__ZTS1S_S]] addrspace(4)*, [[STRUCT__ZTS1S_S]] addrspace(4)* addrspace(4)* [[LHS_ADDR_ASCAST]], align 8, [[TBAA5]] +// CHECK-NEXT: br label [[COND_END:%.*]] +// CHECK: cond.false: +// CHECK-NEXT: br label [[COND_END]] +// CHECK: cond.end: +// CHECK-NEXT: [[COND_LVALUE:%.*]] = phi [[STRUCT__ZTS1S_S]] addrspace(4)* [ [[TMP1]], [[COND_TRUE]] ], [ [[RHS_ASCAST]], [[COND_FALSE]] ] +// CHECK-NEXT: [[TMP2:%.*]] = bitcast [[STRUCT__ZTS1S_S]] addrspace(4)* [[AGG_RESULT:%.*]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP3:%.*]] = bitcast [[STRUCT__ZTS1S_S]] addrspace(4)* [[COND_LVALUE]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 2 [[TMP2]], i8 addrspace(4)* align 2 [[TMP3]], i64 2, i1 false), !tbaa.struct !9 +// CHECK-NEXT: ret void // -// CHECK: [[BFALSE]]: -// CHECK: %[[RHS:.+]] = addrspacecast [[STYPE]]* {{.+}} to [[STYPE]] addrspace(4)* -// CHECK: br label %[[BEND]] -// -// CHECK: [[BEND]]: -// CHECK: %{{.+}} = phi [[STYPE]] addrspace(4)* [ %[[LHS]], %[[BTRUE]] ], [ %[[RHS]], %[[BFALSE]] ] +S foo(bool cond, S &lhs, S rhs) { S val = cond ? lhs : rhs; return val; } diff --git a/clang/test/CodeGenSYCL/address-space-initializer.cpp b/clang/test/CodeGenSYCL/address-space-initializer.cpp new file mode 100644 index 0000000000000..c6946f9e9d9f4 --- /dev/null +++ b/clang/test/CodeGenSYCL/address-space-initializer.cpp @@ -0,0 +1,21 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice \ +// RUN: -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s + +// This test checks that data for big constant initializer lists is placed +// into the global address space by the SYCL compiler. + +struct Test { + Test() : set{ + 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, + 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 + // CHECK: @constinit = {{.*}}addrspace(1) {{.*}}[32 x i32] + // CHECK-SAME: [i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15, + // CHECK-SAME: i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15] + } {} + int set[32]; +}; + +__attribute__((sycl_device)) void foo() { + Test t; + (void)t; +} diff --git a/clang/test/CodeGenSYCL/address-space-new.cpp b/clang/test/CodeGenSYCL/address-space-new.cpp index 5e332e739b77c..0be436d6d5e4c 100644 --- a/clang/test/CodeGenSYCL/address-space-new.cpp +++ b/clang/test/CodeGenSYCL/address-space-new.cpp @@ -31,45 +31,42 @@ void test() { // CHECK: @[[STR:[.a-zA-Z0-9_]+]] = private unnamed_addr addrspace(1) constant [14 x i8] c"Hello, world!\00", align 1 - // CHECK: %i.ascast = addrspacecast i32* %i to i32 addrspace(4)* + // CHECK: %[[GEN:.*]] = addrspacecast i32* %i to i32 addrspace(4)* // CHECK: %[[ARR:[a-zA-Z0-9]+]] = alloca [42 x i32] + // CHECK: [[ARR_ASCAST:%.*]] = addrspacecast [42 x i32]* %[[ARR]] to [42 x i32] addrspace(4)* int i = 0; int *pptr = &i; - // CHECK: %[[GEN:[0-9]+]] = addrspacecast i32* %i to i32 addrspace(4)* - // CHECK: store i32 addrspace(4)* %[[GEN]], i32 addrspace(4)** %pptr + // CHECK: store i32 addrspace(4)* %[[GEN]], i32 addrspace(4)* addrspace(4)* %pptr.ascast bool is_i_ptr = (pptr == &i); - // CHECK: %[[VALPPTR:[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %pptr + // CHECK: %[[VALPPTR:[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %pptr.ascast // CHECK: %cmp{{[0-9]*}} = icmp eq i32 addrspace(4)* %[[VALPPTR]], %i.ascast *pptr = foo; int var23 = 23; char *cp = (char *)&var23; *cp = 41; - // CHECK: store i32 23, i32* %[[VAR:[a-zA-Z0-9]+]] - // CHECK: [[VARAS:[a-zA-Z0-9]+]] = addrspacecast i32* %[[VAR]] to i32 addrspace(4)* - // CHECK: [[VARCAST:[a-zA-Z0-9]+]] = bitcast i32 addrspace(4)* %[[VARAS]] to i8 addrspace(4)* - // CHECK: store i8 addrspace(4)* %[[VARCAST]], i8 addrspace(4)** %{{.*}} + // CHECK: store i32 23, i32 addrspace(4)* %[[VAR:[a-zA-Z0-9]+]] + // CHECK: [[VARCAST:[a-zA-Z0-9]+]] = bitcast i32 addrspace(4)* %[[VARAS:[a-zA-Z0-9]+]] to i8 addrspace(4)* + // CHECK: store i8 addrspace(4)* %[[VARCAST]], i8 addrspace(4)* addrspace(4)* %{{.*}} int arr[42]; char *cpp = (char *)arr; *cpp = 43; - // CHECK: %[[ARRDECAY:[a-zA-Z0-9]+]] = getelementptr inbounds [42 x i32], [42 x i32]* %[[ARR]], i64 0, i64 0 - // CHECK: %[[ARRAS:[a-zA-Z0-9]+]] = addrspacecast i32* %[[ARRDECAY]] to i32 addrspace(4)* - // CHECK: %[[ARRCAST:[a-zA-Z0-9]+]] = bitcast i32 addrspace(4)* %[[ARRAS]] to i8 addrspace(4)* - // CHECK: store i8 addrspace(4)* %[[ARRCAST]], i8 addrspace(4)** %{{.*}} + // CHECK: [[ARRAYDECAY1:%.*]] = getelementptr inbounds [42 x i32], [42 x i32] addrspace(4)* [[ARR_ASCAST]], i64 0, i64 0 + // CHECK: [[ADD_PTR:%.*]] = getelementptr inbounds i32, i32 addrspace(4)* [[ARRAYDECAY1]], i64 10 + // CHECK: store i32 addrspace(4)* [[ADD_PTR]], i32 addrspace(4)* addrspace(4)* %{{.*}} int *aptr = arr + 10; if (aptr < arr + sizeof(arr)) *aptr = 44; - // CHECK: %[[VALAPTR:[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %aptr - // CHECK: %[[ARRDCY2:[a-zA-Z0-9]+]] = getelementptr inbounds [42 x i32], [42 x i32]* %[[ARR]], i64 0, i64 0 - // CHECK: %[[ADDPTR:[a-zA-Z0-9.]+]] = getelementptr inbounds i32, i32* %[[ARRDCY2]], i64 168 - // CHECK: %[[ADDPTRCAST:[a-zA-Z0-9.]+]] = addrspacecast i32* %[[ADDPTR]] to i32 addrspace(4)* - // CHECK: %cmp{{[0-9]+}} = icmp ult i32 addrspace(4)* %[[VALAPTR]], %[[ADDPTRCAST]] + // CHECK: [[TMP13:%.*]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %aptr.ascast + // CHECK: [[ARRAYDECAY2:%.*]] = getelementptr inbounds [42 x i32], [42 x i32] addrspace(4)* [[ARR_ASCAST]], i64 0, i64 0 + // CHECK: [[ADD_PTR3:%.*]] = getelementptr inbounds i32, i32 addrspace(4)* [[ARRAYDECAY2]], i64 168 + // CHECK: [[CMP4:%.*]] = icmp ult i32 addrspace(4)* [[TMP13]], [[ADD_PTR3]] const char *str = "Hello, world!"; - // CHECK: store i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([14 x i8], [14 x i8] addrspace(1)* @[[STR]], i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)** %[[STRVAL:[a-zA-Z0-9]+]], align 8 + // CHECK: store i8 addrspace(4)* getelementptr inbounds ([14 x i8], [14 x i8] addrspace(4)* addrspacecast ([14 x i8] addrspace(1)* @.str to [14 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* addrspace(4)* %[[STRVAL:[a-zA-Z0-9.]+]], align 8 i = str[0]; @@ -79,28 +76,27 @@ void test() { // CHECK: br i1 %[[COND]], label %[[CONDTRUE:[.a-zA-Z0-9]+]], label %[[CONDFALSE:[.a-zA-Z0-9]+]] // CHECK: [[CONDTRUE]]: - // CHECK-NEXT: %[[VALTRUE:[a-zA-Z0-9]+]] = load i8 addrspace(4)*, i8 addrspace(4)** %[[STRVAL]] + // CHECK-NEXT: %[[VALTRUE:[a-zA-Z0-9]+]] = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %[[STRVAL]] // CHECK-NEXT: br label %[[CONDEND:[.a-zA-Z0-9]+]] // CHECK: [[CONDFALSE]]: // CHECK: [[CONDEND]]: - // CHECK-NEXT: phi i8 addrspace(4)* [ %[[VALTRUE]], %[[CONDTRUE]] ], [ addrspacecast (i8 addrspace(1)* getelementptr inbounds ([21 x i8], [21 x i8] addrspace(1)* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), %[[CONDFALSE]] ] + // CHECK-NEXT: phi i8 addrspace(4)* [ %[[VALTRUE]], %[[CONDTRUE]] ], [ getelementptr inbounds ([21 x i8], [21 x i8] addrspace(4)* addrspacecast ([21 x i8] addrspace(1)* @{{.*}} to [21 x i8] addrspace(4)*), i64 0, i64 0), %[[CONDFALSE]] ] const char *select_null = i > 2 ? "Yet another Hello world" : nullptr; (void)select_null; - // CHECK: select i1 %{{.*}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([24 x i8], [24 x i8] addrspace(1)* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* null + // CHECK: select i1 %{{.*}}, i8 addrspace(4)* getelementptr inbounds ([24 x i8], [24 x i8] addrspace(4)* addrspacecast ([24 x i8] addrspace(1)* @{{.*}} to [24 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* null const char *select_str_trivial1 = true ? str : "Another hello world!"; (void)select_str_trivial1; - // CHECK: %[[TRIVIALTRUE:[a-zA-Z0-9]+]] = load i8 addrspace(4)*, i8 addrspace(4)** %[[STRVAL]] - // CHECK: store i8 addrspace(4)* %[[TRIVIALTRUE]], i8 addrspace(4)** %{{.*}}, align 8 + // CHECK: %[[TRIVIALTRUE:[a-zA-Z0-9]+]] = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %[[STRVAL]] + // CHECK: store i8 addrspace(4)* %[[TRIVIALTRUE]], i8 addrspace(4)* addrspace(4)* %{{.*}}, align 8 const char *select_str_trivial2 = false ? str : "Another hello world!"; (void)select_str_trivial2; - // CHECK: store i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([21 x i8], [21 x i8] addrspace(1)* @{{.*}}, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)** %{{.*}} - // - // + // CHECK: store i8 addrspace(4)* getelementptr inbounds ([21 x i8], [21 x i8] addrspace(4)* addrspacecast ([21 x i8] addrspace(1)* @{{.*}} to [21 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* addrspace(4)* %{{.*}} + Y yy; baz(yy); // CHECK: define {{.*}}spir_func void @{{.*}}baz{{.*}} diff --git a/clang/test/CodeGenSYCL/address-space-of-returns.cpp b/clang/test/CodeGenSYCL/address-space-of-returns.cpp index a0b477c4e6d51..4602160feb784 100644 --- a/clang/test/CodeGenSYCL/address-space-of-returns.cpp +++ b/clang/test/CodeGenSYCL/address-space-of-returns.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -fsycl -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck %s struct A { int B[42]; @@ -7,17 +7,16 @@ struct A { const char *ret_char() { return "N"; } -// CHECK: ret i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([2 x i8], [2 x i8] addrspace(1)* @.str, i64 0, i64 0) to i8 addrspace(4)*) +// CHECK: ret i8 addrspace(4)* getelementptr inbounds ([2 x i8], [2 x i8] addrspace(4)* addrspacecast ([2 x i8] addrspace(1)* @.str to [2 x i8] addrspace(4)*), i64 0, i64 0) const char *ret_arr() { - const static char Arr[36] = "Carrots, cabbage, radish, potatoes!"; + static const char Arr[42] = {0}; return Arr; } - -// CHECK: ret i8 addrspace(4)* getelementptr inbounds ([36 x i8], [36 x i8] addrspace(4)* addrspacecast ([36 x i8] addrspace(1)* @{{.*}}ret_arr{{.*}}Arr to [36 x i8] addrspace(4)*), i64 0, i64 0) +// CHECK: ret i8 addrspace(4)* getelementptr inbounds ([42 x i8], [42 x i8] addrspace(4)* addrspacecast ([42 x i8] addrspace(1)* @{{.*}}ret_arr{{.*}}Arr to [42 x i8] addrspace(4)*), i64 0, i64 0) const char &ret_ref() { - const static char a = 'A'; + static const char a = 'A'; return a; } // CHECK: ret i8 addrspace(4)* addrspacecast (i8 addrspace(1)* @{{.*}}ret_ref{{.*}} to i8 addrspace(4)*) @@ -26,7 +25,7 @@ A ret_agg() { A a; return a; } -// CHECK: define {{.*}}spir_func void @{{.*}}ret_agg{{.*}}(%struct.{{.*}}.A addrspace(4)* {{.*}} %agg.result) +// CHECK: define {{.*}}spir_func void @{{.*}}ret_agg{{.*}}(%struct{{.*}}.A addrspace(4)* noalias sret(%struct{{.*}}.A) align 4 %agg.result) template __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { diff --git a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp index cc3ad88b64940..c1b9d6f430b06 100644 --- a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp +++ b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp @@ -21,133 +21,139 @@ void tmpl(T t){} // See Check Lines below. void usages() { - // CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca i32 addrspace(1)* + // CHECK-DAG: [[GLOBA:%[a-zA-Z0-9]+]] = alloca i32 addrspace(1)* + // CHECK-DAG: [[GLOB:%.*]] = addrspacecast i32 addrspace(1)** [[GLOBA]] to i32 addrspace(1)* addrspace(4)* __attribute__((opencl_global)) int *GLOB; - // CHECK-DAG: [[GLOBDEV:%[a-zA-Z0-9]+]] = alloca i32 addrspace(5)* + // CHECK-DAG: [[GLOBDEVA:%[a-zA-Z0-9]+]] = alloca i32 addrspace(5)* + // CHECK-DAG: [[GLOBDEV:%.*]] = addrspacecast i32 addrspace(5)** [[GLOBDEVA]] to i32 addrspace(5)* addrspace(4)* __attribute__((opencl_global_device)) int *GLOBDEV; - // CHECK-DAG: [[GLOBHOST:%[a-zA-Z0-9]+]] = alloca i32 addrspace(6)* + // CHECK-DAG: [[GLOBHOSTA:%[a-zA-Z0-9]+]] = alloca i32 addrspace(6)* + // CHECK-DAG: [[GLOBHOST:%.*]] = addrspacecast i32 addrspace(6)** [[GLOBHOSTA]] to i32 addrspace(6)* addrspace(4)* __attribute__((opencl_global_host)) int *GLOBHOST; - // CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca i32 addrspace(3)* + // CHECK-DAG: [[LOCA:%[a-zA-Z0-9]+]] = alloca i32 addrspace(3)* + // CHECK-DAG: [[LOC:%.*]] = addrspacecast i32 addrspace(3)** [[LOCA]] to i32 addrspace(3)* addrspace(4)* __attribute__((opencl_local)) int *LOC; - // CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca i32 addrspace(4)* + // CHECK-DAG: [[NoASA:%[a-zA-Z0-9]+]] = alloca i32 addrspace(4)* + // CHECK-DAG: [[NoAS:%.*]] = addrspacecast i32 addrspace(4)** [[NoASA]] to i32 addrspace(4)* addrspace(4)* int *NoAS; - // CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca i32* + // CHECK-DAG: [[PRIVA:%[a-zA-Z0-9]+]] = alloca i32* + // CHECK-DAG: [[PRIV:%.*]] = addrspacecast i32** [[PRIVA]] to i32* addrspace(4)* __attribute__((opencl_private)) int *PRIV; bar(*GLOB); - // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[GLOB]] // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* align 4 dereferenceable(4) [[GLOB_CAST]]) bar2(*GLOB); - // CHECK-DAG: [[GLOB_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[GLOB]] // CHECK-DAG: [[GLOB_CAST2:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD2]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* align 4 dereferenceable(4) [[GLOB_CAST2]]) bar(*GLOBDEV); - // CHECK-DAG: [[GLOBDEV_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)** [[GLOBDEV]] + // CHECK-DAG: [[GLOBDEV_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)* addrspace(4)* [[GLOBDEV]] // CHECK-DAG: [[GLOBDEV_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(5)* [[GLOBDEV_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* align 4 dereferenceable(4) [[GLOBDEV_CAST]]) bar2(*GLOBDEV); - // CHECK-DAG: [[GLOBDEV_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)** [[GLOBDEV]] + // CHECK-DAG: [[GLOBDEV_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)* addrspace(4)* [[GLOBDEV]] // CHECK-DAG: [[GLOBDEV_CAST2:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(5)* [[GLOBDEV_LOAD2]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* align 4 dereferenceable(4) [[GLOBDEV_CAST2]]) bar3(*GLOBDEV); - // CHECK-DAG: [[GLOBDEV_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)** [[GLOBDEV]] + // CHECK-DAG: [[GLOBDEV_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)* addrspace(4)* [[GLOBDEV]] // CHECK-DAG: [[GLOBDEV_CAST3:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(5)* [[GLOBDEV_LOAD3]] to i32 addrspace(1)* // CHECK-DAG: call spir_func void @[[GLOB_REF]](i32 addrspace(1)* align 4 dereferenceable(4) [[GLOBDEV_CAST3]]) bar(*GLOBHOST); - // CHECK-DAG: [[GLOBHOST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)** [[GLOBHOST]] + // CHECK-DAG: [[GLOBHOST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)* addrspace(4)* [[GLOBHOST]] // CHECK-DAG: [[GLOBHOST_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(6)* [[GLOBHOST_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* align 4 dereferenceable(4) [[GLOBHOST_CAST]]) bar2(*GLOBHOST); - // CHECK-DAG: [[GLOBHOST_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)** [[GLOBHOST]] + // CHECK-DAG: [[GLOBHOST_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)* addrspace(4)* [[GLOBHOST]] // CHECK-DAG: [[GLOBHOST_CAST2:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(6)* [[GLOBHOST_LOAD2]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* align 4 dereferenceable(4) [[GLOBHOST_CAST2]]) bar3(*GLOBHOST); - // CHECK-DAG: [[GLOBHOST_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)** [[GLOBHOST]] + // CHECK-DAG: [[GLOBHOST_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)* addrspace(4)* [[GLOBHOST]] // CHECK-DAG: [[GLOBHOST_CAST3:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(6)* [[GLOBHOST_LOAD3]] to i32 addrspace(1)* // CHECK-DAG: call spir_func void @[[GLOB_REF]](i32 addrspace(1)* align 4 dereferenceable(4) [[GLOBHOST_CAST3]]) bar(*LOC); - // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)* addrspace(4)* [[LOC]] // CHECK-DAG: call spir_func void [[LOC_REF]](i32 addrspace(3)* align 4 dereferenceable(4) [[LOC_LOAD]]) bar2(*LOC); - // CHECK-DAG: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)* addrspace(4)* [[LOC]] // CHECK-DAG: [[LOC_CAST2:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOC_LOAD2]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* align 4 dereferenceable(4) [[LOC_CAST2]]) bar(*NoAS); - // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* [[NoAS]] // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* align 4 dereferenceable(4) [[NoAS_LOAD]]) bar2(*NoAS); - // CHECK-DAG: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* [[NoAS]] // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* align 4 dereferenceable(4) [[NoAS_LOAD2]]) foo(GLOB); - // CHECK-DAG: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[GLOB]] // CHECK-DAG: [[GLOB_CAST3:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD3]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_PTR]](i32 addrspace(4)* [[GLOB_CAST3]]) foo2(GLOB); - // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[GLOB]] // CHECK-DAG: [[GLOB_CAST4:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD4]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[GLOB_CAST4]]) foo(GLOBDEV); - // CHECK-DAG: [[GLOBDEV_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)** [[GLOBDEV]] + // CHECK-DAG: [[GLOBDEV_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)* addrspace(4)* [[GLOBDEV]] // CHECK-DAG: [[GLOBDEV_CAST4:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(5)* [[GLOBDEV_LOAD4]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_PTR]](i32 addrspace(4)* [[GLOBDEV_CAST4]]) foo2(GLOBDEV); - // CHECK-DAG: [[GLOBDEV_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)** [[GLOBDEV]] + // CHECK-DAG: [[GLOBDEV_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)* addrspace(4)* [[GLOBDEV]] // CHECK-DAG: [[GLOBDEV_CAST5:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(5)* [[GLOBDEV_LOAD5]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[GLOBDEV_CAST5]]) foo3(GLOBDEV); - // CHECK-DAG: [[GLOBDEV_LOAD6:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)** [[GLOBDEV]] + // CHECK-DAG: [[GLOBDEV_LOAD6:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)* addrspace(4)* [[GLOBDEV]] // CHECK-DAG: [[GLOBDEV_CAST6:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(5)* [[GLOBDEV_LOAD6]] to i32 addrspace(1)* // CHECK-DAG: call spir_func void @[[GLOB_PTR]](i32 addrspace(1)* [[GLOBDEV_CAST6]]) foo(GLOBHOST); - // CHECK-DAG: [[GLOBHOST_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)** [[GLOBHOST]] + // CHECK-DAG: [[GLOBHOST_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)* addrspace(4)* [[GLOBHOST]] // CHECK-DAG: [[GLOBHOST_CAST4:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(6)* [[GLOBHOST_LOAD4]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_PTR]](i32 addrspace(4)* [[GLOBHOST_CAST4]]) foo2(GLOBHOST); - // CHECK-DAG: [[GLOBHOST_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)** [[GLOBHOST]] + // CHECK-DAG: [[GLOBHOST_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)* addrspace(4)* [[GLOBHOST]] // CHECK-DAG: [[GLOBHOST_CAST5:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(6)* [[GLOBHOST_LOAD5]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[GLOBHOST_CAST5]]) foo3(GLOBHOST); - // CHECK-DAG: [[GLOBHOST_LOAD6:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)** [[GLOBHOST]] + // CHECK-DAG: [[GLOBHOST_LOAD6:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)* addrspace(4)* [[GLOBHOST]] // CHECK-DAG: [[GLOBHOST_CAST6:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(6)* [[GLOBHOST_LOAD6]] to i32 addrspace(1)* // CHECK-DAG: call spir_func void @[[GLOB_PTR]](i32 addrspace(1)* [[GLOBHOST_CAST6]]) foo(LOC); - // CHECK-DAG: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)* addrspace(4)* [[LOC]] // CHECK-DAG: call spir_func void [[LOC_PTR]](i32 addrspace(3)* [[LOC_LOAD3]]) foo2(LOC); - // CHECK-DAG: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)* addrspace(4)* [[LOC]] // CHECK-DAG: [[LOC_CAST4:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOC_LOAD4]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[LOC_CAST4]]) foo(NoAS); - // CHECK-DAG: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* [[NoAS]] // CHECK-DAG: call spir_func void @[[RAW_PTR]](i32 addrspace(4)* [[NoAS_LOAD3]]) foo2(NoAS); - // CHECK-DAG: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* [[NoAS]] // CHECK-DAG: call spir_func void @[[RAW_PTR2]](i32 addrspace(4)* [[NoAS_LOAD4]]) // Ensure that we still get 5 different template instantiations. tmpl(GLOB); - // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[GLOB]] // CHECK-DAG: call spir_func void [[GLOB_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(1)* [[GLOB_LOAD4]]) tmpl(GLOBDEV); - // CHECK-DAG: [[GLOBDEV_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)** [[GLOBDEV]] + // CHECK-DAG: [[GLOBDEV_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)* addrspace(4)* [[GLOBDEV]] // CHECK-DAG: call spir_func void [[GLOBDEV_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(5)* [[GLOBDEV_LOAD4]]) tmpl(GLOBHOST); - // CHECK-DAG: [[GLOBHOST_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)** [[GLOBHOST]] + // CHECK-DAG: [[GLOBHOST_LOAD4:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)* addrspace(4)* [[GLOBHOST]] // CHECK-DAG: call spir_func void [[GLOBHOST_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(6)* [[GLOBHOST_LOAD4]]) tmpl(LOC); - // CHECK-DAG: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOC]] + // CHECK-DAG: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)* addrspace(4)* [[LOC]] // CHECK-DAG: call spir_func void [[LOC_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(3)* [[LOC_LOAD5]]) tmpl(PRIV); - // CHECK-DAG: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]] + // CHECK-DAG: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load i32*, i32* addrspace(4)* [[PRIV]] // CHECK-DAG: call spir_func void [[PRIV_TMPL:@[a-zA-Z0-9_]+]](i32* [[PRIV_LOAD5]]) tmpl(NoAS); - // CHECK-DAG: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[NoAS]] + // CHECK-DAG: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* [[NoAS]] // CHECK-DAG: call spir_func void [[GEN_TMPL:@[a-zA-Z0-9_]+]](i32 addrspace(4)* [[NoAS_LOAD5]]) } @@ -173,23 +179,23 @@ void usages2() { // CHECK-DAG: [[LOCAL:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)* bar(*PRIV); - // CHECK-DAG: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]] + // CHECK-DAG: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32* addrspace(4)* [[PRIV]] // CHECK-DAG: [[PRIV_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* align 4 dereferenceable(4) [[PRIV_ASCAST]]) bar(*GLOB); - // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]] + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[GLOB]] // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* align 4 dereferenceable(4) [[GLOB_CAST]]) bar(*GLOBDEV); - // CHECK-DAG: [[GLOBDEV_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)** [[GLOBDEV]] + // CHECK-DAG: [[GLOBDEV_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(5)*, i32 addrspace(5)* addrspace(4)* [[GLOBDEV]] // CHECK-DAG: [[GLOBDEV_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(5)* [[GLOBDEV_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* align 4 dereferenceable(4) [[GLOBDEV_CAST]]) bar(*GLOBHOST); - // CHECK-DAG: [[GLOBHOST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)** [[GLOBHOST]] + // CHECK-DAG: [[GLOBHOST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(6)*, i32 addrspace(6)* addrspace(4)* [[GLOBHOST]] // CHECK-DAG: [[GLOBHOST_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(6)* [[GLOBHOST_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* align 4 dereferenceable(4) [[GLOBHOST_CAST]]) bar2(*LOCAL); - // CHECK-DAG: [[LOCAL_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL]] + // CHECK-DAG: [[LOCAL_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)* addrspace(4)* [[LOCAL]] // CHECK-DAG: [[LOCAL_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_LOAD]] to i32 addrspace(4)* // CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* align 4 dereferenceable(4) [[LOCAL_CAST]]) } diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp index b381dad7b6fe3..2d9748e3b138d 100644 --- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp @@ -27,30 +27,33 @@ int main() { // Check alloca for pointer argument // CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)* // Check lambda object alloca -// CHECK: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon" +// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %"class.{{.*}}.anon" +// CHECK: [[ANON:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon"* [[ANONALLOCA]] to %"class.{{.*}}.anon" addrspace(4)* // Check allocas for ranges -// CHECK: [[ARANGE:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MRANGE:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[OID:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::id" +// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ARANGEA]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MRANGEA]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::id" +// CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OIDA]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* // // Check store of kernel pointer argument to alloca -// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)** [[MEM_ARG]].addr, align 8 +// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)* addrspace(4)* [[MEM_ARG]].addr.ascast, align 8 // Check for default constructor of accessor // CHECK: call spir_func {{.*}}accessor // Check accessor GEP -// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0 +// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[ANON]], i32 0, i32 0 // Check load from kernel pointer argument alloca -// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr +// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG]].addr.ascast // Check accessor __init method call -// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) -// CHECK: [[ACCESSORCAST:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[ACCESSOR]] to %"class{{.*}}accessor" addrspace(4)* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACCESSORCAST]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) +// CHECK: [[ARANGE:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ARANGET]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[MRANGE:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MRANGET]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[OID:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OIDT]] to %"struct.{{.*}}.cl::sycl::id"* +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) // Check lambda "()" operator call -// CHECK-OLD: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]]) -// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %"class{{.*}}anon"* {{.*}} to %"class{{.*}}anon" addrspace(4)* -// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} [[ANONCAST]]) +// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}}) diff --git a/clang/test/CodeGenSYCL/device-variables.cpp b/clang/test/CodeGenSYCL/device-variables.cpp index f56a2b432e53a..79820a7ac5213 100644 --- a/clang/test/CodeGenSYCL/device-variables.cpp +++ b/clang/test/CodeGenSYCL/device-variables.cpp @@ -21,23 +21,23 @@ int main() { int some_local_var = 10; kernel([=]() { -// Global variables used directly + // Global variables used directly foo(global_value); -// CHECK: call spir_func void @{{.*}}foo{{.*}}(i32 addrspace(4)* align 4 dereferenceable(4) addrspacecast (i32 addrspace(1)* @{{.*}}global_value to i32 addrspace(4)*)) + // CHECK: call spir_func void @{{.*}}foo{{.*}}(i32 addrspace(4)* align 4 dereferenceable(4) addrspacecast (i32 addrspace(1)* @{{.*}}global_value to i32 addrspace(4)*)) int a = my_array[0]; -// CHECK: [[LOAD:%[0-9]+]] = load i32, i32 addrspace(4)* getelementptr inbounds ([1 x i32], [1 x i32] addrspace(4)* addrspacecast ([1 x i32] addrspace(1)* @{{.*}}my_array to [1 x i32] addrspace(4)*), i64 0, i64 0) -// CHECK: store i32 [[LOAD]], i32* %a + // CHECK: [[LOAD:%[0-9]+]] = load i32, i32 addrspace(4)* getelementptr inbounds ([1 x i32], [1 x i32] addrspace(4)* addrspacecast ([1 x i32] addrspace(1)* @{{.*}}my_array to [1 x i32] addrspace(4)*), i64 0, i64 0) + // CHECK: store i32 [[LOAD]], i32 addrspace(4)* %a int b = some_const; -// Constant used directly -// CHECK: store i32 1, i32* %b + // Constant used directly + // CHECK: store i32 1, i32 addrspace(4)* %b foo(local_value); -// Local variables and constexprs captured by lambda -// CHECK: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* %{{.*}}, i32 0, i32 0 -// CHECK: call spir_func void @{{.*}}foo{{.*}}(i32 addrspace(4)* align 4 dereferenceable(4) [[GEP]]) + // Local variables and constexprs captured by lambda + // CHECK: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* %{{.*}}, i32 0, i32 0 + // CHECK: call spir_func void @{{.*}}foo{{.*}}(i32 addrspace(4)* align 4 dereferenceable(4) [[GEP]]) int some_device_local_var = some_local_var; -// CHECK: [[GEP1:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* %{{.*}}, i32 0, i32 1 -// CHECK: [[LOAD1:%[0-9]+]] = load i32, i32 addrspace(4)* [[GEP1]] -// CHECK: store i32 [[LOAD1]], i32* %some_device_local_var + // CHECK: [[GEP1:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* %{{.*}}, i32 0, i32 1 + // CHECK: [[LOAD1:%[0-9]+]] = load i32, i32 addrspace(4)* [[GEP1]] + // CHECK: store i32 [[LOAD1]], i32 addrspace(4)* %some_device_local_var }); return 0; diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp index bd5f996c64796..02453e1c61cf0 100644 --- a/clang/test/CodeGenSYCL/inheritance.cpp +++ b/clang/test/CodeGenSYCL/inheritance.cpp @@ -43,30 +43,31 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}derived(%struct.{{.*}}.base* byval(%struct.{{.*}}.base) align 4 %_arg__base, %struct.{{.*}}.__wrapper_class* byval(%struct.{{.*}}.__wrapper_class) align 8 %_arg_e, i32 %_arg_a) // Check alloca for kernel paramters -// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = alloca i32, align 4 +// CHECK: %[[ARG_AA:[a-zA-Z0-9_.]+]] = alloca i32, align 4 +// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast i32* %[[ARG_AA]] to i32 addrspace(4)* // Check alloca for local functor object // CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.{{.*}}.derived, align 8 -// CHECK: store i32 %_arg_a, i32* %[[ARG_A]], align 4 +// CHECK: store i32 %_arg_a, i32 addrspace(4)* %[[ARG_A]], align 4 // Initialize 'base' subobject -// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to %struct.{{.*}}.base* -// CHECK: %[[BASE_TO_PTR:.*]] = bitcast %struct.{{.*}}.base* %[[DERIVED_TO_BASE]] to i8* -// CHECK: %[[PARAM_TO_PTR:.*]] = bitcast %struct.{{.*}}.base* %_arg__base to i8* -// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %[[BASE_TO_PTR]], i8* align 4 %[[PARAM_TO_PTR]], i64 12, i1 false) +// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.{{.*}}.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast to %struct.{{.*}}.base addrspace(4)* +// CHECK: %[[BASE_TO_PTR:.*]] = bitcast %struct.{{.*}}.base addrspace(4)* %[[DERIVED_TO_BASE]] to i8 addrspace(4)* +// CHECK: %[[PARAM_TO_PTR:.*]] = bitcast %struct.{{.*}}.base addrspace(4)* %_arg__base.ascast to i8 addrspace(4)* +// CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %[[BASE_TO_PTR]], i8 addrspace(4)* align 4 %[[PARAM_TO_PTR]], i64 12, i1 false) // Initialize 'second_base' subobject // First, derived-to-base cast with offset: -// CHECK: %[[DERIVED_PTR:.*]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to i8* -// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, i8* %[[DERIVED_PTR]], i64 16 -// CHECK: %[[TO_SECOND_BASE:.*]] = bitcast i8* %[[OFFSET_CALC]] to %class.{{.*}}.second_base* +// CHECK: %[[DERIVED_PTR:.*]] = bitcast %struct.{{.*}}.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast to i8 addrspace(4)* +// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, i8 addrspace(4)* %[[DERIVED_PTR]], i64 16 +// CHECK: %[[TO_SECOND_BASE:.*]] = bitcast i8 addrspace(4)* %[[OFFSET_CALC]] to %class.{{.*}}.second_base addrspace(4)* // Initialize 'second_base::e' -// CHECK: %[[SECOND_BASE_PTR:.*]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base* %[[TO_SECOND_BASE]], i32 0, i32 0 -// CHECK: %[[PTR_TO_WRAPPER:.*]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class, %struct.{{.*}}.__wrapper_class* %_arg_e, i32 0, i32 0 -// CHECK: %[[LOAD_PTR:.*]] = load i32 addrspace(1)*, i32 addrspace(1)** %[[PTR_TO_WRAPPER]] +// CHECK: %[[SECOND_BASE_PTR:.*]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base addrspace(4)* %[[TO_SECOND_BASE]], i32 0, i32 0 +// CHECK: %[[PTR_TO_WRAPPER:.*]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class, %struct.{{.*}}.__wrapper_class addrspace(4)* %_arg_e.ascast, i32 0, i32 0 +// CHECK: %[[LOAD_PTR:.*]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* %[[PTR_TO_WRAPPER]] // CHECK: %[[AS_CAST:.*]] = addrspacecast i32 addrspace(1)* %[[LOAD_PTR]] to i32 addrspace(4)* -// CHECK: store i32 addrspace(4)* %[[AS_CAST]], i32 addrspace(4)** %[[SECOND_BASE_PTR]] +// CHECK: store i32 addrspace(4)* %[[AS_CAST]], i32 addrspace(4)* addrspace(4)* %[[SECOND_BASE_PTR]] // Initialize field 'a' -// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived* %[[LOCAL_OBJECT]], i32 0, i32 3 -// CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32* %[[ARG_A]], align 4 -// CHECK: store i32 %[[LOAD_A]], i32* %[[GEP_A]] +// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast, i32 0, i32 3 +// CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32 addrspace(4)* %[[ARG_A]], align 4 +// CHECK: store i32 %[[LOAD_A]], i32 addrspace(4)* %[[GEP_A]] diff --git a/clang/test/CodeGenSYCL/inline_asm.cpp b/clang/test/CodeGenSYCL/inline_asm.cpp index be63c64f07e7a..f06d23f8bfcc6 100644 --- a/clang/test/CodeGenSYCL/inline_asm.cpp +++ b/clang/test/CodeGenSYCL/inline_asm.cpp @@ -6,13 +6,15 @@ template __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [100 x i32], align 4 // CHECK: %[[IDX:.*]] = getelementptr inbounds [100 x i32], [100 x i32]* %[[ARRAY_A]], i64 0, i64 0 + // CHECK: %[[IDX4:.*]] = addrspacecast i32* %[[IDX]] to i32 addrspace(4)* int a[100], i = 0; // CHECK-NEXT: call void asm sideeffect // CHECK: ".decl V52 v_type=G type=d num_elts=16 align=GRF // CHECK: svm_gather.4.1 (M1, 16) $0.0 V52.0 // CHECK: add(M1, 16) V52(0, 0)<1> V52(0, 0)<1; 1, 0> 0x1 // CHECK: svm_scatter.4.1 (M1, 16) $0.0 V52.0", - // CHECK: "rw"(i32* nonnull %[[IDX]]) + // CHECK: "rw"(i32 addrspace(4)* %[[IDX4]]) + // TODO: nonnull attribute missing? asm volatile(".decl V52 v_type=G type=d num_elts=16 align=GRF\n" "svm_gather.4.1 (M1, 16) %0.0 V52.0\n" "add(M1, 16) V52(0, 0)<1> V52(0, 0)<1; 1, 0> 0x1\n" diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp index 9615d3da4a8a0..47ef2d5addd3c 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp @@ -9,9 +9,9 @@ void ivdep_array_no_safelen() { // CHECK: %[[ARRAY_B:[0-9a-z]+]] = alloca [10 x i32] int b[10]; [[intel::ivdep(a)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_ARR:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_ARR:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}{{[[:space:]]}} + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}{{[[:space:]]}} b[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_ARR:[0-9]+]] } @@ -25,9 +25,9 @@ void ivdep_array_with_safelen() { // CHECK: %[[ARRAY_B:[0-9a-z]+]] = alloca [10 x i32] int b[10]; [[intel::ivdep(a, 5)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_ARR_SAFELEN:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_ARR_SAFELEN:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}{{[[:space:]]}} + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}{{[[:space:]]}} b[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_ARR_SAFELEN:[0-9]+]] } @@ -49,13 +49,13 @@ void ivdep_multiple_arrays() { [[intel::ivdep(b, 5)]] [[intel::ivdep(c)]] [[intel::ivdep(d)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_MUL_ARR:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_MUL_ARR:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_MUL_ARR:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_MUL_ARR:[0-9]+]] b[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_C]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_C_MUL_ARR:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_C]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_C_MUL_ARR:[0-9]+]] c[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_D]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_D_MUL_ARR:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_D]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_D_MUL_ARR:[0-9]+]] d[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_MUL_ARR:[0-9]+]] } @@ -71,9 +71,9 @@ void ivdep_array_and_global() { int b[10]; [[intel::ivdep]] [[intel::ivdep(a)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_ARR_AND_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_ARR_AND_GLOB:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_ARR_AND_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_ARR_AND_GLOB:[0-9]+]] b[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_ARR_AND_GLOB:[0-9]+]] } @@ -89,9 +89,9 @@ void ivdep_array_and_inf_global() { int b[10]; [[intel::ivdep]] [[intel::ivdep(a, 8)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_ARR_AND_INF_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_ARR_AND_INF_GLOB:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_ARR_AND_INF_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_ARR_AND_INF_GLOB:[0-9]+]] b[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_ARR_AND_INF_GLOB:[0-9]+]] } @@ -107,9 +107,9 @@ void ivdep_array_and_greater_global() { int b[10]; [[intel::ivdep(9)]] [[intel::ivdep(a, 8)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_ARR_AND_GREAT_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_ARR_AND_GREAT_GLOB:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_ARR_AND_GREAT_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_ARR_AND_GREAT_GLOB:[0-9]+]] b[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_ARR_AND_GREAT_GLOB:[0-9]+]] } @@ -128,11 +128,11 @@ void ivdep_mul_arrays_and_global() { [[intel::ivdep(5)]] [[intel::ivdep(b, 6)]] [[intel::ivdep(c)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_MUL_ARR_AND_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_MUL_ARR_AND_GLOB:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_MUL_ARR_AND_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_MUL_ARR_AND_GLOB:[0-9]+]] b[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_C]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_C_MUL_ARR_AND_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_C]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_C_MUL_ARR_AND_GLOB:[0-9]+]] c[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_MUL_ARR_AND_GLOB:[0-9]+]] } @@ -144,7 +144,7 @@ void ivdep_ptr() { // CHECK: %[[PTR:[0-9a-z]+]] = alloca i32 addrspace(4)* [[intel::ivdep(ptr, 5)]] for (int i = 0; i != 10; ++i) ptr[i] = 0; - // CHECK: %[[PTR_LOAD:[0-9a-z]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %[[PTR]] + // CHECK: %[[PTR_LOAD:[0-9a-z]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %[[PTR]] // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds i32, i32 addrspace(4)* %[[PTR_LOAD]], i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_PTR:[0-9]+]] // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_PTR:[0-9]+]] } @@ -158,14 +158,14 @@ void ivdep_struct() { // CHECK: %[[STRUCT:[0-9a-z]+]] = alloca %struct.{{.+}}.S [[intel::ivdep(s.arr, 5)]] for (int i = 0; i != 10; ++i) s.arr[i] = 0; - // CHECK: %[[STRUCT_ARR:[0-9a-z]+]] = getelementptr inbounds %struct.{{.+}}.S, %struct.{{.+}}.S* %[[STRUCT]], i32 0, i32 1 - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[STRUCT_ARR]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_STRUCT_ARR:[0-9]+]] + // CHECK: %[[STRUCT_ARR:[0-9a-z]+]] = getelementptr inbounds %struct.{{.+}}.S, %struct.{{.+}}.S addrspace(4)* %[[STRUCT]].ascast, i32 0, i32 1 + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[STRUCT_ARR]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_STRUCT_ARR:[0-9]+]] // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_STRUCT_ARR:[0-9]+]] [[intel::ivdep(s.ptr, 5)]] for (int i = 0; i != 10; ++i) s.ptr[i] = 0; - // CHECK: %[[STRUCT_PTR:[0-9a-z]+]] = getelementptr inbounds %struct.{{.+}}.S, %struct.{{.+}}.S* %[[STRUCT]], i32 0, i32 0 - // CHECK: %[[LOAD_STRUCT_PTR:[0-9a-z]+]] = load i32 addrspace(4)*, i32 addrspace(4)** %[[STRUCT_PTR]] + // CHECK: %[[STRUCT_PTR:[0-9a-z]+]] = getelementptr inbounds %struct.{{.+}}.S, %struct.{{.+}}.S addrspace(4)* %[[STRUCT]].ascast, i32 0, i32 0 + // CHECK: %[[LOAD_STRUCT_PTR:[0-9a-z]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %[[STRUCT_PTR]] // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds i32, i32 addrspace(4)* %[[LOAD_STRUCT_PTR]], i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_STRUCT_PTR:[0-9]+]] // CHECK: br label %for.cond{{[0-9]*}}, !llvm.loop ![[MD_LOOP_STRUCT_PTR:[0-9]+]] } diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp index a1511f0c3886d..bc47cf4c9784e 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp @@ -8,8 +8,8 @@ void ivdep_inner_loop_access() { int a[10]; [[intel::ivdep]] for (int i = 0; i != 10; ++i) { [[intel::ivdep(3)]] for (int j = 0; j != 10; ++j) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_ACCESS:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_ACCESS]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_ACCESS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_ACCESS]] a[i] = a[(i + j) % 10]; // CHECK: br label %for.cond{{[0-9]*}}, !llvm.loop ![[MD_INNER_LOOP_INNER_ACCESS:[0-9]+]] } @@ -24,12 +24,12 @@ void ivdep_embedded_global_safelen() { // CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32] int a[10]; [[intel::ivdep]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_OUTER_GLOB_SFLN:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_OUTER_GLOB_SFLN]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_OUTER_GLOB_SFLN:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_OUTER_GLOB_SFLN]] a[i] = a[i % 2]; [[intel::ivdep]] for (int j = 0; j != 10; ++j) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_GLOB_SFLN:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_GLOB_SFLN]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_GLOB_SFLN:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_GLOB_SFLN]] a[i] = a[(i + j) % 10]; // CHECK: br label %for.cond{{[0-9]*}}, !llvm.loop ![[MD_LOOP_INNER_GLOB_SFLN:[0-9]+]] } @@ -44,12 +44,12 @@ void ivdep_embedded_various_safelens() { // CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32] int a[10]; [[intel::ivdep(a, 4)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_OUTER_VAR_SFLN:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_OUTER_VAR_SFLN]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_OUTER_VAR_SFLN:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_OUTER_VAR_SFLN]] a[i] = a[i % 2]; [[intel::ivdep(a, 2)]] for (int j = 0; j != 10; ++j) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_VAR_SFLN:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_VAR_SFLN]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_VAR_SFLN:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_INNER_VAR_SFLN]] a[i] = a[(i + j) % 10]; // CHECK: br label %for.cond{{[0-9]*}}, !llvm.loop ![[MD_LOOP_INNER_VAR_SFLN:[0-9]+]] } @@ -68,15 +68,15 @@ void ivdep_embedded_multiple_arrays() { // CHECK: %[[ARRAY_B:[0-9a-z]+]] = alloca [10 x i32] int b[10]; [[intel::ivdep(a, 3), intel::ivdep(b, 4)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_OUTER_MUL_ARRS:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_OUTER_MUL_ARRS]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_OUTER_MUL_ARRS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_OUTER_MUL_ARRS]] a[i] = a[i % 2]; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_OUTER_MUL_ARRS:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_OUTER_MUL_ARRS]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_OUTER_MUL_ARRS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_OUTER_MUL_ARRS]] b[i] = b[i % 2]; [[intel::ivdep(2)]] for (int j = 0; j != 10; ++j) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_INNER_MUL_ARRS:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_INNER_MUL_ARRS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_INNER_MUL_ARRS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_INNER_MUL_ARRS:[0-9]+]] a[i] = b[(i + j) % 10]; // CHECK: br label %for.cond{{[0-9]*}}, !llvm.loop ![[MD_LOOP_INNER_MUL_ARRS:[0-9]+]] } @@ -97,15 +97,15 @@ void ivdep_embedded_multiple_arrays_global() { // CHECK: %[[ARRAY_B:[0-9a-z]+]] = alloca [10 x i32] int b[10]; [[intel::ivdep(a)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_OUTER_MUL_ARRS_GLOB:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_OUTER_MUL_ARRS_GLOB]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_OUTER_MUL_ARRS_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_OUTER_MUL_ARRS_GLOB]] a[i] = a[i % 2]; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}{{[[:space:]]}} - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}{{[[:space:]]}} + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}{{[[:space:]]}} + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}{{[[:space:]]}} b[i] = b[i % 2]; [[intel::ivdep]] for (int j = 0; j != 10; ++j) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_INNER_MUL_ARRS_GLOB:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_INNER_MUL_ARRS_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_INNER_MUL_ARRS_GLOB:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_INNER_MUL_ARRS_GLOB:[0-9]+]] a[i] = b[(i + j) % 10]; // CHECK: br label %for.cond{{[0-9]*}}, !llvm.loop ![[MD_LOOP_INNER_MUL_ARRS_GLOB:[0-9]+]] } @@ -118,13 +118,13 @@ void ivdep_embedded_multiple_arrays_global() { void ivdep_embedded_multiple_dimensions() { int a[10]; [[intel::ivdep]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_DIM_1_MUL_DIMS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_DIM_1_MUL_DIMS:[0-9]+]] a[i] = i; [[intel::ivdep]] for (int j = 0; j != 10; ++j) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_DIM_2_MUL_DIMS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_DIM_2_MUL_DIMS:[0-9]+]] a[j] += j; [[intel::ivdep]] for (int k = 0; k != 10; ++k) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_DIM_3_MUL_DIMS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_DIM_3_MUL_DIMS:[0-9]+]] a[k] += k; // CHECK: br label %for.cond{{[0-9]*}}, !llvm.loop ![[MD_LOOP_DIM_3_MUL_DIMS:[0-9]+]] } diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp index 3ce59a1383510..f94318bc063dc 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp @@ -9,9 +9,9 @@ void ivdep_no_param() { // CHECK: %[[ARRAY_B:[0-9a-z]+]] = alloca [10 x i32] int b[10]; [[intel::ivdep]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_NO_PARAM:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_NO_PARAM:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_NO_PARAM:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_NO_PARAM:[0-9]+]] b[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_NO_PARAM:[0-9]+]] } @@ -29,12 +29,12 @@ void ivdep_no_param_multiple_geps() { // CHECK: %[[TMP:[0-9a-z]+]] = alloca i32 int t; [[intel::ivdep]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_MUL_GEPS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_MUL_GEPS:[0-9]+]] t = a[i]; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_MUL_GEPS:[0-9]+]] - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_MUL_GEPS]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_MUL_GEPS:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_MUL_GEPS]] a[i] = b[i]; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_MUL_GEPS]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_MUL_GEPS]] b[i] = t; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_MUL_GEPS:[0-9]+]] } @@ -49,9 +49,9 @@ void ivdep_safelen() { // CHECK: %[[ARRAY_B:[0-9a-z]+]] = alloca [10 x i32] int b[10]; [[intel::ivdep(5)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_SAFELEN:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_SAFELEN:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_SAFELEN:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_SAFELEN:[0-9]+]] b[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_SAFELEN:[0-9]+]] } @@ -67,9 +67,9 @@ void ivdep_conflicting_safelen() { int b[10]; [[intel::ivdep(5)]] [[intel::ivdep(4)]] for (int i = 0; i != 10; ++i) { - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_A]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_CONFL_SAFELEN:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_A]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_A_CONFL_SAFELEN:[0-9]+]] a[i] = 0; - // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32]* %[[ARRAY_B]], i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_CONFL_SAFELEN:[0-9]+]] + // CHECK: %{{[0-9a-z]+}} = getelementptr inbounds [10 x i32], [10 x i32] addrspace(4)* %[[ARRAY_B]].ascast, i64 0, i64 %{{[0-9a-z]+}}, !llvm.index.group ![[IDX_GROUP_B_CONFL_SAFELEN:[0-9]+]] b[i] = 0; // CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_CONFL_SAFELEN:[0-9]+]] } diff --git a/clang/test/CodeGenSYCL/intel-fpga-local.cpp b/clang/test/CodeGenSYCL/intel-fpga-local.cpp index 55cde17b43358..75d4b32a01307 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-local.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-local.cpp @@ -46,72 +46,72 @@ void attrs_on_static() { } void attrs_on_var() { - // CHECK-DEVICE: %[[VAR_NUMBANKS:[0-9]+]] = bitcast{{.*}}%numbanks - // CHECK-DEVICE: %[[VAR_NUMBANKS1:numbanks[0-9]+]] = bitcast{{.*}}%numbanks + // CHECK-DEVICE: %[[VAR_NUMBANKS:numbanks.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %numbanks.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_NUMBANKS1:numbanks.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_NUMBANKS]] to i8* // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[VAR_NUMBANKS1]],{{.*}}[[ANN_numbanks_4]] int numbanks [[intel::numbanks(4)]]; - // CHECK-DEVICE: %[[VAR_REGISTER:[0-9]+]] = bitcast{{.*}}%reg - // CHECK-DEVICE: %[[VAR_REGISTER1:reg[0-9]+]] = bitcast{{.*}}%reg + // CHECK-DEVICE: %[[VAR_REGISTER:reg.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %reg.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_REGISTER1:reg.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_REGISTER]] to i8* // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[VAR_REGISTER1]],{{.*}}[[ANN_register]] int reg [[intel::fpga_register]]; - // CHECK-DEVICE: %[[VAR_MEMORY:[0-9]+]] = bitcast{{.*}}%memory - // CHECK-DEVICE: %[[VAR_MEMORY1:memory[0-9]+]] = bitcast{{.*}}%memory + // CHECK-DEVICE: %[[VAR_MEMORY:memory.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %memory.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_MEMORY1:memory.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_MEMORY]] to i8* // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[VAR_MEMORY1]],{{.*}}[[ANN_memory_default]] int memory [[intel::fpga_memory]]; - // CHECK-DEVICE: %[[VAR_SIZE_MLAB:[0-9]+]] = bitcast{{.*}}size_mlab - // CHECK-DEVICE: %[[VAR_SIZE_MLAB1:size_mlab[0-9]+]] = bitcast{{.*}}size_mlab + // CHECK-DEVICE: %[[VAR_SIZE_MLAB:size_mlab.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %size_mlab.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_SIZE_MLAB1:size_mlab.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_SIZE_MLAB]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_SIZE_MLAB1]],{{.*}}[[ANN_mlab_sizeinfo_500]] [[intel::fpga_memory("MLAB")]] int size_mlab[500]; - // CHECK-DEVICE: %[[VAR_size_blockram:[0-9]+]] = bitcast{{.*}}size_blockram - // CHECK-DEVICE: %[[VAR_size_blockram1:size_blockram[0-9]+]] = bitcast{{.*}}size_blockram - // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_size_blockram1]],{{.*}}[[ANN_blockram_sizeinfo_10_2]] + // CHECK-DEVICE: %[[VAR_SIZE_BLOCKRAM:size_blockram.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %size_blockram.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_SIZE_BLOCKRAM1:size_blockram.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_SIZE_BLOCKRAM]] to i8* + // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_SIZE_BLOCKRAM1]],{{.*}}[[ANN_blockram_sizeinfo_10_2]] [[intel::fpga_memory("BLOCK_RAM")]] int size_blockram[10][2]; - // CHECK-DEVICE: %[[VAR_BANKWIDTH:[0-9]+]] = bitcast{{.*}}%bankwidth - // CHECK-DEVICE: %[[VAR_BANKWIDTH1:bankwidth[a-z0-9]+]] = bitcast{{.*}}%bankwidth + // CHECK-DEVICE: %[[VAR_BANKWIDTH:bankwidth.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %bankwidth.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_BANKWIDTH1:bankwidth.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_BANKWIDTH]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_BANKWIDTH1]],{{.*}}[[ANN_bankwidth_4]] int bankwidth [[intel::bankwidth(4)]]; - // CHECK-DEVICE: %[[VAR_PRIV_COPIES:[0-9]+]] = bitcast{{.*}}%priv_copies - // CHECK-DEVICE: %[[VAR_PRIV_COPIES1:priv_copies[0-9]+]] = bitcast{{.*}}%priv_copies + // CHECK-DEVICE: %[[VAR_PRIV_COPIES:priv_copies.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %priv_copies.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_PRIV_COPIES1:priv_copies.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_PRIV_COPIES]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_PRIV_COPIES1]],{{.*}}[[ANN_private_copies_8]] int priv_copies [[intel::private_copies(8)]]; - // CHECK-DEVICE: %[[VAR_SINGLEPUMP:[0-9]+]] = bitcast{{.*}}%singlepump - // CHECK-DEVICE: %[[VAR_SINGLEPUMP1:singlepump[0-9]+]] = bitcast{{.*}}%singlepump + // CHECK-DEVICE: %[[VAR_SINGLEPUMP:singlepump.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %singlepump.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_SINGLEPUMP1:singlepump.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_SINGLEPUMP]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_SINGLEPUMP1]],{{.*}}[[ANN_singlepump]] int singlepump [[intel::singlepump]]; - // CHECK-DEVICE: %[[VAR_DOUBLEPUMP:[0-9]+]] = bitcast{{.*}}%doublepump - // CHECK-DEVICE: %[[VAR_DOUBLEPUMP1:doublepump[0-9]+]] = bitcast{{.*}}%doublepump + // CHECK-DEVICE: %[[VAR_DOUBLEPUMP:doublepump.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %doublepump.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_DOUBLEPUMP1:doublepump.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_DOUBLEPUMP]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_DOUBLEPUMP1]],{{.*}}[[ANN_doublepump]] int doublepump [[intel::doublepump]]; - // CHECK-DEVICE: %[[VAR_MERGE_DEPTH:[0-9]+]] = bitcast{{.*}}%merge_depth - // CHECK-DEVICE: %[[VAR_MERGE_DEPTH1:merge_depth[0-9]+]] = bitcast{{.*}}%merge_depth + // CHECK-DEVICE: %[[VAR_MERGE_DEPTH:merge_depth.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %merge_depth.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_MERGE_DEPTH1:merge_depth.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_MERGE_DEPTH]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_MERGE_DEPTH1]],{{.*}}[[ANN_merge_depth]] int merge_depth [[intel::merge("foo", "depth")]]; - // CHECK-DEVICE: %[[VAR_MERGE_WIDTH:[0-9]+]] = bitcast{{.*}}%merge_width - // CHECK-DEVICE: %[[VAR_MERGE_WIDTH1:merge_width[0-9]+]] = bitcast{{.*}}%merge_width + // CHECK-DEVICE: %[[VAR_MERGE_WIDTH:merge_width.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %merge_width.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_MERGE_WIDTH1:merge_width.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_MERGE_WIDTH]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_MERGE_WIDTH1]],{{.*}}[[ANN_merge_width]] int merge_width [[intel::merge("bar", "width")]]; - // CHECK-DEVICE: %[[VAR_MAXREPL:[0-9]+]] = bitcast{{.*}}%max_repl - // CHECK-DEVICE: %[[VAR_MAXREPL1:max_repl[0-9]+]] = bitcast{{.*}}%max_repl + // CHECK-DEVICE: %[[VAR_MAXREPL:max_repl.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %max_repl.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_MAXREPL1:max_repl.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_MAXREPL]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_MAXREPL1]],{{.*}}[[ANN_max_replicates_2]] int max_repl [[intel::max_replicates(2)]]; - // CHECK-DEVICE: %[[VAR_DUALPORT:[0-9]+]] = bitcast{{.*}}%dualport - // CHECK-DEVICE: %[[VAR_DUALPORT1:dualport[0-9]+]] = bitcast{{.*}}%dualport + // CHECK-DEVICE: %[[VAR_DUALPORT:dualport.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %dualport.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_DUALPORT1:dualport.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_DUALPORT]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_DUALPORT1]],{{.*}}[[ANN_simple_dual_port]] int dualport [[intel::simple_dual_port]]; - // CHECK-DEVICE: %[[VAR_BANKBITS:[0-9]+]] = bitcast{{.*}}%bankbits - // CHECK-DEVICE: %[[VAR_BANKBITS1:bankbits[0-9]+]] = bitcast{{.*}}%bankbits + // CHECK-DEVICE: %[[VAR_BANKBITS:bankbits.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %bankbits.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_BANKBITS1:bankbits.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_BANKBITS]] to i8* // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[VAR_BANKBITS1]],{{.*}}[[ANN_bankbits_4_5]] int bankbits [[intel::bank_bits(4, 5)]]; - // CHECK-DEVICE: %[[VAR_BANKBITS_NUMBANKS:[0-9]+]] = bitcast{{.*}}%bankbits_numbanks_mlab - // CHECK-DEVICE: %[[VAR_BANKBITS_NUMBANKS1:bankbits_numbanks_mlab[0-9]+]] = bitcast{{.*}}%bankbits_numbanks_mlab + // CHECK-DEVICE: %[[VAR_BANKBITS_NUMBANKS:bankbits_numbanks_mlab.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %bankbits_numbanks_mlab.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_BANKBITS_NUMBANKS1:bankbits_numbanks_mlab.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_BANKBITS_NUMBANKS]] to i8* // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[VAR_BANKBITS_NUMBANKS1]],{{.*}}[[ANN_bankbits_numbanks_mlab]] [[intel::bank_bits(5, 4, 3), intel::numbanks(8), intel::fpga_memory("MLAB")]] int bankbits_numbanks_mlab; - // CHECK-DEVICE: %[[VAR_BANK_BITS_WIDTH:[0-9]+]] = bitcast{{.*}}%bank_bits_width - // CHECK-DEVICE: %[[VAR_BANK_BITS_WIDTH1:bank_bits_width[0-9]+]] = bitcast{{.*}}%bank_bits_width + // CHECK-DEVICE: %[[VAR_BANK_BITS_WIDTH:bank_bits_width.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %bank_bits_width.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_BANK_BITS_WIDTH1:bank_bits_width.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_BANK_BITS_WIDTH]] to i8* // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[VAR_BANK_BITS_WIDTH1]],{{.*}}[[ANN_bankbits_bankwidth]] [[intel::bank_bits(0), intel::bankwidth(16)]] int bank_bits_width[10][2]; - // CHECK-DEVICE: %[[VAR_FP2D:[0-9]+]] = bitcast{{.*}}%force_p2d - // CHECK-DEVICE: %[[VAR_FP2D1:force_p2d[0-9]+]] = bitcast{{.*}}%force_p2d + // CHECK-DEVICE: %[[VAR_FP2D:force_p2d.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %force_p2d.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[VAR_FP2D1:force_p2d.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[VAR_FP2D]] to i8* // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[VAR_FP2D1]],{{.*}}[[ANN_force_pow2_depth_0]] int force_p2d [[intel::force_pow2_depth(0)]]; } @@ -136,49 +136,49 @@ void attrs_on_struct() { } s; // CHECK-DEVICE: %[[FIELD_NUMBANKS:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_NUMBANKS]]{{.*}}[[ANN_numbanks_4]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_NUMBANKS]]{{.*}}[[ANN_numbanks_4]] s.numbanks = 0; // CHECK-DEVICE: %[[FIELD_REGISTER:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_REGISTER]]{{.*}}[[ANN_register]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_REGISTER]]{{.*}}[[ANN_register]] s.reg = 0; // CHECK-DEVICE: %[[FIELD_MEM_DEFAULT:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_MEM_DEFAULT]]{{.*}}[[ANN_memory_default]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MEM_DEFAULT]]{{.*}}[[ANN_memory_default]] s.memory = 0; // CHECK-DEVICE: %[[FIELD_MEM_BLOCKRAM:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_MEM_BLOCKRAM]]{{.*}}[[ANN_memory_blockram]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MEM_BLOCKRAM]]{{.*}}[[ANN_memory_blockram]] s.memory_blockram = 0; // CHECK-DEVICE: %[[FIELD_MEM_MLAB:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_MEM_MLAB]]{{.*}}[[ANN_memory_mlab]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MEM_MLAB]]{{.*}}[[ANN_memory_mlab]] s.memory_mlab = 0; // CHECK-DEVICE: %[[FIELD_BANKWIDTH:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_BANKWIDTH]]{{.*}}[[ANN_bankwidth_4]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_BANKWIDTH]]{{.*}}[[ANN_bankwidth_4]] s.bankwidth = 0; // CHECK-DEVICE: %[[FIELD_PRIV_COPIES:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_PRIV_COPIES]]{{.*}}[[ANN_private_copies_8]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_PRIV_COPIES]]{{.*}}[[ANN_private_copies_8]] s.privatecopies = 0; // CHECK-DEVICE: %[[FIELD_SINGLEPUMP:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_SINGLEPUMP]]{{.*}}[[ANN_singlepump]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_SINGLEPUMP]]{{.*}}[[ANN_singlepump]] s.singlepump = 0; // CHECK-DEVICE: %[[FIELD_DOUBLEPUMP:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_DOUBLEPUMP]]{{.*}}[[ANN_doublepump]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_DOUBLEPUMP]]{{.*}}[[ANN_doublepump]] s.doublepump = 0; // CHECK-DEVICE: %[[FIELD_MERGE_DEPTH:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_MERGE_DEPTH]]{{.*}}[[ANN_merge_depth]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MERGE_DEPTH]]{{.*}}[[ANN_merge_depth]] s.merge_depth = 0; // CHECK-DEVICE: %[[FIELD_MERGE_WIDTH:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_MERGE_WIDTH]]{{.*}}[[ANN_merge_width]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MERGE_WIDTH]]{{.*}}[[ANN_merge_width]] s.merge_width = 0; // CHECK-DEVICE: %[[FIELD_MAX_REPLICATES:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_MAX_REPLICATES]]{{.*}}[[ANN_max_replicates_2]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MAX_REPLICATES]]{{.*}}[[ANN_max_replicates_2]] s.maxreplicates = 0; // CHECK-DEVICE: %[[FIELD_DUALPORT:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_DUALPORT]]{{.*}}[[ANN_simple_dual_port]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_DUALPORT]]{{.*}}[[ANN_simple_dual_port]] s.dualport = 0; // CHECK-DEVICE: %[[FIELD_BANKBITS:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_BANKBITS]]{{.*}}[[ANN_bankbits_4_5]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_BANKBITS]]{{.*}}[[ANN_bankbits_4_5]] s.bankbits = 0; // CHECK-DEVICE: %[[FIELD_FP2D:.*]] = getelementptr inbounds %struct.{{.*}}.attrs_on_struct{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_FP2D]]{{.*}}[[ANN_force_pow2_depth_1]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_FP2D]]{{.*}}[[ANN_force_pow2_depth_1]] s.force_p2d = 0; } @@ -187,23 +187,29 @@ void attrs_on_struct() { template void attrs_with_template_param() { - // CHECK-DEVICE: %[[TEMPL_NUMBANKS:numbanks[0-9]+]] = bitcast{{.*}}%numbanks - // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[TEMPL_NUMBANKS]],{{.*}}[[ANN_numbanks_4]] + // CHECK-DEVICE: %[[TEMPL_NUMBANKS:numbanks.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %numbanks.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[TEMPL_NUMBANKS1:numbanks.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[TEMPL_NUMBANKS]] to i8* + // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[TEMPL_NUMBANKS1]],{{.*}}[[ANN_numbanks_4]] int numbanks [[intel::numbanks(A)]]; - // CHECK-DEVICE: %[[TEMPL_BANKWIDTH:bankwidth[a-z0-9]+]] = bitcast{{.*}}%bankwidth - // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[TEMPL_BANKWIDTH]],{{.*}}[[ANN_bankwidth_4]] + // CHECK-DEVICE: %[[TEMPL_BANKWIDTH:bankwidth.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %bankwidth.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[TEMPL_BANKWIDTH1:bankwidth.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[TEMPL_BANKWIDTH]] to i8* + // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[TEMPL_BANKWIDTH1]],{{.*}}[[ANN_bankwidth_4]] int bankwidth [[intel::bankwidth(A)]]; - // CHECK-DEVICE: %[[TEMPL_PRIV_COPIES:priv_copies[0-9]+]] = bitcast{{.*}}%priv_copies - // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[TEMPL_PRIV_COPIES]],{{.*}}[[ANN_private_copies_4]] + // CHECK-DEVICE: %[[TEMPL_PRIV_COPIES:priv_copies.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %priv_copies.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[TEMPL_PRIV_COPIES1:priv_copies.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[TEMPL_PRIV_COPIES]] to i8* + // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[TEMPL_PRIV_COPIES1]],{{.*}}[[ANN_private_copies_4]] int priv_copies [[intel::private_copies(A)]]; - // CHECK-DEVICE: %[[TEMPL_MAXREPL:max_repl[0-9]+]] = bitcast{{.*}}%max_repl - // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[TEMPL_MAXREPL]],{{.*}}[[ANN_max_replicates_4]] + // CHECK-DEVICE: %[[TEMPL_MAXREPL:max_repl.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %max_repl.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[TEMPL_MAXREPL1:max_repl.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[TEMPL_MAXREPL]] to i8* + // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[TEMPL_MAXREPL1]],{{.*}}[[ANN_max_replicates_4]] int max_repl [[intel::max_replicates(A)]]; - // CHECK-DEVICE: %[[TEMPL_BANKBITS:bankbits[0-9]+]] = bitcast{{.*}}%bankbits - // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[TEMPL_BANKBITS]],{{.*}}[[ANN_bankbits_4_5]] + // CHECK-DEVICE: %[[TEMPL_BANKBITS:bankbits.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %bankbits.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[TEMPL_BANKBITS1:bankbits.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[TEMPL_BANKBITS]] to i8* + // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[TEMPL_BANKBITS1]],{{.*}}[[ANN_bankbits_4_5]] int bankbits [[intel::bank_bits(A, B)]]; - // CHECK-DEVICE: %[[TEMPL_FP2D:force_p2d[0-9]+]] = bitcast{{.*}}%force_p2d - // CHECK-DEVICE: @llvm.var.annotation{{.*}}%[[TEMPL_FP2D]]{{.*}}[[ANN_force_pow2_depth_1]] + // CHECK-DEVICE: %[[TEMPL_FP2D:force_p2d.ascast[0-9]+]] = bitcast {{.*}} addrspace(4)* %force_p2d.ascast to i8 addrspace(4)* + // CHECK-DEVICE: %[[TEMPL_FP2D1:force_p2d.ascast[0-9]+]] = addrspacecast i8 addrspace(4)* %[[TEMPL_FP2D]] to i8* + // CHECK-DEVICE: llvm.var.annotation{{.*}}%[[TEMPL_FP2D1]],{{.*}}[[ANN_force_pow2_depth_1]] int force_p2d [[intel::force_pow2_depth(C)]]; struct templ_on_struct_fields { @@ -216,22 +222,22 @@ void attrs_with_template_param() { } s; // CHECK-DEVICE: %[[FIELD_NUMBANKS:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_NUMBANKS]]{{.*}}[[ANN_numbanks_4]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_NUMBANKS]]{{.*}}[[ANN_numbanks_4]] s.numbanks = 0; // CHECK-DEVICE: %[[FIELD_BANKWIDTH:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_BANKWIDTH]]{{.*}}[[ANN_bankwidth_4]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_BANKWIDTH]]{{.*}}[[ANN_bankwidth_4]] s.bankwidth = 0; // CHECK-DEVICE: %[[FIELD_PRIV_COPIES:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_PRIV_COPIES]]{{.*}}[[ANN_private_copies_4]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_PRIV_COPIES]]{{.*}}[[ANN_private_copies_4]] s.privatecopies = 0; // CHECK-DEVICE: %[[FIELD_MAX_REPLICATES:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_MAX_REPLICATES]]{{.*}}[[ANN_max_replicates_4]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_MAX_REPLICATES]]{{.*}}[[ANN_max_replicates_4]] s.maxreplicates = 0; // CHECK-DEVICE: %[[FIELD_BANKBITS:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_BANKBITS]]{{.*}}[[ANN_bankbits_4_5]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_BANKBITS]]{{.*}}[[ANN_bankbits_4_5]] s.bankbits = 0; // CHECK-DEVICE: %[[FIELD_FP2D:.*]] = getelementptr inbounds %struct.{{.*}}.templ_on_struct_fields{{.*}} - // CHECK-DEVICE: call i32* @llvm.ptr.annotation.p0i32{{.*}}%[[FIELD_FP2D]]{{.*}}[[ANN_force_pow2_depth_1]] + // CHECK-DEVICE: call i32 addrspace(4)* @llvm.ptr.annotation.p4i32{{.*}}%[[FIELD_FP2D]]{{.*}}[[ANN_force_pow2_depth_1]] s.force_p2d = 0; } diff --git a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp index 74f3f8269c79f..72f4d934e2e33 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp @@ -7,7 +7,7 @@ // CHECK: br label %for.cond2, !llvm.loop ![[MD_MC_2:[0-9]+]] // CHECK: br label %for.cond, !llvm.loop ![[MD_LC:[0-9]+]] // CHECK: br label %for.cond2, !llvm.loop ![[MD_LC_2:[0-9]+]] -// CHECK: br label %for.cond12, !llvm.loop ![[MD_LC_3:[0-9]+]] +// CHECK: br label %for.cond13, !llvm.loop ![[MD_LC_3:[0-9]+]] // CHECK: br label %for.cond, !llvm.loop ![[MD_MI:[0-9]+]] // CHECK: br label %for.cond2, !llvm.loop ![[MD_MI_2:[0-9]+]] // CHECK: br label %for.cond, !llvm.loop ![[MD_SI:[0-9]+]] diff --git a/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp b/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp index 797d3de3a1199..141462c4c2733 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp @@ -27,44 +27,43 @@ void foo(float *A, int *B, State *C, State &D) { // CHECK-DAG: [[F:%.*]] = alloca double // CHECK-DAG: [[f:%.*]] = alloca double addrspace(4)* - // CHECK-DAG: [[A:%[0-9]+]] = load float addrspace(4)*, float addrspace(4)** [[Aaddr]] + // CHECK-DAG: [[A:%[0-9]+]] = load float addrspace(4)*, float addrspace(4)* addrspace(4)* [[Aaddr]] // CHECK-DAG: [[PTR1:%[0-9]+]] = call float addrspace(4)* @llvm.ptr.annotation{{.*}}[[A]]{{.*}}[[ANN1]]{{.*}}[[ATT:#[0-9]+]] - // CHECK-DAG: store float addrspace(4)* [[PTR1]], float addrspace(4)** %x + // CHECK-DAG: store float addrspace(4)* [[PTR1]], float addrspace(4)* addrspace(4)* %x x = __builtin_intel_fpga_mem(A, PARAM_1 | PARAM_2, 0); - // CHECK-DAG: [[B:%[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[Baddr]] + // CHECK-DAG: [[B:%[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* [[Baddr]] // CHECK-DAG: [[PTR2:%[0-9]+]] = call i32 addrspace(4)* @llvm.ptr.annotation{{.*}}[[B]]{{.*}}[[ANN1]]{{.*}}[[ATT:#[0-9]+]] - // CHECK-DAG: store i32 addrspace(4)* [[PTR2]], i32 addrspace(4)** %y + // CHECK-DAG: store i32 addrspace(4)* [[PTR2]], i32 addrspace(4)* addrspace(4)* %y y = __builtin_intel_fpga_mem(B, PARAM_1 | PARAM_2, 0); - // CHECK-DAG: [[C:%[0-9]+]] = load [[STRUCT]] addrspace(4)*, [[STRUCT]] addrspace(4)** [[Caddr]] + // CHECK-DAG: [[C:%[0-9]+]] = load [[STRUCT]] addrspace(4)*, [[STRUCT]] addrspace(4)* addrspace(4)* [[Caddr]] // CHECK-DAG: [[PTR3:%[0-9]+]] = call [[STRUCT]] addrspace(4)* @llvm.ptr.annotation{{.*}}[[C]]{{.*}}[[ANN1]]{{.*}}[[ATT:#[0-9]+]] - // CHECK-DAG: store [[STRUCT]] addrspace(4)* [[PTR3]], [[STRUCT]] addrspace(4)** %z + // CHECK-DAG: store [[STRUCT]] addrspace(4)* [[PTR3]], [[STRUCT]] addrspace(4)* addrspace(4)* %z z = __builtin_intel_fpga_mem(C, PARAM_1 | PARAM_2, 0); - // CHECK-DAG: [[A2:%[0-9]+]] = load float addrspace(4)*, float addrspace(4)** [[Aaddr]] + // CHECK-DAG: [[A2:%[0-9]+]] = load float addrspace(4)*, float addrspace(4)* addrspace(4)* [[Aaddr]] // CHECK-DAG: [[PTR4:%[0-9]+]] = call float addrspace(4)* @llvm.ptr.annotation{{.*}}[[A2]]{{.*}}[[ANN2]]{{.*}}[[ATT:#[0-9]+]] - // CHECK-DAG: store float addrspace(4)* [[PTR4]], float addrspace(4)** %x + // CHECK-DAG: store float addrspace(4)* [[PTR4]], float addrspace(4)* addrspace(4)* %x x = __builtin_intel_fpga_mem(A, PARAM_1 | PARAM_2, 127); - // CHECK-DAG: [[B2:%[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)** [[Baddr]] + // CHECK-DAG: [[B2:%[0-9]+]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* [[Baddr]] // CHECK-DAG: [[PTR5:%[0-9]+]] = call i32 addrspace(4)* @llvm.ptr.annotation{{.*}}[[B2]]{{.*}}[[ANN2]]{{.*}}[[ATT:#[0-9]+]] - // CHECK-DAG: store i32 addrspace(4)* [[PTR5]], i32 addrspace(4)** %y + // CHECK-DAG: store i32 addrspace(4)* [[PTR5]], i32 addrspace(4)* addrspace(4)* %y y = __builtin_intel_fpga_mem(B, PARAM_1 | PARAM_2, 127); - // CHECK-DAG: [[C2:%[0-9]+]] = load [[STRUCT]] addrspace(4)*, [[STRUCT]] addrspace(4)** [[Caddr]] + // CHECK-DAG: [[C2:%[0-9]+]] = load [[STRUCT]] addrspace(4)*, [[STRUCT]] addrspace(4)* addrspace(4)* [[Caddr]] // CHECK-DAG: [[PTR6:%[0-9]+]] = call [[STRUCT]] addrspace(4)* @llvm.ptr.annotation{{.*}}[[C2]]{{.*}}[[ANN2]]{{.*}}[[ATT:#[0-9]+]] - // CHECK-DAG: store [[STRUCT]] addrspace(4)* [[PTR6]], [[STRUCT]] addrspace(4)** %z + // CHECK-DAG: store [[STRUCT]] addrspace(4)* [[PTR6]], [[STRUCT]] addrspace(4)* addrspace(4)* %z z = __builtin_intel_fpga_mem(C, PARAM_1 | PARAM_2, 127); - // CHECK-DAG: [[D:%[0-9]+]] = load [[STRUCT]] addrspace(4)*, [[STRUCT]] addrspace(4)** [[Daddr]] + // CHECK-DAG: [[D:%[0-9]+]] = load [[STRUCT]] addrspace(4)*, [[STRUCT]] addrspace(4)* addrspace(4)* [[Daddr]] // CHECK-DAG: [[PTR7:%[0-9]+]] = call [[STRUCT]] addrspace(4)* @llvm.ptr.annotation{{.*}}[[D]]{{.*}}[[ANN2]]{{.*}}[[ATT:#[0-9]+]] - // CHECK-DAG: store [[STRUCT]] addrspace(4)* [[PTR7]], [[STRUCT]] addrspace(4)** %z + // CHECK-DAG: store [[STRUCT]] addrspace(4)* [[PTR7]], [[STRUCT]] addrspace(4)* addrspace(4)* %z z = __builtin_intel_fpga_mem(&D, PARAM_1 | PARAM_2, 127); - // CHECK-DAG: [[PTR8:%[0-9]+]] = call double* @llvm.ptr.annotation{{.*}}[[F]]{{.*}}[[ANN2]]{{.*}}[[ATT:#[0-9]+]] - // CHECK-DAG: [[ASCAST:%[0-9]+]] = addrspacecast double* [[PTR8]] to double addrspace(4)* - // CHECK-DAG: store double addrspace(4)* [[ASCAST]], double addrspace(4)** [[f]] + // CHECK-DAG: [[PTR8:%[0-9]+]] = call double addrspace(4)* @llvm.ptr.annotation{{.*}}[[F]]{{.*}}[[ANN2]]{{.*}}[[ATT:#[0-9]+]] + // CHECK-DAG: store double addrspace(4)* [[PTR8]], double addrspace(4)* addrspace(4)* [[f]] f = __builtin_intel_fpga_mem(&F, PARAM_1 | PARAM_2, 127); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-reg.cpp b/clang/test/CodeGenSYCL/intel-fpga-reg.cpp index 066ccc4064933..380402592e1e1 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-reg.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-reg.cpp @@ -1,16 +1,15 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --functions ["['foo']"] // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s struct st { int a; float b; }; -// CHECK: %[[T_ST:struct[a-zA-Z0-9_.]*.st]] = type { i32, float } union un { int a; char c[4]; }; -// CHECK: %[[T_UN:union[a-zA-Z0-9_.]*.un]] = type { i32 } class A { public: @@ -23,123 +22,273 @@ class A { private: int m_val; }; -// CHECK: %[[T_CL:class[a-zA-Z0-9_.]*.A]] = type { i32 } typedef int myInt; -// CHECK: @.str = private unnamed_addr constant [25 x i8] c"__builtin_intel_fpga_reg\00", section "llvm.metadata" - +// CHECK-LABEL: define {{.*}} @_Z3foov( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[A:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[A_ASCAST:%.*]] = addrspacecast i32* [[A]] to i32 addrspace(4)* +// CHECK-NEXT: [[MYA:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[MYA_ASCAST:%.*]] = addrspacecast i32* [[MYA]] to i32 addrspace(4)* +// CHECK-NEXT: [[B:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[B_ASCAST:%.*]] = addrspacecast i32* [[B]] to i32 addrspace(4)* +// CHECK-NEXT: [[MYB:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[MYB_ASCAST:%.*]] = addrspacecast i32* [[MYB]] to i32 addrspace(4)* +// CHECK-NEXT: [[C:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[C_ASCAST:%.*]] = addrspacecast i32* [[C]] to i32 addrspace(4)* +// CHECK-NEXT: [[D:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[D_ASCAST:%.*]] = addrspacecast i32* [[D]] to i32 addrspace(4)* +// CHECK-NEXT: [[E:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[E_ASCAST:%.*]] = addrspacecast i32* [[E]] to i32 addrspace(4)* +// CHECK-NEXT: [[F:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[F_ASCAST:%.*]] = addrspacecast i32* [[F]] to i32 addrspace(4)* +// CHECK-NEXT: [[I:%.*]] = alloca [[STRUCT__ZTS2ST_ST:%.*]], align 4 +// CHECK-NEXT: [[I_ASCAST:%.*]] = addrspacecast %struct._ZTS2st.st* [[I]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: [[I2:%.*]] = alloca [[STRUCT__ZTS2ST_ST]], align 4 +// CHECK-NEXT: [[I2_ASCAST:%.*]] = addrspacecast %struct._ZTS2st.st* [[I2]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: [[II:%.*]] = alloca [[STRUCT__ZTS2ST_ST]], align 4 +// CHECK-NEXT: [[II_ASCAST:%.*]] = addrspacecast %struct._ZTS2st.st* [[II]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: [[AGG_TEMP:%.*]] = alloca [[STRUCT__ZTS2ST_ST]], align 4 +// CHECK-NEXT: [[AGG_TEMP_ASCAST:%.*]] = addrspacecast %struct._ZTS2st.st* [[AGG_TEMP]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: [[III:%.*]] = alloca [[STRUCT__ZTS2ST_ST]], align 4 +// CHECK-NEXT: [[III_ASCAST:%.*]] = addrspacecast %struct._ZTS2st.st* [[III]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: [[REF_TMP:%.*]] = alloca [[STRUCT__ZTS2ST_ST]], align 4 +// CHECK-NEXT: [[REF_TMP_ASCAST:%.*]] = addrspacecast %struct._ZTS2st.st* [[REF_TMP]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: [[AGG_TEMP2:%.*]] = alloca [[STRUCT__ZTS2ST_ST]], align 4 +// CHECK-NEXT: [[AGG_TEMP2_ASCAST:%.*]] = addrspacecast %struct._ZTS2st.st* [[AGG_TEMP2]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: [[IIII:%.*]] = alloca [[STRUCT__ZTS2ST_ST]] addrspace(4)*, align 8 +// CHECK-NEXT: [[IIII_ASCAST:%.*]] = addrspacecast [[STRUCT__ZTS2ST_ST]] addrspace(4)** [[IIII]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* addrspace(4)* +// CHECK-NEXT: [[U1:%.*]] = alloca [[UNION__ZTS2UN_UN:%.*]], align 4 +// CHECK-NEXT: [[U1_ASCAST:%.*]] = addrspacecast %union._ZTS2un.un* [[U1]] to [[UNION__ZTS2UN_UN]] addrspace(4)* +// CHECK-NEXT: [[U2:%.*]] = alloca [[UNION__ZTS2UN_UN]], align 4 +// CHECK-NEXT: [[U2_ASCAST:%.*]] = addrspacecast %union._ZTS2un.un* [[U2]] to [[UNION__ZTS2UN_UN]] addrspace(4)* +// CHECK-NEXT: [[U3:%.*]] = alloca [[UNION__ZTS2UN_UN]] addrspace(4)*, align 8 +// CHECK-NEXT: [[U3_ASCAST:%.*]] = addrspacecast [[UNION__ZTS2UN_UN]] addrspace(4)** [[U3]] to [[UNION__ZTS2UN_UN]] addrspace(4)* addrspace(4)* +// CHECK-NEXT: [[REF_TMP3:%.*]] = alloca [[UNION__ZTS2UN_UN]], align 4 +// CHECK-NEXT: [[REF_TMP3_ASCAST:%.*]] = addrspacecast %union._ZTS2un.un* [[REF_TMP3]] to [[UNION__ZTS2UN_UN]] addrspace(4)* +// CHECK-NEXT: [[AGG_TEMP4:%.*]] = alloca [[UNION__ZTS2UN_UN]], align 4 +// CHECK-NEXT: [[AGG_TEMP4_ASCAST:%.*]] = addrspacecast %union._ZTS2un.un* [[AGG_TEMP4]] to [[UNION__ZTS2UN_UN]] addrspace(4)* +// CHECK-NEXT: [[CA:%.*]] = alloca [[CLASS__ZTS1A_A:%.*]], align 4 +// CHECK-NEXT: [[CA_ASCAST:%.*]] = addrspacecast %class._ZTS1A.A* [[CA]] to [[CLASS__ZTS1A_A]] addrspace(4)* +// CHECK-NEXT: [[CB:%.*]] = alloca [[CLASS__ZTS1A_A]], align 4 +// CHECK-NEXT: [[CB_ASCAST:%.*]] = addrspacecast %class._ZTS1A.A* [[CB]] to [[CLASS__ZTS1A_A]] addrspace(4)* +// CHECK-NEXT: [[AGG_TEMP5:%.*]] = alloca [[CLASS__ZTS1A_A]], align 4 +// CHECK-NEXT: [[AGG_TEMP5_ASCAST:%.*]] = addrspacecast %class._ZTS1A.A* [[AGG_TEMP5]] to [[CLASS__ZTS1A_A]] addrspace(4)* +// CHECK-NEXT: [[AP:%.*]] = alloca i32 addrspace(4)*, align 8 +// CHECK-NEXT: [[AP_ASCAST:%.*]] = addrspacecast i32 addrspace(4)** [[AP]] to i32 addrspace(4)* addrspace(4)* +// CHECK-NEXT: [[BP:%.*]] = alloca i32 addrspace(4)*, align 8 +// CHECK-NEXT: [[BP_ASCAST:%.*]] = addrspacecast i32 addrspace(4)** [[BP]] to i32 addrspace(4)* addrspace(4)* +// CHECK-NEXT: [[TMP0:%.*]] = bitcast i32* [[A]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP0]]) [[ATTR6:#.*]] +// CHECK-NEXT: store i32 123, i32 addrspace(4)* [[A_ASCAST]], align 4, [[TBAA9:!tbaa !.*]] +// CHECK-NEXT: [[TMP1:%.*]] = bitcast i32* [[MYA]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP1]]) [[ATTR6]] +// CHECK-NEXT: store i32 321, i32 addrspace(4)* [[MYA_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP2:%.*]] = bitcast i32* [[B]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP2]]) [[ATTR6]] +// CHECK-NEXT: [[TMP3:%.*]] = load i32, i32 addrspace(4)* [[A_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.annotation.i32(i32 [[TMP3]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY:\[[0-9]+ x i8\]]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 265) +// CHECK-NEXT: store i32 [[TMP4]], i32 addrspace(4)* [[B_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP5:%.*]] = bitcast i32* [[MYB]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP5]]) [[ATTR6]] +// CHECK-NEXT: [[TMP6:%.*]] = load i32, i32 addrspace(4)* [[MYA_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP7:%.*]] = call i32 @llvm.annotation.i32(i32 [[TMP6]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 266) +// CHECK-NEXT: store i32 [[TMP7]], i32 addrspace(4)* [[MYB_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP8:%.*]] = bitcast i32* [[C]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP8]]) [[ATTR6]] +// CHECK-NEXT: [[TMP9:%.*]] = call i32 @llvm.annotation.i32(i32 1073741824, i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 267) +// CHECK-NEXT: [[TMP10:%.*]] = bitcast i32 [[TMP9]] to float +// CHECK-NEXT: [[CONV:%.*]] = fptosi float [[TMP10]] to i32 +// CHECK-NEXT: store i32 [[CONV]], i32 addrspace(4)* [[C_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP11:%.*]] = bitcast i32* [[D]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP11]]) [[ATTR6]] +// CHECK-NEXT: [[TMP12:%.*]] = load i32, i32 addrspace(4)* [[B_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP12]], 12 +// CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.annotation.i32(i32 [[ADD]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 268) +// CHECK-NEXT: [[TMP14:%.*]] = call i32 @llvm.annotation.i32(i32 [[TMP13]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 268) +// CHECK-NEXT: store i32 [[TMP14]], i32 addrspace(4)* [[D_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP15:%.*]] = bitcast i32* [[E]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP15]]) [[ATTR6]] +// CHECK-NEXT: [[TMP16:%.*]] = load i32, i32 addrspace(4)* [[A_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP17:%.*]] = load i32, i32 addrspace(4)* [[B_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[ADD1:%.*]] = add nsw i32 [[TMP16]], [[TMP17]] +// CHECK-NEXT: [[TMP18:%.*]] = call i32 @llvm.annotation.i32(i32 [[ADD1]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 269) +// CHECK-NEXT: [[TMP19:%.*]] = call i32 @llvm.annotation.i32(i32 [[TMP18]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 269) +// CHECK-NEXT: store i32 [[TMP19]], i32 addrspace(4)* [[E_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP20:%.*]] = bitcast i32* [[F]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP20]]) [[ATTR6]] +// CHECK-NEXT: [[TMP21:%.*]] = load i32, i32 addrspace(4)* [[A_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP22:%.*]] = call i32 @llvm.annotation.i32(i32 [[TMP21]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 271) +// CHECK-NEXT: store i32 [[TMP22]], i32 addrspace(4)* [[F_ASCAST]], align 4, [[TBAA9]] +// CHECK-NEXT: [[TMP23:%.*]] = bitcast %struct._ZTS2st.st* [[I]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[TMP23]]) [[ATTR6]] +// CHECK-NEXT: [[TMP24:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[I_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p1i8.i64(i8 addrspace(4)* align 4 [[TMP24]], i8 addrspace(1)* align 4 bitcast (%struct._ZTS2st.st addrspace(1)* @__const._Z3foov.i to i8 addrspace(1)*), i64 8, i1 false) +// CHECK-NEXT: [[TMP25:%.*]] = bitcast %struct._ZTS2st.st* [[I2]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[TMP25]]) [[ATTR6]] +// CHECK-NEXT: [[TMP26:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[I2_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP27:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[I_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP26]], i8 addrspace(4)* align 4 [[TMP27]], i64 8, i1 false), !tbaa.struct !11 +// CHECK-NEXT: [[TMP28:%.*]] = bitcast %struct._ZTS2st.st* [[II]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[TMP28]]) [[ATTR6]] +// CHECK-NEXT: [[TMP29:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[AGG_TEMP_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP30:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[I_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP29]], i8 addrspace(4)* align 4 [[TMP30]], i64 8, i1 false), !tbaa.struct !11 +// CHECK-NEXT: [[TMP31:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[AGG_TEMP_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP32:%.*]] = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 addrspace(4)* [[TMP31]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 275, i8* null) +// CHECK-NEXT: [[TMP33:%.*]] = bitcast i8 addrspace(4)* [[TMP32]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: [[TMP34:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[II_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP35:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[TMP33]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP34]], i8 addrspace(4)* align 4 [[TMP35]], i64 8, i1 false) +// CHECK-NEXT: [[TMP36:%.*]] = bitcast %struct._ZTS2st.st* [[III]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[TMP36]]) [[ATTR6]] +// CHECK-NEXT: [[TMP37:%.*]] = bitcast %struct._ZTS2st.st* [[REF_TMP]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[TMP37]]) [[ATTR6]] +// CHECK-NEXT: [[TMP38:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[AGG_TEMP2_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP39:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[II_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP38]], i8 addrspace(4)* align 4 [[TMP39]], i64 8, i1 false), !tbaa.struct !11 +// CHECK-NEXT: [[TMP40:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[AGG_TEMP2_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP41:%.*]] = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 addrspace(4)* [[TMP40]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 277, i8* null) +// CHECK-NEXT: [[TMP42:%.*]] = bitcast i8 addrspace(4)* [[TMP41]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: [[TMP43:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[REF_TMP_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP44:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[TMP42]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP43]], i8 addrspace(4)* align 4 [[TMP44]], i64 8, i1 false) +// CHECK-NEXT: [[TMP45:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[III_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP46:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[REF_TMP_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP45]], i8 addrspace(4)* align 4 [[TMP46]], i64 8, i1 false), !tbaa.struct !11 +// CHECK-NEXT: [[TMP47:%.*]] = bitcast %struct._ZTS2st.st* [[REF_TMP]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 8, i8* [[TMP47]]) [[ATTR6]] +// CHECK-NEXT: [[TMP48:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)** [[IIII]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[TMP48]]) [[ATTR6]] +// CHECK-NEXT: [[TMP49:%.*]] = ptrtoint [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[III_ASCAST]] to i64 +// CHECK-NEXT: [[TMP50:%.*]] = call i64 @llvm.annotation.i64(i64 [[TMP49]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 279) +// CHECK-NEXT: [[TMP51:%.*]] = inttoptr i64 [[TMP50]] to [[STRUCT__ZTS2ST_ST]] addrspace(4)* +// CHECK-NEXT: store [[STRUCT__ZTS2ST_ST]] addrspace(4)* [[TMP51]], [[STRUCT__ZTS2ST_ST]] addrspace(4)* addrspace(4)* [[IIII_ASCAST]], align 8, [[TBAA5:!tbaa !.*]] +// CHECK-NEXT: [[TMP52:%.*]] = bitcast %union._ZTS2un.un* [[U1]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP52]]) [[ATTR6]] +// CHECK-NEXT: [[TMP53:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)* [[U1_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p1i8.i64(i8 addrspace(4)* align 4 [[TMP53]], i8 addrspace(1)* align 4 bitcast (%union._ZTS2un.un addrspace(1)* @__const._Z3foov.u1 to i8 addrspace(1)*), i64 4, i1 false) +// CHECK-NEXT: [[TMP54:%.*]] = bitcast %union._ZTS2un.un* [[U2]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP54]]) [[ATTR6]] +// CHECK-NEXT: [[TMP55:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)** [[U3]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[TMP55]]) [[ATTR6]] +// CHECK-NEXT: [[TMP56:%.*]] = bitcast %union._ZTS2un.un* [[REF_TMP3]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP56]]) [[ATTR6]] +// CHECK-NEXT: [[TMP57:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)* [[AGG_TEMP4_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP58:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)* [[U1_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP57]], i8 addrspace(4)* align 4 [[TMP58]], i64 4, i1 false), !tbaa.struct !14 +// CHECK-NEXT: [[TMP59:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)* [[AGG_TEMP4_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP60:%.*]] = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 addrspace(4)* [[TMP59]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 283, i8* null) +// CHECK-NEXT: [[TMP61:%.*]] = bitcast i8 addrspace(4)* [[TMP60]] to [[UNION__ZTS2UN_UN]] addrspace(4)* +// CHECK-NEXT: [[TMP62:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)* [[REF_TMP3_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP63:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)* [[TMP61]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP62]], i8 addrspace(4)* align 4 [[TMP63]], i64 8, i1 false) +// CHECK-NEXT: [[TMP64:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)* [[U2_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP65:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)* [[REF_TMP3_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP64]], i8 addrspace(4)* align 4 [[TMP65]], i64 4, i1 false), !tbaa.struct !14 +// CHECK-NEXT: [[TMP66:%.*]] = bitcast %union._ZTS2un.un* [[REF_TMP3]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP66]]) [[ATTR6]] +// CHECK-NEXT: [[TMP67:%.*]] = ptrtoint [[UNION__ZTS2UN_UN]] addrspace(4)* [[U2_ASCAST]] to i64 +// CHECK-NEXT: [[TMP68:%.*]] = call i64 @llvm.annotation.i64(i64 [[TMP67]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 285) +// CHECK-NEXT: [[TMP69:%.*]] = inttoptr i64 [[TMP68]] to [[UNION__ZTS2UN_UN]] addrspace(4)* +// CHECK-NEXT: store [[UNION__ZTS2UN_UN]] addrspace(4)* [[TMP69]], [[UNION__ZTS2UN_UN]] addrspace(4)* addrspace(4)* [[U3_ASCAST]], align 8, [[TBAA5]] +// CHECK-NEXT: [[TMP70:%.*]] = bitcast %class._ZTS1A.A* [[CA]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP70]]) [[ATTR6]] +// CHECK-NEXT: call spir_func void @_ZN1AC1Ei(%class._ZTS1A.A addrspace(4)* dereferenceable_or_null(4) [[CA_ASCAST]], i32 213) [[ATTR7:#.*]] +// CHECK-NEXT: [[TMP71:%.*]] = bitcast %class._ZTS1A.A* [[CB]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 4, i8* [[TMP71]]) [[ATTR6]] +// CHECK-NEXT: [[TMP72:%.*]] = bitcast [[CLASS__ZTS1A_A]] addrspace(4)* [[AGG_TEMP5_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP73:%.*]] = bitcast [[CLASS__ZTS1A_A]] addrspace(4)* [[CA_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP72]], i8 addrspace(4)* align 4 [[TMP73]], i64 4, i1 false), !tbaa.struct !16 +// CHECK-NEXT: [[TMP74:%.*]] = bitcast [[CLASS__ZTS1A_A]] addrspace(4)* [[AGG_TEMP5_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP75:%.*]] = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 addrspace(4)* [[TMP74]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 288, i8* null) +// CHECK-NEXT: [[TMP76:%.*]] = bitcast i8 addrspace(4)* [[TMP75]] to [[CLASS__ZTS1A_A]] addrspace(4)* +// CHECK-NEXT: [[TMP77:%.*]] = bitcast [[CLASS__ZTS1A_A]] addrspace(4)* [[CB_ASCAST]] to i8 addrspace(4)* +// CHECK-NEXT: [[TMP78:%.*]] = bitcast [[CLASS__ZTS1A_A]] addrspace(4)* [[TMP76]] to i8 addrspace(4)* +// CHECK-NEXT: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[TMP77]], i8 addrspace(4)* align 4 [[TMP78]], i64 8, i1 false) +// CHECK-NEXT: [[TMP79:%.*]] = bitcast i32 addrspace(4)** [[AP]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[TMP79]]) [[ATTR6]] +// CHECK-NEXT: store i32 addrspace(4)* [[A_ASCAST]], i32 addrspace(4)* addrspace(4)* [[AP_ASCAST]], align 8, [[TBAA5]] +// CHECK-NEXT: [[TMP80:%.*]] = bitcast i32 addrspace(4)** [[BP]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[TMP80]]) [[ATTR6]] +// CHECK-NEXT: [[TMP81:%.*]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* [[AP_ASCAST]], align 8, [[TBAA5]] +// CHECK-NEXT: [[TMP82:%.*]] = ptrtoint i32 addrspace(4)* [[TMP81]] to i64 +// CHECK-NEXT: [[TMP83:%.*]] = call i64 @llvm.annotation.i64(i64 [[TMP82]], i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([[STR1TY]], [[STR1TY]]* @.str.1, i32 0, i32 0), i32 291) +// CHECK-NEXT: [[TMP84:%.*]] = inttoptr i64 [[TMP83]] to i32 addrspace(4)* +// CHECK-NEXT: store i32 addrspace(4)* [[TMP84]], i32 addrspace(4)* addrspace(4)* [[BP_ASCAST]], align 8, [[TBAA5]] +// CHECK-NEXT: [[TMP85:%.*]] = bitcast i32 addrspace(4)** [[BP]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 8, i8* [[TMP85]]) [[ATTR6]] +// CHECK-NEXT: [[TMP86:%.*]] = bitcast i32 addrspace(4)** [[AP]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 8, i8* [[TMP86]]) [[ATTR6]] +// CHECK-NEXT: [[TMP87:%.*]] = bitcast %class._ZTS1A.A* [[CB]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP87]]) [[ATTR6]] +// CHECK-NEXT: [[TMP88:%.*]] = bitcast %class._ZTS1A.A* [[CA]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP88]]) [[ATTR6]] +// CHECK-NEXT: [[TMP89:%.*]] = bitcast [[UNION__ZTS2UN_UN]] addrspace(4)** [[U3]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 8, i8* [[TMP89]]) [[ATTR6]] +// CHECK-NEXT: [[TMP90:%.*]] = bitcast %union._ZTS2un.un* [[U2]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP90]]) [[ATTR6]] +// CHECK-NEXT: [[TMP91:%.*]] = bitcast %union._ZTS2un.un* [[U1]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP91]]) [[ATTR6]] +// CHECK-NEXT: [[TMP92:%.*]] = bitcast [[STRUCT__ZTS2ST_ST]] addrspace(4)** [[IIII]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 8, i8* [[TMP92]]) [[ATTR6]] +// CHECK-NEXT: [[TMP93:%.*]] = bitcast %struct._ZTS2st.st* [[III]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 8, i8* [[TMP93]]) [[ATTR6]] +// CHECK-NEXT: [[TMP94:%.*]] = bitcast %struct._ZTS2st.st* [[II]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 8, i8* [[TMP94]]) [[ATTR6]] +// CHECK-NEXT: [[TMP95:%.*]] = bitcast %struct._ZTS2st.st* [[I2]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 8, i8* [[TMP95]]) [[ATTR6]] +// CHECK-NEXT: [[TMP96:%.*]] = bitcast %struct._ZTS2st.st* [[I]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 8, i8* [[TMP96]]) [[ATTR6]] +// CHECK-NEXT: [[TMP97:%.*]] = bitcast i32* [[F]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP97]]) [[ATTR6]] +// CHECK-NEXT: [[TMP98:%.*]] = bitcast i32* [[E]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP98]]) [[ATTR6]] +// CHECK-NEXT: [[TMP99:%.*]] = bitcast i32* [[D]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP99]]) [[ATTR6]] +// CHECK-NEXT: [[TMP100:%.*]] = bitcast i32* [[C]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP100]]) [[ATTR6]] +// CHECK-NEXT: [[TMP101:%.*]] = bitcast i32* [[MYB]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP101]]) [[ATTR6]] +// CHECK-NEXT: [[TMP102:%.*]] = bitcast i32* [[B]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP102]]) [[ATTR6]] +// CHECK-NEXT: [[TMP103:%.*]] = bitcast i32* [[MYA]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP103]]) [[ATTR6]] +// CHECK-NEXT: [[TMP104:%.*]] = bitcast i32* [[A]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.end.p0i8(i64 4, i8* [[TMP104]]) [[ATTR6]] +// CHECK-NEXT: ret void +// 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 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 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 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 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 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 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) 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 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 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 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 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) 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 } template diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp index 2bd2b52945015..f116fff36e9b2 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp @@ -39,40 +39,49 @@ int main() { // CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8 // CHECK lambda object alloca -// CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4 +// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %"class.{{.*}}.anon", align 4 +// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %"class.{{.*}}.anon"* [[LOCAL_OBJECTA]] to %"class.{{.*}}.anon" addrspace(4)* // CHECK allocas for ranges -// CHECK: [[ACC_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MEM_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[OFFSET1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" -// CHECK: [[ACC_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" +// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE1A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[MEM_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[MEM_RANGE1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE1A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[OFFSET1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" +// CHECK: [[OFFSET1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET1A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* +// CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE2A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE2A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" +// CHECK: [[OFFSET2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET2A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* // CHECK accessor array default inits -// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0 +// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0 // Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP. -// CHECK: [[ELEM1_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[BEGIN]] to [[ACCESSOR]] addrspace(4)* // CTOR Call #1 -// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM1_ASCAST]]) -// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[BEGIN]], i64 1 -// CHECK: [[ELEM2_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[ELEM2_GEP]] to [[ACCESSOR]] addrspace(4)* +// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[BEGIN]]) +// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]] addrspace(4)* [[BEGIN]], i64 1 // CTOR Call #2 -// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_ASCAST]]) +// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_GEP]]) // CHECK acc[0] __init method call -// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0 +// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0 // CHECK load from kernel pointer argument alloca -// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]] -// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[INDEX1]] to [[ACCESSOR]] addrspace(4)* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]]) +// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]] +// CHECK: [[ACC_RANGE1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ACC_RANGE1AS]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[MEM_RANGE1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MEM_RANGE1AS]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[OFFSET1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OFFSET1AS]] to %"struct.{{.*}}.cl::sycl::id"* +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[INDEX1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]]) // CHECK acc[1] __init method call -// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY2]], i64 0, i64 1 +// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY2]], i64 0, i64 1 // CHECK load from kernel pointer argument alloca -// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG2]] -// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[INDEX2]] to [[ACCESSOR]] addrspace(4)* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]]) +// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG2]] +// CHECK: [[ACC_RANGE2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ACC_RANGE2AS]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[MEM_RANGE2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MEM_RANGE2AS]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[OFFSET2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OFFSET2AS]] to %"struct.{{.*}}.cl::sycl::id"* +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[INDEX2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]]) diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp index 672e80f0edc38..2ca0b0ec3609f 100644 --- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp @@ -29,54 +29,62 @@ int main() { // CHECK kernel_C parameters // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_C // CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]], +// CHECK-SAME: %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]], +// CHECK-SAME: %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]], +// CHECK-SAME: %"struct{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]], // CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+4]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]], -// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]]) +// CHECK-SAME: %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]], +// CHECK-SAME: %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]], +// CHECK-SAME: %"struct{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]]) // Check alloca for pointer arguments // CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8 // CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8 // Check lambda object alloca -// CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4 +// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %"class{{.*}}.anon", align 4 +// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %"class{{.*}}.anon"* %0 to %"class{{.*}}.anon" addrspace(4)* // Check allocas for ranges -// CHECK: [[ACC_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MEM_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[OFFSET1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" -// CHECK: [[ACC_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" +// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE1A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[MEM_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[MEM_RANGE1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE1A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[OFFSET1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" +// CHECK: [[OFFSET1AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET1A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* +// CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE2A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" +// CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE2A]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)* +// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" +// CHECK: [[OFFSET2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET2A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)* // CHECK accessor array default inits -// CHECK: [[ACCESSOR_WRAPPER:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_WRAPPER]], i32 0, i32 0 -// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0 -// Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP. -// CHECK: [[ELEM1_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[BEGIN]] to [[ACCESSOR]] addrspace(4)* +// CHECK: [[ACCESSOR_WRAPPER:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, %struct{{.*}}.struct_acc_t addrspace(4)* [[ACCESSOR_WRAPPER]], i32 0, i32 0 +// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0 // CTOR Call #1 -// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM1_ASCAST]]) -// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[BEGIN]], i64 1 -// CHECK: [[ELEM2_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[ELEM2_GEP]] to [[ACCESSOR]] addrspace(4)* +// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[BEGIN]]) +// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]] addrspace(4)* [[BEGIN]], i64 1 // CTOR Call #2 -// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_ASCAST]]) +// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_GEP]]) // Check acc[0] __init method call -// CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[GEP_MEMBER_ACC1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA1]], i32 0, i32 0 -// CHECK: [[ARRAY_IDX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC1]], i64 0, i64 0 -// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr -// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX1]] to [[ACCESSOR]] addrspace(4)* -// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]]) +// CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[GEP_MEMBER_ACC1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, %struct{{.*}}.struct_acc_t addrspace(4)* [[GEP_LAMBDA1]], i32 0, i32 0 +// CHECK: [[ARRAY_IDX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[GEP_MEMBER_ACC1]], i64 0, i64 0 +// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]].addr +// CHECK: [[ACC_RANGE1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ACC_RANGE1AS]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[MEM_RANGE1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MEM_RANGE1AS]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[OFFSET1:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OFFSET1AS]] to %"struct.{{.*}}.cl::sycl::id"* +// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ARRAY_IDX1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]]) // Check acc[1] __init method call -// CHECK: [[GEP_LAMBDA2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[GEP_MEMBER_ACC2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA2]], i32 0, i32 0 -// CHECK: [[ARRAY_IDX2:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC2]], i64 0, i64 1 -// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr -// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX2]] to [[ACCESSOR]] addrspace(4)* -// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]]) +// CHECK: [[GEP_LAMBDA2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[GEP_MEMBER_ACC2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, %struct{{.*}}.struct_acc_t addrspace(4)* [[GEP_LAMBDA2]], i32 0, i32 0 +// CHECK: [[ARRAY_IDX2:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[GEP_MEMBER_ACC2]], i64 0, i64 1 +// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]].addr +// CHECK: [[ACC_RANGE2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[ACC_RANGE2AS]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[MEM_RANGE2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range" addrspace(4)* [[MEM_RANGE2AS]] to %"struct.{{.*}}.cl::sycl::range"* +// CHECK: [[OFFSET2:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id" addrspace(4)* [[OFFSET2AS]] to %"struct.{{.*}}.cl::sycl::id"* +// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ARRAY_IDX2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]]) diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp index 4c7ffe0d3c591..a6c0a4f3f1cda 100644 --- a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp @@ -46,81 +46,84 @@ int main() { // Check kernel_B parameters // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_B -// CHECK-SAME:(%struct.{{.*}}.__wrapper_class* byval(%struct.{{.*}}.__wrapper_class) align 4 %[[ARR_ARG:.*]]) +// CHECK-SAME:(%struct{{.*}}.__wrapper_class* byval(%struct{{.*}}.__wrapper_class) align 4 %[[ARR_ARG:.*]]) // Check local lambda object alloca -// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = alloca %"class.{{.*}}.anon", align 4 +// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %"class{{.*}}.anon", align 4 +// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %"class{{.*}}.anon"* %[[LOCAL_OBJECTA]] to %"class{{.*}}.anon" addrspace(4)* // Check for Array init loop -// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* %[[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class, %struct.{{.*}}.__wrapper_class* %[[ARR_ARG]], i32 0, i32 0 -// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x i32], [2 x i32]* %[[LAMBDA_PTR]], i64 0, i64 0 +// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct{{.*}}.__wrapper_class, %struct{{.*}}.__wrapper_class addrspace(4)* %[[ARR_ARG]].ascast, i32 0, i32 0 +// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x i32], [2 x i32] addrspace(4)* %[[LAMBDA_PTR]], i64 0, i64 0 // CHECK: br label %[[ARRAYINITBODY:.+]] // The loop body itself // CHECK: [[ARRAYINITBODY]]: // CHECK: %[[ARRAYINDEX:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX:.+]], %[[ARRAYINITBODY]] ] -// CHECK: %[[TARG_ARRAY_ELEM:.+]] = getelementptr inbounds i32, i32* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]] -// CHECK: %[[SRC_ELEM:.+]] = getelementptr inbounds [2 x i32], [2 x i32]* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]] -// CHECK: %[[SRC_VAL:.+]] = load i32, i32* %[[SRC_ELEM]] -// CHECK: store i32 %[[SRC_VAL]], i32* %[[TARG_ARRAY_ELEM]] +// CHECK: %[[TARG_ARRAY_ELEM:.+]] = getelementptr inbounds i32, i32 addrspace(4)* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]] +// CHECK: %[[SRC_ELEM:.+]] = getelementptr inbounds [2 x i32], [2 x i32] addrspace(4)* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]] +// CHECK: %[[SRC_VAL:.+]] = load i32, i32 addrspace(4)* %[[SRC_ELEM]] +// CHECK: store i32 %[[SRC_VAL]], i32 addrspace(4)* %[[TARG_ARRAY_ELEM]] // CHECK: %[[NEXTINDEX]] = add nuw i64 %[[ARRAYINDEX]], 1 // CHECK: %[[ISDONE:.+]] = icmp eq i64 %[[NEXTINDEX]], 2 // CHECK: br i1 %[[ISDONE]], label %{{.*}}, label %[[ARRAYINITBODY]] // Check kernel_C parameters // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_C -// CHECK-SAME:(%struct.{{.*}}.__wrapper_class{{.*}}* byval(%struct.{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]]) +// CHECK-SAME:(%struct{{.*}}.__wrapper_class{{.*}}* byval(%struct{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]]) // Check local lambda object alloca -// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = alloca %"class.{{.*}}.anon{{.*}}", align 4 +// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %"class{{.*}}.anon{{.*}}", align 4 +// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %"class{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECTA]] to %"class{{.*}}.anon{{.*}}" addrspace(4)* // Check for Array init loop -// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class.{{.*}}.anon{{.*}}", %"class.{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class{{.*}}, %struct.{{.*}}.__wrapper_class{{.*}}* %[[ARR_ARG]], i32 0, i32 0 -// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* %[[LAMBDA_PTR]], i64 0, i64 0 +// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class{{.*}}.anon{{.*}}", %"class{{.*}}.anon{{.*}}" addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct{{.*}}.__wrapper_class{{.*}}, %struct{{.*}}.__wrapper_class{{.*}} addrspace(4)* %[[ARR_ARG]].ascast, i32 0, i32 0 +// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x %struct{{.*}}.foo], [2 x %struct{{.*}}.foo] addrspace(4)* %[[LAMBDA_PTR]], i64 0, i64 0 // CHECK: br label %[[ARRAYINITBODY:.+]] // The loop body itself // CHECK: [[ARRAYINITBODY]]: // CHECK: %[[ARRAYINDEX:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX:.+]], %[[ARRAYINITBODY]] ] -// CHECK: %[[TARG_ARRAY_ELEM:.+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]] -// CHECK: %[[SRC_ELEM:.+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]] -// CHECK: %[[TARG_PTR:.+]] = bitcast %struct.{{.*}}.foo* %[[TARG_ARRAY_ELEM]] to i8* -// CHECK: %[[SRC_PTR:.+]] = bitcast %struct.{{.*}}.foo* %[[SRC_ELEM]] to i8* -// call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[TARG_PTR]], i8* align %[[SRC_PTR]], i64 24, i1 false) +// CHECK: %[[TARG_ARRAY_ELEM:.+]] = getelementptr inbounds %struct{{.*}}.foo, %struct{{.*}}.foo addrspace(4)* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]] +// CHECK: %[[SRC_ELEM:.+]] = getelementptr inbounds [2 x %struct{{.*}}.foo], [2 x %struct{{.*}}.foo] addrspace(4)* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]] +// CHECK: %[[TARG_PTR:.+]] = bitcast %struct{{.*}}.foo addrspace(4)* %[[TARG_ARRAY_ELEM]] to i8 addrspace(4)* +// CHECK: %[[SRC_PTR:.+]] = bitcast %struct{{.*}}.foo addrspace(4)* %[[SRC_ELEM]] to i8 addrspace(4)* +// call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 %[[TARG_PTR]], i8 addrspace(4)* align %[[SRC_PTR]], i64 24, i1 false) // CHECK: %[[NEXTINDEX]] = add nuw i64 %[[ARRAYINDEX]], 1 // CHECK: %[[ISDONE:.+]] = icmp eq i64 %[[NEXTINDEX]], 2 // CHECK: br i1 %[[ISDONE]], label %{{.*}}, label %[[ARRAYINITBODY]] // Check kernel_D parameters // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_D -// CHECK-SAME:(%struct.{{.*}}.__wrapper_class{{.*}}* byval(%struct.{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]]) +// CHECK-SAME:(%struct{{.*}}.__wrapper_class{{.*}}* byval(%struct{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]]) // Check local lambda object alloca -// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = alloca %"class.{{.*}}.anon{{.*}}", align 4 +// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %"class{{.*}}.anon{{.*}}", align 4 +// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %"class{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECTA]] to %"class{{.*}}.anon{{.*}}" addrspace(4)* // Check for Array init loop -// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class.{{.*}}.anon{{.*}}", %"class.{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class{{.*}}, %struct.{{.*}}.__wrapper_class{{.*}}* %[[ARR_ARG]], i32 0, i32 0 -// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]]* %[[LAMBDA_PTR]], i64 0, i64 0 +// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class{{.*}}.anon{{.*}}", %"class{{.*}}.anon{{.*}}" addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct{{.*}}.__wrapper_class{{.*}}, %struct{{.*}}.__wrapper_class{{.*}} addrspace(4)* %[[ARR_ARG]].ascast, i32 0, i32 0 +// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]] addrspace(4)* %[[LAMBDA_PTR]], i64 0, i64 0 // CHECK: br label %[[ARRAYINITBODY:.+]] // Check Outer loop. // CHECK: [[ARRAYINITBODY]]: // CHECK: %[[ARRAYINDEX:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX:.+]], %[[ARRAYINITEND:.+]] ] -// CHECK: %[[TARG_OUTER_ELEM:.+]] = getelementptr inbounds [1 x i32], [1 x i32]* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]] -// CHECK: %[[SRC_OUTER_ELEM:.+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]]* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]] -// CHECK: %[[ARRAY_BEGIN_INNER:.+]] = getelementptr inbounds [1 x i32], [1 x i32]* %[[TARG_OUTER_ELEM]], i64 0, i64 0 +// CHECK: %[[TARG_OUTER_ELEM:.+]] = getelementptr inbounds [1 x i32], [1 x i32] addrspace(4)* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]] +// CHECK: %[[SRC_OUTER_ELEM:.+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]] addrspace(4)* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]] +// CHECK: %[[ARRAY_BEGIN_INNER:.+]] = getelementptr inbounds [1 x i32], [1 x i32] addrspace(4)* %[[TARG_OUTER_ELEM]], i64 0, i64 0 // CHECK: br label %[[ARRAYINITBODY_INNER:.+]] // Check Inner Loop // CHECK: [[ARRAYINITBODY_INNER]]: // CHECK: %[[ARRAYINDEX_INNER:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX_INNER:.+]], %[[ARRAYINITBODY_INNER:.+]] ] -// CHECK: %[[TARG_INNER_ELEM:.+]] = getelementptr inbounds i32, i32* %[[ARRAY_BEGIN_INNER]], i64 %[[ARRAYINDEX_INNER]] -// CHECK: %[[SRC_INNER_ELEM:.+]] = getelementptr inbounds [1 x i32], [1 x i32]* %[[SRC_OUTER_ELEM]], i64 0, i64 %[[ARRAYINDEX_INNER]] -// CHECK: %[[SRC_LOAD:.+]] = load i32, i32* %[[SRC_INNER_ELEM]] -// CHECK: store i32 %[[SRC_LOAD]], i32* %[[TARG_INNER_ELEM]] +// CHECK: %[[TARG_INNER_ELEM:.+]] = getelementptr inbounds i32, i32 addrspace(4)* %[[ARRAY_BEGIN_INNER]], i64 %[[ARRAYINDEX_INNER]] +// CHECK: %[[SRC_INNER_ELEM:.+]] = getelementptr inbounds [1 x i32], [1 x i32] addrspace(4)* %[[SRC_OUTER_ELEM]], i64 0, i64 %[[ARRAYINDEX_INNER]] +// CHECK: %[[SRC_LOAD:.+]] = load i32, i32 addrspace(4)* %[[SRC_INNER_ELEM]] +// CHECK: store i32 %[[SRC_LOAD]], i32 addrspace(4)* %[[TARG_INNER_ELEM]] // CHECK: %[[NEXTINDEX_INNER]] = add nuw i64 %[[ARRAYINDEX_INNER]], 1 // CHECK: %[[ISDONE_INNER:.+]] = icmp eq i64 %[[NEXTINDEX_INNER]], 1 // CHECK: br i1 %[[ISDONE_INNER]], label %[[ARRAYINITEND]], label %[[ARRAYINITBODY_INNER]] diff --git a/clang/test/CodeGenSYCL/local_var_big_init_as.cpp b/clang/test/CodeGenSYCL/local_var_big_init_as.cpp deleted file mode 100644 index 81461549d6a26..0000000000000 --- a/clang/test/CodeGenSYCL/local_var_big_init_as.cpp +++ /dev/null @@ -1,30 +0,0 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice \ -// RUN: -emit-llvm -o - %s | FileCheck %s - -// This test checks that data for big constant initializer lists is placed -// into the global address space by the SYCL compiler. - -struct Test { - Test() : set{ - 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, - 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15} {}; - int set[32]; -}; -// CHECK-DAG: @constinit = private unnamed_addr addrspace(1) constant -// CHECK: [32 x i32] -// CHECK: [i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, -// CHECK: i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15, -// CHECK: i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, -// CHECK: i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15 -// CHECK: ], align 4 -// CHECK-NOT: @constinit = private unnamed_addr addrspace(0) -// CHECK-NOT: @constinit = private unnamed_addr addrspace(2) -// CHECK-NOT: @constinit = private unnamed_addr addrspace(3) -// CHECK-NOT: @constinit = private unnamed_addr addrspace(4) - -__attribute__((sycl_device)) void bar(Test &x); - -__attribute__((sycl_device)) void zoo() { - Test mc; - bar(mc); -} diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp index a5aa3ed837c4e..dfe58a127592e 100644 --- a/clang/test/CodeGenSYCL/sampler.cpp +++ b/clang/test/CodeGenSYCL/sampler.cpp @@ -2,14 +2,14 @@ // CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%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 -// CHECK-NEXT: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8 +// CHECK: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 8 +// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon"* [[ANON]] to %"class.{{.*}}.anon" addrspace(4)* +// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8 // CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %"class.{{.*}}.anon"* [[ANON]] to i8* // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4 -// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0 -// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8 -// CHECK-NEXT: [[GEPCAST:%[0-9]+]] = addrspacecast %"class{{.*}}.cl::sycl::sampler"* [[GEP]] to %"class{{.*}}.cl::sycl::sampler" addrspace(4)* -// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{[^,]*}} [[GEPCAST]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) +// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[ANONCAST]], i32 0, i32 0 +// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8 +// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{[^,]*}} [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) // // CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED:%[a-zA-Z0-9_]+]], i32 [[ARG_A:%[a-zA-Z0-9_]+]]) @@ -17,22 +17,23 @@ // Check alloca // CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8 // CHECK: [[ARG_A]].addr = alloca i32, align 4 -// CHECK: [[LAMBDA:%[0-9]+]] = alloca %"class.{{.*}}.anon.0", align 8 +// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %"class.{{.*}}.anon.0", align 8 +// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon.0"* [[LAMBDAA]] to %"class.{{.*}}.anon.0" addrspace(4)* // Check argument store -// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG_WRAPPED]].addr, align 8 -// CHECK: store i32 [[ARG_A]], i32* [[ARG_A]].addr, align 4 +// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8 +// CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4 // Initialize 'a' -// CHECK: [[GEP_LAMBDA:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0"* [[LAMBDA]], i32 0, i32 0 -// CHECK: [[GEP_A:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper* [[GEP_LAMBDA]], i32 0, i32 1 -// CHECK: [[LOAD_A:%[0-9]+]] = load i32, i32* [[ARG_A]].addr, align 4 -// CHECK: store i32 [[LOAD_A]], i32* [[GEP_A]], align 8 +// CHECK: [[GEP_LAMBDA:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0" addrspace(4)* [[LAMBDA]], i32 0, i32 0 +// CHECK: [[GEP_A:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper addrspace(4)* [[GEP_LAMBDA]], i32 0, i32 1 +// CHECK: [[LOAD_A:%[0-9]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4 +// CHECK: store i32 [[LOAD_A]], i32 addrspace(4)* [[GEP_A]], align 8 // Initialize wrapped sampler 'smpl' -// CHECK: [[GEP_LAMBDA_0:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0"* %0, i32 0, i32 0 -// CHECK: [[GEP_SMPL:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper* [[GEP_LAMBDA_0]], i32 0, i32 0 -// CHECK: [[LOAD_SMPL:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG_WRAPPED]].addr, align 8 +// CHECK: [[GEP_LAMBDA_0:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0" addrspace(4)* [[LAMBDA]], i32 0, i32 0 +// CHECK: [[GEP_SMPL:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper addrspace(4)* [[GEP_LAMBDA_0]], i32 0, i32 0 +// CHECK: [[LOAD_SMPL:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8 // CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SMPL]]) // #include "Inputs/sycl.hpp" diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp index 2544b4e22f007..7efdf06c09fe6 100644 --- a/clang/test/CodeGenSYCL/spir-calling-conv.cpp +++ b/clang/test/CodeGenSYCL/spir-calling-conv.cpp @@ -9,7 +9,7 @@ int main() { // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function() - // CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %2) + // CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %{{.+}}) // 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 121270b57642f..ba519e539e7d9 100644 --- a/clang/test/CodeGenSYCL/spir-enum.cpp +++ b/clang/test/CodeGenSYCL/spir-enum.cpp @@ -22,9 +22,8 @@ int main() { // CHECK: define {{.*}}spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 %_arg_) - // CHECK: getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* - // CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %4) - + // CHECK: getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* + // CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %{{.+}}) test( enum_type::B ); return 0; diff --git a/clang/test/CodeGenSYCL/union-kernel-param.cpp b/clang/test/CodeGenSYCL/union-kernel-param.cpp index 45e950b8e48d2..ee22e92a80347 100644 --- a/clang/test/CodeGenSYCL/union-kernel-param.cpp +++ b/clang/test/CodeGenSYCL/union-kernel-param.cpp @@ -33,9 +33,10 @@ int main() { // Check lambda object alloca // CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4 -// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion* [[L_STRUCT_ADDR]] to i8* -// CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion* [[MEM_ARG]] to i8* -// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[MEMCPY_DST]], i8* align 4 [[MEMCPY_SRC]], i64 12, i1 false) -// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon"* [[LOCAL_OBJECT]] to %"class.{{.*}}.anon" addrspace(4)* -// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} [[ACC_CAST1]]) +// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast %"class.{{.*}}.anon"* [[LOCAL_OBJECT]] to %"class.{{.*}}.anon" addrspace(4)* +// CHECK: [[MEM_ARGAS:%.*]] = addrspacecast %union.{{.*}}.MyUnion* [[MEM_ARG]] to %union.{{.*}}.MyUnion addrspace(4)* +// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECTAS]], i32 0, i32 0 +// CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion addrspace(4)* [[L_STRUCT_ADDR]] to i8 addrspace(4)* +// CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion addrspace(4)* [[MEM_ARGAS]] to i8 addrspace(4)* +// CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[MEMCPY_DST]], i8 addrspace(4)* align 4 [[MEMCPY_SRC]], i64 12, i1 false) +// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} [[LOCAL_OBJECTAS]]) diff --git a/clang/test/CodeGenSYCL/unique-stable-name.cpp b/clang/test/CodeGenSYCL/unique-stable-name.cpp index 64b06f0a2fd7e..f2c38737b95e5 100644 --- a/clang/test/CodeGenSYCL/unique-stable-name.cpp +++ b/clang/test/CodeGenSYCL/unique-stable-name.cpp @@ -41,36 +41,36 @@ int main() { kernel_single_task( []() { printf(__builtin_unique_stable_name(int)); - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(4)* addrspacecast ([[INT_SIZE]] addrspace(1)* @[[INT]] to [[INT_SIZE]] addrspace(4)* auto x = [](){}; printf(__builtin_unique_stable_name(x)); printf(__builtin_unique_stable_name(decltype(x))); - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(4)* addrspacecast ([[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] to [[LAMBDA_X_SIZE]] addrspace(4)* + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(4)* addrspacecast ([[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] to [[LAMBDA_X_SIZE]] addrspace(4)* DEF_IN_MACRO(); - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_X]] - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(1)* @[[MACRO_Y]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(4)* addrspacecast ([[MACRO_SIZE]] addrspace(1)* @[[MACRO_X]] to [[MACRO_SIZE]] addrspace(4)* + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(4)* addrspacecast ([[MACRO_SIZE]] addrspace(1)* @[[MACRO_Y]] to [[MACRO_SIZE]] addrspace(4)* MACRO_CALLS_MACRO(); - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_X]] - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_Y]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(4)* addrspacecast ([[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_X]] to [[MACRO_MACRO_SIZE]] addrspace(4)* + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(4)* addrspacecast ([[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_Y]] to [[MACRO_MACRO_SIZE]] addrspace(4)* template_param(); // CHECK: define linkonce_odr spir_func void @_Z14template_paramIiEvv - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(1)* @[[INT]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(4)* addrspacecast ([[INT_SIZE]] addrspace(1)* @[[INT]] to [[INT_SIZE]] addrspace(4)* template_param(); // CHECK: define internal spir_func void @"_Z14template_paramIZZ4mainENK3 - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(4)* addrspacecast ([[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] to [[LAMBDA_X_SIZE]] addrspace(4)* lambda_in_dependent_function(); // CHECK: define linkonce_odr spir_func void @_Z28lambda_in_dependent_functionIiEvv - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_INT]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]] addrspace(4)* addrspacecast ([[DEP_INT_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_INT]] to [[DEP_INT_SIZE]] addrspace(4)* lambda_in_dependent_function(); // CHECK: define internal spir_func void @"_Z28lambda_in_dependent_functionIZZ4mainENK3$_0clEvEUlvE_Evv - // CHECK: call spir_func void @printf(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_X]] + // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]] addrspace(4)* addrspacecast ([[DEP_LAMBDA_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_X]] to [[DEP_LAMBDA_SIZE]] addrspace(4)* }); }