Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel_ROCM
Commits
ae2d7d2b
Commit
ae2d7d2b
authored
Oct 12, 2024
by
danyao12
Browse files
code revert
parent
e2ea64d9
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
7 additions
and
7 deletions
+7
-7
example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py
example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py
+7
-7
No files found.
example/ck_tile/01_fmha/codegen/ops/fmha_bwd.py
View file @
ae2d7d2b
...
...
@@ -1012,14 +1012,14 @@ class FmhaBwdDQDKDVKernel:
def
get_fmha_bwd_dq_dk_dv_tile_ppl_dict_from_dtype
(
dtype
:
str
)
->
Optional
[
dict
]:
if
dtype
==
'fp16'
or
dtype
==
'bf16'
:
return
{
#
'32' : [FmhaBwdDQDKDVTileSize( 32, 128, 32, 32, 32, 32, 64, 32, 32, 1, 4, 1, 4, 1, 1, 2, 2, 1, 16, 16, 32, 16, 16, 16, 1),
#
"kr_ktr_vr_iglp", "kr_ktr_vr"],
#
'64' : [FmhaBwdDQDKDVTileSize( 32, 128, 64, 32, 64, 32, 32, 64, 64, 1, 4, 1, 4, 1, 1, 1, 4, 1, 16, 16, 32, 16, 16, 16, 1),
#
"kr_ktr_vr_iglp", "kr_ktr_vr"],
'32'
:
[
FmhaBwdDQDKDVTileSize
(
32
,
128
,
32
,
32
,
32
,
32
,
64
,
32
,
32
,
1
,
4
,
1
,
4
,
1
,
1
,
2
,
2
,
1
,
16
,
16
,
32
,
16
,
16
,
16
,
1
),
"kr_ktr_vr_iglp"
,
"kr_ktr_vr"
],
'64'
:
[
FmhaBwdDQDKDVTileSize
(
32
,
128
,
64
,
32
,
64
,
32
,
32
,
64
,
64
,
1
,
4
,
1
,
4
,
1
,
1
,
1
,
4
,
1
,
16
,
16
,
32
,
16
,
16
,
16
,
1
),
"kr_ktr_vr_iglp"
,
"kr_ktr_vr"
],
'128'
:
[
FmhaBwdDQDKDVTileSize
(
16
,
128
,
128
,
16
,
128
,
16
,
32
,
128
,
128
,
1
,
4
,
1
,
4
,
1
,
1
,
1
,
4
,
1
,
16
,
16
,
32
,
16
,
16
,
16
,
1
),
"kr_ktr_vr_iglp"
,
"kr_ktr_vr"
],
#
'256' : [FmhaBwdDQDKDVTileSize( 16, 64, 256, 16, 256, 16, 32, 256, 256, 1, 4, 1, 4, 1, 1, 1, 4, 1, 16, 16, 32, 16, 16, 16, 1),
#
"kr_ktr_vr_iglp", "kr_ktr_vr"]
'256'
:
[
FmhaBwdDQDKDVTileSize
(
16
,
64
,
256
,
16
,
256
,
16
,
32
,
256
,
256
,
1
,
4
,
1
,
4
,
1
,
1
,
1
,
4
,
1
,
16
,
16
,
32
,
16
,
16
,
16
,
1
),
"kr_ktr_vr_iglp"
,
"kr_ktr_vr"
]
}
else
:
return
None
...
...
@@ -1062,7 +1062,7 @@ def get_bwd_dq_dk_dv_blobs(kernel_filter : Optional[str], receipt, mask_impl) ->
continue
if
receipt
==
3
:
cond
=
dtype
in
[
'fp16'
,
'bf16'
]
cond
&=
bias
in
[
'no'
]
cond
&=
bias
in
[
'no'
,
'alibi'
]
cond
&=
dpad
==
dvpad
cond
&=
deterministic
==
"f"
if
not
cond
:
...
...
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