diff --git a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp index 5887232ff83c..238dbe7131c0 100644 --- a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp @@ -323,6 +323,11 @@ bool ThreadSanitizerOnSpirv::instrumentAllocInst( InstrumentationIRBuilder::ensureDebugInfo(*AtExit, *F); for (auto *Inst : AllocaInsts) { AllocaInst *AI = cast(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, diff --git a/llvm/test/Instrumentation/ThreadSanitizer/SPIRV/cleanup_private_shadow.ll b/llvm/test/Instrumentation/ThreadSanitizer/SPIRV/cleanup_private_shadow.ll index 08a1ac30e909..28ac1f0c980d 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 }