Skip to content

[CK_TILE] Fix FP8 MXGEMM numerical error in async load path #4704

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
ThomasNing merged 6 commits into from
Feb 23, 2026
Merged

[CK_TILE] Fix FP8 MXGEMM numerical error in async load path #4704

Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
d4c4b36
[CK_TILE] Fix FP8 MXGEMM numerical error in async load path
AviralGoelAMD Feb 19, 2026
8dd4faf
Finish the fix
Feb 19, 2026
3ec0228
fix the error
Feb 20, 2026
32b10fb
clang format
Feb 20, 2026
fafd833
Merge branch 'develop' into ck/aviralgoel/fix-fp8-mxgemm-imm-offset
Feb 20, 2026
711663a
Merge branch 'develop' into ck/aviralgoel/fix-fp8-mxgemm-imm-offset
AviralGoelAMD Feb 22, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 6 additions & 12 deletions projects/composablekernel/include/ck_tile/core/tensor/tile_window.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -627,16 +627,9 @@ struct tile_window_with_static_distribution
const auto lds_coord =
make_tensor_coordinate(tensor_descriptor, lds_bottom_tensor_thread_idx);

constexpr auto IMM_RANGE =
(1 << 12) / sizeof(typename Base::DataType) * Traits::PackedSize;
constexpr auto imm_total = lds_ys_offset;
constexpr auto imm_valid = imm_total % IMM_RANGE;
constexpr auto imm_overflow = imm_total - imm_valid;

// Calculate SMEM address using base pointer
CK_TILE_LDS_ADDR LdsDataType* smem = lds_base_ptr +
lds_coord.get_offset() / Traits::PackedSize +
imm_overflow / Traits::PackedSize;
CK_TILE_LDS_ADDR LdsDataType* smem =
lds_base_ptr + (lds_coord.get_offset() + lds_ys_offset) / Traits::PackedSize;

const auto dram_ys_offset = [&]() {
if constexpr(static_move_ys)
Expand All @@ -656,13 +649,14 @@ struct tile_window_with_static_distribution
offset + dram_ys_offset,
bool_constant<oob_conditional_check>{});
else
{
this->get_bottom_tensor_view().template async_get_vectorized_elements<vector_t>(
smem,
bottom_tensor_thread_coord.get_offset() + offset,
dram_ys_offset - imm_valid,
number<imm_valid>{},
dram_ys_offset,
number<0>{},
bool_constant<oob_conditional_check>{});

}
// Move thread coordinate if not last access
if constexpr(iCoordAccess != (NumAccessPerCoord - 1))
{
Expand Down
Loading