From a9310055fa0d7fa1a3a310eb76da4c7ab2b02f62 Mon Sep 17 00:00:00 2001 From: "Sabianin, Maksim" Date: Thu, 2 Oct 2025 06:57:37 -0700 Subject: [PATCH 1/3] [SYCL][ESIMD] Decorate dump methods according to LLVM guideline. --- llvm/include/llvm/SYCLPostLink/ModuleSplitter.h | 12 +++++++----- .../SYCLPostLink/ESIMDPostSplitProcessing.cpp | 16 +--------------- 2 files changed, 8 insertions(+), 20 deletions(-) diff --git a/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h b/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h index 8ad837beb7da..7cf28d6a5b20 100644 --- a/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h +++ b/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h @@ -303,11 +303,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 10b1c66bbe04..f0df7fa2b3a5 100644 --- a/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp +++ b/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp @@ -27,17 +27,6 @@ #include #include -#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); From 4c826b7c51da329749d0c7b64a64dc4ea8be66a6 Mon Sep 17 00:00:00 2001 From: "Sabianin, Maksim" Date: Mon, 6 Oct 2025 05:12:16 -0700 Subject: [PATCH 2/3] decorate ModuleSplitter::dump method --- llvm/include/llvm/SYCLPostLink/ModuleSplitter.h | 5 ++++- llvm/lib/SYCLPostLink/ModuleSplitter.cpp | 8 +++++--- 2 files changed, 9 insertions(+), 4 deletions(-) diff --git a/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h b/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h index 7cf28d6a5b20..1aaf345b7257 100644 --- a/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h +++ b/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h @@ -238,8 +238,11 @@ class ModuleDesc { #ifndef NDEBUG void verifyESIMDProperty() const; - void dump() const; #endif // NDEBUG + +#if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP) + void dump() const; +#endif }; // Module split support interface. diff --git a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp b/llvm/lib/SYCLPostLink/ModuleSplitter.cpp index 24112fab71c1..9ed0a162ed53 100644 --- a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp +++ b/llvm/lib/SYCLPostLink/ModuleSplitter.cpp @@ -561,7 +561,7 @@ Error ModuleSplitterBase::verifyNoCrossModuleDeviceGlobalUsage() { return Error::success(); } -#ifndef NDEBUG +#if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP) const char *toString(SyclEsimdSplitStatus S) { switch (S) { @@ -608,7 +608,7 @@ void dumpEntryPoints(const Module &M, bool OnlyKernelsAreEntryPoints, llvm::errs() << "}\n"; } -#endif // NDEBUG +#endif void ModuleDesc::assignMergedProperties(const ModuleDesc &MD1, const ModuleDesc &MD2) { @@ -847,7 +847,9 @@ void ModuleDesc::verifyESIMDProperty() const { // } //} } +#endif // NDEBUG +#if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP) void ModuleDesc::dump() const { llvm::errs() << "split_module::ModuleDesc[" << Name << "] {\n"; llvm::errs() << " ESIMD:" << toString(EntryPoints.Props.HasESIMD) @@ -856,7 +858,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 From 160b87b93926527fc3b9121ae6cbed13892ada31 Mon Sep 17 00:00:00 2001 From: "Sabianin, Maksim" Date: Mon, 6 Oct 2025 05:20:59 -0700 Subject: [PATCH 3/3] add missed decoration and group code together --- .../llvm/SYCLPostLink/ModuleSplitter.h | 3 +- llvm/lib/SYCLPostLink/ModuleSplitter.cpp | 98 +++++++++---------- 2 files changed, 50 insertions(+), 51 deletions(-) diff --git a/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h b/llvm/include/llvm/SYCLPostLink/ModuleSplitter.h index 1aaf345b7257..144f99b753e7 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" @@ -241,7 +242,7 @@ class ModuleDesc { #endif // NDEBUG #if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP) - void dump() const; + LLVM_DUMP_METHOD void dump() const; #endif }; diff --git a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp b/llvm/lib/SYCLPostLink/ModuleSplitter.cpp index 9ed0a162ed53..30647bd82604 100644 --- a/llvm/lib/SYCLPostLink/ModuleSplitter.cpp +++ b/llvm/lib/SYCLPostLink/ModuleSplitter.cpp @@ -561,55 +561,6 @@ Error ModuleSplitterBase::verifyNoCrossModuleDeviceGlobalUsage() { return Error::success(); } -#if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP) - -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 ""; -} - -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 - void ModuleDesc::assignMergedProperties(const ModuleDesc &MD1, const ModuleDesc &MD2) { EntryPoints.Props = MD1.EntryPoints.Props.merge(MD2.EntryPoints.Props); @@ -850,7 +801,54 @@ void ModuleDesc::verifyESIMDProperty() const { #endif // NDEBUG #if !defined(NDEBUG) || defined(LLVM_ENABLE_DUMP) -void ModuleDesc::dump() const { +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 ""; +} + +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")