Skip to content

Commit 1775084

Browse files
Xiao YUXiao YU
authored andcommitted
Update fp8 PV index
1 parent 418752b commit 1775084

1 file changed

Lines changed: 2 additions & 2 deletions

File tree

csrc/rocm/attention.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -817,8 +817,8 @@ __launch_bounds__(NUM_THREADS, 5) void paged_attention_ll4mi_QKV_mfma16_kernel(
817817
const int offset =
818818
rowid * ELEMS16_ELEMS8_RATIO * ELEMS8_ELEMS4_RATIO +
819819
j * ELEMS8_ELEMS4_RATIO + i;
820-
const int offset1 = offset % ROWS_PER_WARP;
821-
const int offset2 = offset / ROWS_PER_WARP / 2;
820+
const int offset1 = (offset % ROWS_PER_WARP) / 2;
821+
const int offset2 = offset / ROWS_PER_WARP;
822822
// output format is 16 qheads across 16 lanes, 16 head elems
823823
// spread across 4 rows
824824
tmp_out = gcn_mfma16x16x32_instr<__hip_fp8_e4m3, 0, 0, 0>(

0 commit comments

Comments
 (0)