[sycl-post-link] Update spec constant pattern match for Device ASan (#20680)

Address Sanitizer pass changed the way to get stack offset from
PtrIntAdd to GEP instruction. We need to update the pattern match rule
correspondingly.

Author: zhaomaosu

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<AllocaInst>(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<Function>(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);

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
 }