Skip to content

Commit 15f4c3b

Browse files
committed
fix: fix fmha_v2 sync bug
Signed-off-by: Mingyang Jiang <[email protected]>
1 parent 18132b5 commit 15f4c3b

File tree

3 files changed

+3
-4
lines changed

3 files changed

+3
-4
lines changed

cpp/kernels/fmha_v2/src/fmha/hopper/fragment.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -486,7 +486,7 @@ struct Softmax_saver_tma
486486
int lane = threadIdx.x % Cta_tile::THREADS_PER_WARP;
487487
if (lane % 4 < 2)
488488
{
489-
values = p_sum[lane % 2] == 0.f ? 1.f : 1.0f / p_sum[lane % 2];
489+
values = p_sum[lane % 2];
490490
}
491491
else
492492
{

cpp/tensorrt_llm/kernels/mlaChunkedPrefill.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -159,8 +159,8 @@ inline __device__ void dequantCopy(
159159
// merged_softmax_sum [q_total_len, H, 2] (float, max/sum)
160160
template <typename T>
161161
__global__ void mergeAttnWithSoftmaxKernel(T* merged_attn, float2* merged_softmax_stats, T const* pre_attn,
162-
float2 const* pre_softmax_stats, T const* curr_attn, float2 const* curr_softmax_stats, float bmm1_scale,
163-
int64_t const* cu_q_seq_len, int64_t const* merge_op, int const num_heads, int const head_size)
162+
float2 const* pre_softmax_stats, T const* curr_attn, float2 const* curr_softmax_stats, int64_t const* cu_q_seq_len,
163+
int64_t const* merge_op, int const num_heads, int const head_size)
164164
{
165165
using KT = MergeSoftmaxTraits<T>;
166166
int const batch_idx = static_cast<int>(blockIdx.y);

tensorrt_llm/_torch/attention_backend/trtllm.py

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1449,7 +1449,6 @@ def merge_attention_for_mla(
14491449
temp_attn: torch.Tensor,
14501450
softmax_stats: torch.Tensor,
14511451
temp_softmax_stats: torch.Tensor,
1452-
bmm1_scale: float,
14531452
merge_op: torch.Tensor,
14541453
metadata: TrtllmAttentionMetadata,
14551454
) -> None:

0 commit comments

Comments
 (0)