Skip to content

Commit 92b0738

Browse files
[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 <[email protected]> * cleanup docstring * clang-format --------- Co-authored-by: Aviral Goel <[email protected]>
1 parent 0f8e33f commit 92b0738

File tree

1 file changed

+24
-0
lines changed

1 file changed

+24
-0
lines changed

include/ck_tile/core/utility/debug.hpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -153,4 +153,28 @@ struct CK_PRINTF_WARP0 : public CK_PRINTF<ConvertTo, FMT, PREFIX, SUFFIX>
153153
base_t::operator()(buf);
154154
}
155155
};
156+
157+
/*
158+
* RAII struct which inserts start/end markers into the generated assembly.
159+
*
160+
* Usage:
161+
* - Create an `AsmScopeMarker` object at the beginning of a scope or code block.
162+
* - Its constructor will emit a "CK_ASM_SCOPE_START" marker into the assembly.
163+
* - When the object goes out of scope (end of block, return, exception, etc.),
164+
* the destructor will emit a "CK_ASM_SCOPE_END" marker.
165+
*
166+
* Example:
167+
* {
168+
* [[maybe_unused]] AsmScopeMarker marker; // Emits CK_ASM_SCOPE_START
169+
* // ... code you want to delimit in assembly ...
170+
* } // marker goes out of scope → Emits CK_ASM_SCOPE_END
171+
*
172+
*/
173+
struct AsmScopeMarker
174+
{
175+
// in some future version of clang we might be able to use string_view to customize
176+
CK_TILE_HOST_DEVICE AsmScopeMarker() { asm volatile(";;# CK_ASM_SCOPE_START"); }
177+
CK_TILE_HOST_DEVICE ~AsmScopeMarker() { asm volatile(";;# CK_ASM_SCOPE_END"); }
178+
};
179+
156180
} // namespace ck_tile

0 commit comments

Comments
 (0)