[DevTSAN] Don't insert cleanup instruction for dynamic allocas (#18090)

For dynamic allocas, sometime it will not dominate exit BB, we need to
skip them.

Author: zhaomaosu

llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp

@@ -323,6 +323,11 @@ bool ThreadSanitizerOnSpirv::instrumentAllocInst(
     InstrumentationIRBuilder::ensureDebugInfo(*AtExit, *F);
     for (auto *Inst : AllocaInsts) {
       AllocaInst *AI = cast<AllocaInst>(Inst);
+      // For dynamic allocas, sometime it will not dominate exit BB, we need to
+      // skip them.
+      if (!AI->isStaticAlloca())
+        continue;
+
       if (auto AllocSize = AI->getAllocationSize(DL)) {
         AtExit->CreateCall(
             TsanCleanupPrivate,

llvm/test/Instrumentation/ThreadSanitizer/SPIRV/cleanup_private_shadow.ll

@@ -5,10 +5,24 @@ target triple = "spir64-unknown-unknown"
 %"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" }
 %"class.sycl::_V1::detail::array" = type { [1 x i64] }
 
-define spir_kernel void @test() {
+%struct.Dimensions = type { i32, i32, i32, i32, i32, i32 }
+
+define spir_kernel void @test(i32 %val) {
 entry:
   %agg.tmp = alloca %"class.sycl::_V1::range", align 8
+  %cmp = icmp eq i32 %val, 1
+  br i1 %cmp, label %for.body.preheader, label %exit
+
+for.body.preheader:                     ; preds = %entry
+  br label %for.body
+
+for.body:                               ; preds = %for.body.preheader
+  %device-byval-temp.ascast234298 = alloca %struct.Dimensions, i32 0, align 8, addrspace(4)
+  br label %exit
+
+exit:
 ; CHECK: [[REG1:%[0-9]+]] = ptrtoint ptr %agg.tmp to i64
 ; CHECK-NEXT: call void @__tsan_cleanup_private(i64 [[REG1]], i32 8)
+; CHECK-NOT: ptrtoint ptr %device-byval-temp.ascast234298 to i64
   ret void
 }