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
23602a0d
"git@developer.sourcefind.cn:gaoqiong/composable_kernel.git" did not exist on "f16356d4ff7d7f4f0c28c48c89a55d2372b2a3f7"
Commit
23602a0d
authored
Dec 17, 2024
by
Po Yen Chen
Browse files
Use better num_splits heuristic
parent
289d5eb0
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
10 additions
and
8 deletions
+10
-8
example/ck_tile/01_fmha/fmha_fwd.hpp
example/ck_tile/01_fmha/fmha_fwd.hpp
+10
-8
No files found.
example/ck_tile/01_fmha/fmha_fwd.hpp
View file @
23602a0d
...
...
@@ -12,6 +12,7 @@
#include "mask.hpp"
#include "rotary.hpp"
#include <array>
#include <type_traits>
#include <utility>
#include <variant>
...
...
@@ -827,26 +828,27 @@ Int num_splits_heuristic(Int batch_nhead_mblocks, Int num_SMs, Int max_splits)
max_splits
=
std
::
min
({
max_splits
,
num_SMs
});
constexpr
std
::
array
<
Int
,
5
>
num_splits_array
=
{
1
,
2
,
4
,
8
,
16
};
float
max_efficiency
=
0.
f
;
std
::
vector
<
float
>
efficiency
;
efficiency
.
reserve
(
max_splits
);
std
::
array
<
float
,
num_splits_array
.
size
()
>
efficiency
;
for
(
Int
num_splits
=
1
;
num_splits
<=
max_splits
;
num_splits
*=
2
)
for
(
size_t
idx
=
0
;
idx
<
num_splits_array
.
size
()
&&
num_splits_array
[
idx
]
<=
max_splits
;
++
idx
)
{
float
n_blocks
=
float
(
batch_nhead_mblocks
*
num_splits
)
/
num_SMs
;
float
n_blocks
=
float
(
batch_nhead_mblocks
*
num_splits
_array
[
idx
]
)
/
num_SMs
;
float
eff
=
n_blocks
/
std
::
ceil
(
n_blocks
);
if
(
eff
>
max_efficiency
)
{
max_efficiency
=
eff
;
}
efficiency
.
push_back
(
eff
)
;
efficiency
[
idx
]
=
eff
;
}
for
(
Int
num_splits
=
1
;
num_splits
<=
max_splits
;
num_splits
++
)
for
(
size_t
idx
=
0
;
idx
<
num_splits_array
.
size
()
&&
num_splits_array
[
idx
]
<=
max_splits
;
++
idx
)
{
if
(
efficiency
[
num_splits
-
1
]
>=
0.85
*
max_efficiency
)
if
(
efficiency
[
idx
]
>=
0.85
*
max_efficiency
)
{
return
num_splits
;
return
num_splits
_array
[
idx
]
;
}
}
return
1
;
...
...
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