minor change

root · root · commit 9d84889927f9 · 2026-04-17T06:55:10.000Z
diff --git a/csrc/xpu/gdn_attn/xe_2/chunk_gated_delta_rule_kernels_xe2.hpp b/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

Original file line number	Diff line number	Diff line change
`@@ -1016,8 +1016,8 @@ CUTE_DEVICE void chunk_fwd_o_kernel(`
`1016`	`1016`	`float g_cumsum_value =`
`1017`	`1017`	`a[(chunk_offset + e) + v_head_id * total_virtual_seqlen];`
`1018`	`1018`	`g_slm_ptr[e] = g_cumsum_value;`
`1019`		`- g_multi_slm_ptr[e] = sycl::native::exp(g_last_value - g_cumsum_value);`
`1020`	`1019`	`g_exp_slm_ptr[e] = sycl::native::exp(g_cumsum_value);`
	`1020`	`+ g_multi_slm_ptr[e] = g_last_value_exp / g_exp_slm_ptr[e];`
`1021`	`1021`	`}`
`1022`	`1022`
`1023`	`1023`	`CUTE_UNROLL`