From 2c7ef80ff207e0e529a8c86dbdbf5f003408ce84 Mon Sep 17 00:00:00 2001
From: "Zhao, Maosu" <maosu.zhao@intel.com>
Date: Fri, 18 Apr 2025 07:47:26 +0200
Subject: [PATCH 1/2] [DevTSAN] Don't insert cleanup instruction for dynamic
 allocas

---
 .../Instrumentation/ThreadSanitizer.cpp          |  3 +++
 .../SPIRV/cleanup_private_shadow.ll              | 16 +++++++++++++++-
 2 files changed, 18 insertions(+), 1 deletion(-)

diff --git a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
index 5887232ff83c9..57aed92f86b13 100644
--- a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
+++ b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
@@ -323,6 +323,9 @@ bool ThreadSanitizerOnSpirv::instrumentAllocInst(
     InstrumentationIRBuilder::ensureDebugInfo(*AtExit, *F);
     for (auto *Inst : AllocaInsts) {
       AllocaInst *AI = cast<AllocaInst>(Inst);
+      if (!AI->isStaticAlloca())
+        continue;
+
       if (auto AllocSize = AI->getAllocationSize(DL)) {
         AtExit->CreateCall(
             TsanCleanupPrivate,
diff --git a/llvm/test/Instrumentation/ThreadSanitizer/SPIRV/cleanup_private_shadow.ll b/llvm/test/Instrumentation/ThreadSanitizer/SPIRV/cleanup_private_shadow.ll
index 08a1ac30e9092..28ac1f0c980d4 100644
--- a/llvm/test/Instrumentation/ThreadSanitizer/SPIRV/cleanup_private_shadow.ll
+++ b/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
 }

From cdc89dfda78cae2f7286469b7ccfeb43e508653e Mon Sep 17 00:00:00 2001
From: "Zhao, Maosu" <maosu.zhao@intel.com>
Date: Fri, 18 Apr 2025 07:56:54 +0200
Subject: [PATCH 2/2] add comments

---
 llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
index 57aed92f86b13..238dbe7131c0e 100644
--- a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
+++ b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
@@ -323,6 +323,8 @@ 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;