Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
Menu
Open sidebar
OpenDAS
FlashMLA
Commits
0ce8ee82
"docs-source/vscode:/vscode.git/clone" did not exist on "1bcd08d784a0b58dfd3535ecac2d565be56f54a9"
Commit
0ce8ee82
authored
Jan 26, 2026
by
zhanghj2
Browse files
fix 关闭attn sink情况下的错误
parent
200f01d5
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
1 addition
and
1 deletion
+1
-1
csrc/smxx/decode/combine/combine.cu
csrc/smxx/decode/combine/combine.cu
+1
-1
No files found.
csrc/smxx/decode/combine/combine.cu
View file @
0ce8ee82
...
@@ -86,7 +86,7 @@ flash_fwd_mla_combine_kernel(const CombineParams params) {
...
@@ -86,7 +86,7 @@ flash_fwd_mla_combine_kernel(const CombineParams params) {
float
global_lse
=
(
sum_lse
==
0.
f
||
sum_lse
==
-
INFINITY
)
?
INFINITY
:
logf
(
sum_lse
)
+
max_lse
;
float
global_lse
=
(
sum_lse
==
0.
f
||
sum_lse
==
-
INFINITY
)
?
INFINITY
:
logf
(
sum_lse
)
+
max_lse
;
if
(
lane_idx
==
0
)
if
(
lane_idx
==
0
)
gLse
(
warp_idx
)
=
global_lse
;
gLse
(
warp_idx
)
=
global_lse
;
float
o_scale
=
0
.0
f
;
float
o_scale
=
1
.0
f
;
if
(
params
.
attn_sink
!=
nullptr
)
{
if
(
params
.
attn_sink
!=
nullptr
)
{
int
q_head_idx
=
h_block_idx
*
BLOCK_SIZE_M
+
warp_idx
;
int
q_head_idx
=
h_block_idx
*
BLOCK_SIZE_M
+
warp_idx
;
float
attn_sink
=
__ldg
(
params
.
attn_sink
+
q_head_idx
);
float
attn_sink
=
__ldg
(
params
.
attn_sink
+
q_head_idx
);
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment