From c1a4d4c5f1ee846e8edb16be4572b44ca068f4a6 Mon Sep 17 00:00:00 2001 From: Max Podkorytov <4273004+tenpercent@users.noreply.github.com> Date: Mon, 8 Sep 2025 22:02:02 -0700 Subject: [PATCH] [Util] add a RAII stuct which inserts markers into generated asm (#2748) * add asm scope raii printer * add comment * clang-format * compress * add Aviral's suggestion to extend the docstring Thanks~ Co-authored-by: Aviral Goel * cleanup docstring * clang-format --------- Co-authored-by: Aviral Goel [ROCm/composable_kernel commit: 92b07380d3a9978608d3a2e36c39162e9d4784ce] --- include/ck_tile/core/utility/debug.hpp | 24 ++++++++++++++++++++++++ 1 file changed, 24 insertions(+) diff --git a/include/ck_tile/core/utility/debug.hpp b/include/ck_tile/core/utility/debug.hpp index 15f0718dc2..9f0f931bc8 100644 --- a/include/ck_tile/core/utility/debug.hpp +++ b/include/ck_tile/core/utility/debug.hpp @@ -153,4 +153,28 @@ struct CK_PRINTF_WARP0 : public CK_PRINTF base_t::operator()(buf); } }; + +/* + * RAII struct which inserts start/end markers into the generated assembly. + * + * Usage: + * - Create an `AsmScopeMarker` object at the beginning of a scope or code block. + * - Its constructor will emit a "CK_ASM_SCOPE_START" marker into the assembly. + * - When the object goes out of scope (end of block, return, exception, etc.), + * the destructor will emit a "CK_ASM_SCOPE_END" marker. + * + * Example: + * { + * [[maybe_unused]] AsmScopeMarker marker; // Emits CK_ASM_SCOPE_START + * // ... code you want to delimit in assembly ... + * } // marker goes out of scope → Emits CK_ASM_SCOPE_END + * + */ +struct AsmScopeMarker +{ + // in some future version of clang we might be able to use string_view to customize + CK_TILE_HOST_DEVICE AsmScopeMarker() { asm volatile(";;# CK_ASM_SCOPE_START"); } + CK_TILE_HOST_DEVICE ~AsmScopeMarker() { asm volatile(";;# CK_ASM_SCOPE_END"); } +}; + } // namespace ck_tile