[sycl-post-link] Fix spec constant pattern match for DeviceSanitizer (#14740)

Adjust spec constant pattern match for base alloca + offset case in
device sanitizer.
Address sanitizer merges static allocas into a large layout base alloca
and original alloca is replaced with base + offset.

Author: wenju-he

llvm/include/llvm/SYCLLowerIR/SpecConstants.h

@@ -59,7 +59,7 @@ class SpecConstantsPass : public PassInfoMixin<SpecConstantsPass> {
   enum class HandlingMode { default_values, emulation, native };
 
 public:
-  SpecConstantsPass(HandlingMode Mode) : Mode(Mode) {}
+  SpecConstantsPass(HandlingMode Mode = HandlingMode::emulation) : Mode(Mode) {}
   PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM);
 
   // Searches given module for occurrences of specialization constant-specific
@@ -73,7 +73,7 @@ class SpecConstantsPass : public PassInfoMixin<SpecConstantsPass> {
                                            std::vector<char> &DefaultValues);
 
 private:
-  HandlingMode Mode = HandlingMode::emulation;
+  HandlingMode Mode;
 };
 
 bool checkModuleContainsSpecConsts(const Module &M);

llvm/lib/Passes/PassBuilder.cpp

@@ -131,6 +131,7 @@
 #include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h"
 #include "llvm/SYCLLowerIR/SYCLPropagateJointMatrixUsage.h"
 #include "llvm/SYCLLowerIR/SYCLVirtualFunctionsAnalysis.h"
+#include "llvm/SYCLLowerIR/SpecConstants.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/ErrorHandling.h"

llvm/lib/Passes/PassRegistry.def

@@ -164,6 +164,7 @@ MODULE_PASS("lower-slm-reservation-calls", ESIMDLowerSLMReservationCalls())
 MODULE_PASS("record-sycl-aspect-names", RecordSYCLAspectNamesPass())
 MODULE_PASS("sycl-virtual-functions-analysis",
             SYCLVirtualFunctionsAnalysisPass())
+MODULE_PASS("spec-constants", SpecConstantsPass())
 #undef MODULE_PASS
 
 #ifndef MODULE_PASS_WITH_PARAMS

llvm/lib/SYCLLowerIR/SpecConstants.cpp

@@ -20,6 +20,7 @@
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/IR/Operator.h"
+#include "llvm/IR/PatternMatch.h"
 #include "llvm/TargetParser/Triple.h"
 
 #include <vector>
@@ -101,12 +102,16 @@ StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo,
     // so that %1 is trivially known to be the address of the @.str literal.
 
     Value *TmpPtr = L->getPointerOperand();
-    AssertRelease((isa<AddrSpaceCastInst>(TmpPtr) &&
-                   isa<AllocaInst>(cast<AddrSpaceCastInst>(TmpPtr)
-                                       ->getPointerOperand()
-                                       ->stripPointerCasts())) ||
-                      isa<AllocaInst>(TmpPtr),
-                  "unexpected instruction type");
+    auto ValueIsAlloca = [](Value *V) {
+      if (auto *ASC = dyn_cast<AddrSpaceCastInst>(V))
+        V = ASC->getPointerOperand()->stripPointerCasts();
+      using namespace PatternMatch;
+      Value *X;
+      if (match(V, m_IntToPtr(m_Add(m_PtrToInt(m_Value(X)), m_ConstantInt()))))
+        V = X;
+      return isa<AllocaInst>(V);
+    };
+    AssertRelease(ValueIsAlloca(TmpPtr), "unexpected instruction type");
 
     // find the store of the literal address into TmpPtr
     StoreInst *Store = nullptr;

llvm/test/SYCLLowerIR/SpecConstants/literal-address-alloca-asan.ll

@@ -0,0 +1,33 @@
+; RUN: opt -passes=spec-constants %s -S -o - | FileCheck %s
+
+; Check there is no assert error when literal address is loaded from an alloca
+; with offset.
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
+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"
+
+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
+
+; 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)
+  %conv = sitofp i32 %call.i8 to double
+  ret void
+}
+
+declare spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4))