From 8a34c5087aa723eafd9995a3af814fcae8334c4d Mon Sep 17 00:00:00 2001 From: Andrew Barnes Date: Thu, 9 Apr 2026 03:12:22 -0400 Subject: [PATCH] [ROCm] Remove unnecessary fp8 roundtrip in gather cache NHD dequant (#39122) Signed-off-by: Bortlesboat --- vllm/v1/attention/backends/rocm_aiter_fa.py | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/vllm/v1/attention/backends/rocm_aiter_fa.py b/vllm/v1/attention/backends/rocm_aiter_fa.py index abbbcafb51..a4cdce3f5d 100644 --- a/vllm/v1/attention/backends/rocm_aiter_fa.py +++ b/vllm/v1/attention/backends/rocm_aiter_fa.py @@ -112,10 +112,12 @@ if current_platform.is_rocm(): if DEQUANT: k_scale = tl.load(k_scale_ptr) v_scale = tl.load(v_scale_ptr) - k_dtype = k_reg.dtype - v_dtype = v_reg.dtype - k_reg = (k_reg.to(tl.float32) * k_scale).to(k_dtype) - v_reg = (v_reg.to(tl.float32) * v_scale).to(v_dtype) + k_reg = (k_reg.to(tl.float32) * k_scale).to( + key_ptr_offset.dtype.element_ty + ) + v_reg = (v_reg.to(tl.float32) * v_scale).to( + value_ptr_offset.dtype.element_ty + ) tl.store(key_ptr_offset + col_offsets, k_reg) tl.store(value_ptr_offset + col_offsets, v_reg)