mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-13 17:55:48 +00:00
Fix gfx12 async tile-load fallback warnings
This commit is contained in:
@@ -164,6 +164,7 @@ CK_TILE_DEVICE void async_load_tile_with_offset(LdsTileWindow_&& lds_tile,
|
||||
// gfx12 does not support the direct global-to-LDS async buffer load used by
|
||||
// tile_window::async_load*. Keep the architecture choice at the tile-load layer by using a
|
||||
// regular buffer load followed by an explicit LDS store instead of hiding it in the primitive.
|
||||
static_cast<void>(smy);
|
||||
static_assert(!static_move_ys,
|
||||
"gfx12 synchronous async_load_tile fallback does not support static_move_ys");
|
||||
auto tile = load_tile_with_offset(tile_window, offset, number<i_access>{}, occ);
|
||||
|
||||
@@ -495,7 +495,9 @@ struct tile_window_with_static_distribution
|
||||
{
|
||||
using LdsTileWindow = remove_cvref_t<LdsTileWindow_>;
|
||||
// using LdsTensorView = typename LdsTileWindow::BottomTensorView;
|
||||
#if !defined(__gfx12__)
|
||||
using LdsDataType = typename LdsTileWindow::DataType;
|
||||
#endif
|
||||
// using LdsDescriptor = typename LdsTileWindow::BottomTensorDesc;
|
||||
|
||||
// issues * warps * lanes
|
||||
|
||||
@@ -488,7 +488,9 @@ struct tile_window_linear
|
||||
bool_constant<pre_nop> = {}) const
|
||||
{
|
||||
using LdsTileWindow = remove_cvref_t<LdsTileWindow_>;
|
||||
#if !defined(__gfx12__)
|
||||
using LdsDataType = typename LdsTileWindow::DataType;
|
||||
#endif
|
||||
|
||||
// currently we only support everything is non linear dim
|
||||
// actually it's not performant if we have linear dim(e.g. fast changing)
|
||||
|
||||
Reference in New Issue
Block a user