fix

Signed-off-by: yangqun <qun.yang@intel.com>

Author: YangQun1

csrc/xpu/gdn_attn/xe_2/chunk_gated_delta_rule_kernels_xe2.hpp

@@ -1016,8 +1016,8 @@ CUTE_DEVICE void chunk_fwd_o_kernel(
         float g_cumsum_value =
             a[(chunk_offset + e) + v_head_id * total_virtual_seqlen];
         g_slm_ptr[e] = g_cumsum_value;
+        g_multi_slm_ptr[e] = sycl::native::exp(g_last_value - g_cumsum_value);
         g_exp_slm_ptr[e] = sycl::native::exp(g_cumsum_value);
-        g_multi_slm_ptr[e] = g_last_value_exp / g_exp_slm_ptr[e];
       }
 
       CUTE_UNROLL
@@ -1148,8 +1148,7 @@ CUTE_DEVICE void chunk_fwd_o_kernel(
           // Fused gemm: W×S[dv] -> tSrU_d, Q×S[dv] -> tSrO_c
           // S is loaded once and reused in registers for both
           gemm_TTS_fused_2A(
-              W_tensor, Q_tensor, S_tensor,
-              tSrU_d, tSrO_c, 0, 0, dv, mma);
+              W_tensor, Q_tensor, S_tensor, tSrU_d, tSrO_c, 0, 0, dv, mma);
 
           // --- WS epilogue: U_new[dv] = U_old[dv] - W×S[dv] ---
           auto tCrU_d = thr_copy_U_d.partition_sg_fragment_S(gU_C);

csrc/xpu/gdn_attn/xe_2/gemm.hpp

@@ -407,12 +407,9 @@ CUTE_DEVICE void gemm_TTS_fused_2A(
 
   auto wg_tile = mma.tile_mnk();
 
-  Tensor gA1 = local_tile(
-      cA1, select<0, 2>(wg_tile), make_coord(wg_m1, _));
-  Tensor gA2 = local_tile(
-      cA2, select<0, 2>(wg_tile), make_coord(wg_m2, _));
-  Tensor gB = local_tile(
-      cB, select<1, 2>(wg_tile), make_coord(wg_n, _));
+  Tensor gA1 = local_tile(cA1, select<0, 2>(wg_tile), make_coord(wg_m1, _));
+  Tensor gA2 = local_tile(cA2, select<0, 2>(wg_tile), make_coord(wg_m2, _));
+  Tensor gB = local_tile(cB, select<1, 2>(wg_tile), make_coord(wg_n, _));
 
   auto copy_a1 = get_block_2d_copy_A<void>(mma, A1);
   auto copy_a2 = get_block_2d_copy_A<void>(mma, A2);