mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-05 22:22:27 +00:00

Files

Aviral Goel 6aa1cd8212 [rocm-libraries] ROCm/rocm-libraries#4704 (commit 17662f9)

[CK_TILE] Fix FP8 MXGEMM numerical error in async load path
 (#4704)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Summary

Fixes FP8 MXGEMM producing half the expected result (e.g., 128 instead
of 256 with all 1s input).

**Bug introduced in:** `b7de1e14cea70681a23cd1a136df42910c776e4a` -
"[CK_TILE] Add blockscale GEMM support for EightWarps on gfx950 (#4280)"

  ## Root Cause

In the `static_move_ys=true` code path in `tile_window.hpp`, the IMM
optimization computes `lds_ys_offset` using a default-constructed tensor
descriptor:

  ```cpp
  make_tensor_coordinate(decltype(tensor_descriptor){}, idx_ys_offset)
```

  This default-constructed descriptor has different strides than the actual DRAM tensor descriptor used for dram_ys_offset. When these offsets are mixed in the address   calculation:

```cpp

  imm_valid = lds_ys_offset % IMM_RANGE;     // From wrong descriptor
  wave_offset = dram_ys_offset - imm_valid;  // From correct descriptor
```

  The final address wave_offset + imm_valid ≠ dram_ys_offset, causing incorrect memory accesses.

  Fix
```cpp
Set imm_valid = 0 to bypass the IMM optimization and ensure the full
offset is passed through wave_offset:

constexpr auto imm_valid = 0; // Avoids inconsistency between
lds_ys_offset and dram_ys_offset
```

  This disables the 12-bit immediate field optimization in the buffer_load_lds instruction but guarantees correctness. A proper fix would require making the DRAM tensor   descriptor constexpr, which is not feasible since tensor strides depend on runtime parameters (LDA, LDB).

2026-02-23 20:30:32 +00:00

algorithm

[Compiler] Addressing new compiler warnings (#3640 )

2026-02-02 09:39:48 -08:00

arch

[rocm-libraries] ROCm/rocm-libraries#4584 (commit 42efd1d)

2026-02-21 01:15:57 +00:00

container

[rocm-libraries] ROCm/rocm-libraries#4280 (commit b7de1e1)

2026-02-09 03:55:52 +00:00

numeric

[Compiler] Addressing new compiler warnings (#3640 )

2026-02-02 09:39:48 -08:00

tensor

[rocm-libraries] ROCm/rocm-libraries#4704 (commit 17662f9)

2026-02-23 20:30:32 +00:00

utility

[Compiler] Addressing new compiler warnings (#3640 )

2026-02-02 09:39:48 -08:00

config.hpp

[rocm-libraries] ROCm/rocm-libraries#4640 (commit 37b8c81)

2026-02-18 15:00:26 +00:00

README.md

introducing ck_tile! (#1216 )

2024-04-15 19:27:12 -05:00

README.md

ck_tile/core

ck_tile/core contains every basic functions and structures to create a GPU kernel using ck_tile. User should only include ck_tile/core.hpp this single header to use all the functionality. Everything is under ck_tile namespace. The coding style under this folder should be similar to std (snake_case for structure/function, Camel for template types...)

algorithm/
    coordinate transform and some other reusable algorithm
arch/
    contains some basic device building block like mma, buffer addressing, etc...
container/
    contains basic container data structure, array/sequence/tuple/...
numeric/
    data type, and data type related math
tensor/
    tensor descriptors and tile level API
utility/
    other utility function for both host/device