mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-25 15:24:39 +00:00
Merge commit '75570d0fa8fcaa4cc22f944d21c47cfec1af1943' into develop
This commit is contained in:
@@ -175,6 +175,9 @@ struct sequence
|
||||
return sequence<type::get(number<Ids>{})...>{};
|
||||
}
|
||||
|
||||
CK_TILE_HOST_DEVICE static constexpr auto sum() { return (Is + ... + 0); }
|
||||
CK_TILE_HOST_DEVICE static constexpr auto product() { return (Is * ... * 1); }
|
||||
|
||||
// modify element at index "I" with value "X"
|
||||
template <index_t I, index_t X>
|
||||
CK_TILE_HOST_DEVICE static constexpr auto modify(number<I>, number<X>)
|
||||
|
||||
@@ -153,4 +153,28 @@ struct CK_PRINTF_WARP0 : public CK_PRINTF<ConvertTo, FMT, PREFIX, SUFFIX>
|
||||
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
|
||||
|
||||
Reference in New Issue
Block a user