diff --git a/llvm/lib/SYCLLowerIR/SpecConstants.cpp b/llvm/lib/SYCLLowerIR/SpecConstants.cpp index 54ef694872c08..177ce9f2d4e5b 100644 --- a/llvm/lib/SYCLLowerIR/SpecConstants.cpp +++ b/llvm/lib/SYCLLowerIR/SpecConstants.cpp @@ -122,7 +122,7 @@ StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo, V = ASC->getPointerOperand()->stripPointerCasts(); using namespace PatternMatch; Value *X; - if (match(V, m_IntToPtr(m_Add(m_PtrToInt(m_Value(X)), m_ConstantInt())))) + if (match(V, m_PtrAdd(m_Value(X), m_Constant()))) V = X; return isa(V); }; @@ -541,7 +541,8 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName, // types? Is it necessary? FunctionCallee FC = M->getOrInsertFunction(FunctionName, FT); - auto *Call = CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore); + auto *Call = CallInst::Create(FT, FC.getCallee(), Args, "", + InsertBefore->getIterator()); if (IsSPIROrSPIRV) { cast(FC.getCallee())->setCallingConv(CallingConv::SPIR_FUNC); Call->setCallingConv(CallingConv::SPIR_FUNC); @@ -724,9 +725,8 @@ Value *createLoadFromBuffer(CallInst *InsertBefore, Value *Buffer, if (SCType->isIntegerTy(1)) // No bitcast to i1 before load BitCast = GEP; else - BitCast = - new BitCastInst(GEP, PointerType::get(SCType, GEP->getAddressSpace()), - "bc", InsertBefore->getIterator()); + BitCast = new BitCastInst(GEP, PointerType::get(C, GEP->getAddressSpace()), + "bc", InsertBefore->getIterator()); // When we encounter i1 spec constant, we still load the whole byte Value *Load = new LoadInst(SCType->isIntegerTy(1) ? Int8Ty : SCType, BitCast, @@ -831,8 +831,8 @@ void updatePaddingInLastMDNode(LLVMContext &Ctx, /// type. void createStoreInstructionIntoSpecConstValue(Value *Dst, Value *V, CallInst *InsertBefore) { - Type *PointerType = - PointerType::get(V->getType(), Dst->getType()->getPointerAddressSpace()); + Type *PointerType = PointerType::get( + V->getContext(), Dst->getType()->getPointerAddressSpace()); IRBuilder B(InsertBefore); Value *Bitcast = B.CreateBitCast(Dst, PointerType); B.CreateStore(V, Bitcast); diff --git a/llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll b/llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll index 1b904abfa0f3a..df2651efa145e 100644 --- a/llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll +++ b/llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll @@ -9,23 +9,21 @@ target triple = "spir64-unknown-unknown" %"class.sycl::_V1::specialization_id" = type { i32 } @_ZL9test_id_1 = addrspace(1) constant %"class.sycl::_V1::specialization_id" { i32 42 } -@__usid_str = constant [36 x i8] c"uide7faddc6b4d2fe92____ZL9test_id_1\00" +@__usid_str = private unnamed_addr addrspace(4) constant [36 x i8] c"uide7faddc6b4d2fe92____ZL9test_id_1\00" define spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlNS0_14kernel_handlerEE_clES4_(ptr addrspace(4) %this1.i7) { entry: - %MyAlloca = alloca i8, i64 224, align 32 - %0 = ptrtoint ptr %MyAlloca to i64 - %1 = add i64 %0, 96 - %2 = inttoptr i64 %1 to ptr - %SymbolicID.ascast.i = addrspacecast ptr %2 to ptr addrspace(4) - store ptr addrspace(4) addrspacecast (ptr @__usid_str to ptr addrspace(4)), ptr addrspace(4) %SymbolicID.ascast.i, align 8 - %3 = load ptr addrspace(4), ptr addrspace(4) %SymbolicID.ascast.i, align 8 - %4 = load ptr addrspace(4), ptr addrspace(4) %this1.i7, align 8 + %MyAlloca = alloca [256 x i8], align 32 + %0 = getelementptr i8, ptr %MyAlloca, i64 96 + %SymbolicID.ascast.i = addrspacecast ptr %0 to ptr addrspace(4) + store ptr addrspace(4) @__usid_str, ptr addrspace(4) %SymbolicID.ascast.i, align 8 + %1 = load ptr addrspace(4), ptr addrspace(4) %SymbolicID.ascast.i, align 8 + %2 = load ptr addrspace(4), ptr addrspace(4) %this1.i7, align 8 ; CHECK-NOT: call spir_func noundef i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_( ; CHECK: %conv = sitofp i32 %load to double - %call.i8 = call spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(ptr addrspace(4) %3, ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL9test_id_1 to ptr addrspace(4)), ptr addrspace(4) %4) + %call.i8 = call spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(ptr addrspace(4) %1, ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL9test_id_1 to ptr addrspace(4)), ptr addrspace(4) %2) %conv = sitofp i32 %call.i8 to double ret void }