diff --git a/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h b/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h
index 8ad837beb7da5..144f99b753e7e 100644
--- a/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h
+++ b/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h
@@ -18,6 +18,7 @@
 #include "llvm/IR/Function.h"
 #include "llvm/IR/Module.h"
 #include "llvm/SYCLLowerIR/SYCLDeviceRequirements.h"
+#include "llvm/Support/Compiler.h"
 #include "llvm/Support/Error.h"
 #include "llvm/Support/PropertySetIO.h"
 
@@ -238,8 +239,11 @@ class ModuleDesc {
 
 #ifndef NDEBUG
   void verifyESIMDProperty() const;
-  void dump() const;
 #endif // NDEBUG
+
+#if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP)
+  LLVM_DUMP_METHOD void dump() const;
+#endif
 };
 
 // Module split support interface.
@@ -303,11 +307,13 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
                       bool EmitOnlyKernelsAsEntryPoints,
                       bool AllowDeviceImageDependencies);
 
-#ifndef NDEBUG
-void dumpEntryPoints(const EntryPointSet &C, const char *Msg = "", int Tab = 0);
-void dumpEntryPoints(const Module &M, bool OnlyKernelsAreEntryPoints = false,
-                     const char *Msg = "", int Tab = 0);
-#endif // NDEBUG
+#if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP)
+LLVM_DUMP_METHOD void dumpEntryPoints(const EntryPointSet &C,
+                                      const char *Msg = "", int Tab = 0);
+LLVM_DUMP_METHOD void dumpEntryPoints(const Module &M,
+                                      bool OnlyKernelsAreEntryPoints = false,
+                                      const char *Msg = "", int Tab = 0);
+#endif
 
 struct SplitModule {
   std::string ModuleFilePath;
diff --git a/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp b/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp
index 10b1c66bbe044..f0df7fa2b3a59 100644
--- a/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp
+++ b/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp
@@ -27,17 +27,6 @@
 #include <string>
 #include <vector>
 
-#ifdef NDEBUG
-#define DUMP_ENTRY_POINTS(...)
-#else
-constexpr int DebugESIMDPostSplit = 0;
-
-#define DUMP_ENTRY_POINTS(...)                                                 \
-  if (DebugESIMDPostSplit > 0) {                                               \
-    llvm::module_split::dumpEntryPoints(__VA_ARGS__);                          \
-  }
-#endif // NDEBUG
-
 using namespace llvm;
 using namespace llvm::module_split;
 
@@ -134,11 +123,9 @@ llvm::sycl::handleESIMD(ModuleDesc MDesc,
 
   SplitOccurred |= Result.size() > 1;
 
-  for (ModuleDesc &MD : Result) {
-    DUMP_ENTRY_POINTS(MD.entries(), MD.Name.c_str(), 4);
+  for (ModuleDesc &MD : Result)
     if (Options.LowerESIMD && MD.isESIMD())
       Modified |= lowerESIMDConstructs(MD, Options);
-  }
 
   if (Options.SplitESIMD || Result.size() == 1)
     return std::move(Result);
@@ -164,7 +151,6 @@ llvm::sycl::handleESIMD(ModuleDesc MDesc,
   Linked.rebuildEntryPoints(Names);
   Result.clear();
   Result.emplace_back(std::move(Linked));
-  DUMP_ENTRY_POINTS(Result.back().entries(), Result.back().Name.c_str(), 4);
   Modified = true;
 
   return std::move(Result);
diff --git a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp b/llvm/lib/SYCLPostLink/ModuleSplitter.cpp
index 24112fab71c16..30647bd826040 100644
--- a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp
+++ b/llvm/lib/SYCLPostLink/ModuleSplitter.cpp
@@ -561,55 +561,6 @@ Error ModuleSplitterBase::verifyNoCrossModuleDeviceGlobalUsage() {
   return Error::success();
 }
 
-#ifndef NDEBUG
-
-const char *toString(SyclEsimdSplitStatus S) {
-  switch (S) {
-  case SyclEsimdSplitStatus::ESIMD_ONLY:
-    return "ESIMD_ONLY";
-  case SyclEsimdSplitStatus::SYCL_ONLY:
-    return "SYCL_ONLY";
-  case SyclEsimdSplitStatus::SYCL_AND_ESIMD:
-    return "SYCL_AND_ESIMD";
-  }
-  return "<UNKNOWN_STATUS>";
-}
-
-void tab(int N) {
-  for (int I = 0; I < N; ++I) {
-    llvm::errs() << "  ";
-  }
-}
-
-void dumpEntryPoints(const EntryPointSet &C, const char *msg, int Tab) {
-  tab(Tab);
-  llvm::errs() << "ENTRY POINTS"
-               << " " << msg << " {\n";
-  for (const Function *F : C) {
-    tab(Tab);
-    llvm::errs() << "  " << F->getName() << "\n";
-  }
-  tab(Tab);
-  llvm::errs() << "}\n";
-}
-
-void dumpEntryPoints(const Module &M, bool OnlyKernelsAreEntryPoints,
-                     const char *msg, int Tab) {
-  tab(Tab);
-  llvm::errs() << "ENTRY POINTS (Module)"
-               << " " << msg << " {\n";
-  for (const auto &F : M) {
-    if (isEntryPoint(F, OnlyKernelsAreEntryPoints)) {
-      tab(Tab);
-      llvm::errs() << "  " << F.getName() << "\n";
-    }
-  }
-  tab(Tab);
-  llvm::errs() << "}\n";
-}
-
-#endif // NDEBUG
-
 void ModuleDesc::assignMergedProperties(const ModuleDesc &MD1,
                                         const ModuleDesc &MD2) {
   EntryPoints.Props = MD1.EntryPoints.Props.merge(MD2.EntryPoints.Props);
@@ -847,8 +798,57 @@ void ModuleDesc::verifyESIMDProperty() const {
   //  }
   //}
 }
+#endif // NDEBUG
+
+#if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP)
+LLVM_DUMP_METHOD const char *toString(SyclEsimdSplitStatus S) {
+  switch (S) {
+  case SyclEsimdSplitStatus::ESIMD_ONLY:
+    return "ESIMD_ONLY";
+  case SyclEsimdSplitStatus::SYCL_ONLY:
+    return "SYCL_ONLY";
+  case SyclEsimdSplitStatus::SYCL_AND_ESIMD:
+    return "SYCL_AND_ESIMD";
+  }
+  return "<UNKNOWN_STATUS>";
+}
 
-void ModuleDesc::dump() const {
+LLVM_DUMP_METHOD void tab(int N) {
+  for (int I = 0; I < N; ++I) {
+    llvm::errs() << "  ";
+  }
+}
+
+LLVM_DUMP_METHOD void dumpEntryPoints(const EntryPointSet &C, const char *msg,
+                                      int Tab) {
+  tab(Tab);
+  llvm::errs() << "ENTRY POINTS"
+               << " " << msg << " {\n";
+  for (const Function *F : C) {
+    tab(Tab);
+    llvm::errs() << "  " << F->getName() << "\n";
+  }
+  tab(Tab);
+  llvm::errs() << "}\n";
+}
+
+LLVM_DUMP_METHOD void dumpEntryPoints(const Module &M,
+                                      bool OnlyKernelsAreEntryPoints,
+                                      const char *msg, int Tab) {
+  tab(Tab);
+  llvm::errs() << "ENTRY POINTS (Module)"
+               << " " << msg << " {\n";
+  for (const auto &F : M) {
+    if (isEntryPoint(F, OnlyKernelsAreEntryPoints)) {
+      tab(Tab);
+      llvm::errs() << "  " << F.getName() << "\n";
+    }
+  }
+  tab(Tab);
+  llvm::errs() << "}\n";
+}
+
+LLVM_DUMP_METHOD void ModuleDesc::dump() const {
   llvm::errs() << "split_module::ModuleDesc[" << Name << "] {\n";
   llvm::errs() << "  ESIMD:" << toString(EntryPoints.Props.HasESIMD)
                << ", SpecConstMet:" << (Props.SpecConstsMet ? "YES" : "NO")
@@ -856,7 +856,7 @@ void ModuleDesc::dump() const {
   dumpEntryPoints(entries(), EntryPoints.GroupId.c_str(), 1);
   llvm::errs() << "}\n";
 }
-#endif // NDEBUG
+#endif
 
 void ModuleDesc::saveSplitInformationAsMetadata() {
   // Add metadata to the module so we can identify what kind of SYCL/ESIMD split