From a2178aa0cbbe675ce6508284720559940c6fd248 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he@intel.com>
Date: Sat, 8 Aug 2026 04:50:30 +0200
Subject: [PATCH] [DeviceSanitizer] Fix circular library dependency causing
 undefined symbol

Commit 0941896bca24 ("[DeviceSanitizer] Make sanitizer metadata globals
per-module unique (#22567)") made SanitizerPostOptimizer.cpp in
LLVMSYCLLowerIR call sycl::isModuleUsingMsan(), which is defined in
LLVMSYCLPostLink. LLVMSYCLPostLink already links against LLVMSYCLLowerIR,
so this created a circular library dependency. It doesn't show up when
linking against full static LLVM libs, but breaks with `ld.lld: error:
undefined symbol:llvm::sycl::isModuleUsingMsan(llvm::Module const&)`
when building LLVMSYCLLowerIR as a standalone shared library (-slibs build).

Fix by moving isModuleUsingAsan/Msan/Tsan down into
SYCLLowerIR/SYCLUtils.{h,cpp}, which has no dependency on SYCLPostLink.

Co-authored-by: Claude Sonnet 5 <noreply@anthropic.com>
---
 llvm/include/llvm/SYCLLowerIR/SYCLUtils.h     |  8 ++++++++
 .../SYCLPostLink/ComputeModuleRuntimeInfo.h   |  3 ---
 llvm/lib/SYCLLowerIR/SYCLUtils.cpp            | 19 +++++++++++++++++++
 .../SYCLLowerIR/SanitizerPostOptimizer.cpp    |  2 +-
 .../SYCLPostLink/ComputeModuleRuntimeInfo.cpp | 18 ------------------
 .../lib/rtc/DeviceCompilation.cpp             |  1 +
 6 files changed, 29 insertions(+), 22 deletions(-)

diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h
index c9ebcdae53f4b..6238b494e9124 100644
--- a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h
+++ b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h
@@ -19,7 +19,15 @@
 #include <functional>
 
 namespace llvm {
+
+class Module;
+
 namespace sycl {
+
+bool isModuleUsingAsan(const Module &M);
+bool isModuleUsingMsan(const Module &M);
+bool isModuleUsingTsan(const Module &M);
+
 namespace utils {
 constexpr char ATTR_SYCL_MODULE_ID[] = "sycl-module-id";
 constexpr char ATTR_SYCL_OPTLEVEL[] = "sycl-optlevel";
diff --git a/llvm/include/llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h b/llvm/include/llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h
index 1821b605b36eb..bb3426d0c36ab 100644
--- a/llvm/include/llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h
+++ b/llvm/include/llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h
@@ -28,9 +28,6 @@ struct GlobalBinImageProps {
   bool EmitImportedSymbols;
   bool EmitDeviceGlobalPropSet;
 };
-bool isModuleUsingAsan(const Module &M);
-bool isModuleUsingMsan(const Module &M);
-bool isModuleUsingTsan(const Module &M);
 using PropSetRegTy = llvm::util::PropertySetRegistry;
 using EntryPointSet = SetVector<Function *>;
 
diff --git a/llvm/lib/SYCLLowerIR/SYCLUtils.cpp b/llvm/lib/SYCLLowerIR/SYCLUtils.cpp
index 84354acf7d416..86ecf2cb61ee0 100644
--- a/llvm/lib/SYCLLowerIR/SYCLUtils.cpp
+++ b/llvm/lib/SYCLLowerIR/SYCLUtils.cpp
@@ -14,6 +14,25 @@
 
 namespace llvm {
 namespace sycl {
+
+bool isModuleUsingAsan(const Module &M) {
+  return any_of(M.globals(), [](const GlobalVariable &GV) {
+    return GV.getName().starts_with("__AsanKernelMetadata");
+  });
+}
+
+bool isModuleUsingMsan(const Module &M) {
+  return any_of(M.globals(), [](const GlobalVariable &GV) {
+    return GV.getName().starts_with("__MsanKernelMetadata");
+  });
+}
+
+bool isModuleUsingTsan(const Module &M) {
+  return any_of(M.globals(), [](const GlobalVariable &GV) {
+    return GV.getName().starts_with("__TsanKernelMetadata");
+  });
+}
+
 namespace utils {
 
 using namespace llvm::esimd;
diff --git a/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp b/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp
index 9759d99b00476..c0318044d89b6 100644
--- a/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp
+++ b/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp
@@ -13,7 +13,7 @@
 //===----------------------------------------------------------------------===//
 
 #include "llvm/SYCLLowerIR/SanitizerPostOptimizer.h"
-#include "llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h"
+#include "llvm/SYCLLowerIR/SYCLUtils.h"
 
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/InstVisitor.h"
diff --git a/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp b/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp
index 10ccf74cc1c7b..0edee0156eda5 100644
--- a/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp
+++ b/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp
@@ -43,24 +43,6 @@ getSYCLESIMDSplitStatusFromMetadata(const Module &M) {
 }
 } // namespace
 
-bool isModuleUsingAsan(const Module &M) {
-  return any_of(M.globals(), [](const GlobalVariable &GV) {
-    return GV.getName().starts_with("__AsanKernelMetadata");
-  });
-}
-
-bool isModuleUsingMsan(const Module &M) {
-  return any_of(M.globals(), [](const GlobalVariable &GV) {
-    return GV.getName().starts_with("__MsanKernelMetadata");
-  });
-}
-
-bool isModuleUsingTsan(const Module &M) {
-  return any_of(M.globals(), [](const GlobalVariable &GV) {
-    return GV.getName().starts_with("__TsanKernelMetadata");
-  });
-}
-
 // Gets 1- to 3-dimension work-group related information for function Func.
 // Returns an empty vector if not present.
 template <typename T>
diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
index 5c44068413c5a..6cfb99e9f3319 100644
--- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
+++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
@@ -44,6 +44,7 @@
 #include <llvm/SYCLLowerIR/LowerInvokeSimd.h>
 #include <llvm/SYCLLowerIR/SYCLDeviceLibBF16.h>
 #include <llvm/SYCLLowerIR/SYCLJointMatrixTransform.h>
+#include <llvm/SYCLLowerIR/SYCLUtils.h>
 #include <llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h>
 #include <llvm/SYCLPostLink/ModuleSplitter.h>
 #include <llvm/Support/BLAKE3.h>