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
cde3b677
Commit
cde3b677
authored
Feb 03, 2025
by
Qianfeng Zhang
Browse files
Move the definition of v_tiles out from the loop
parent
08598523
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
6 additions
and
6 deletions
+6
-6
include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_async.hpp
.../ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_async.hpp
+6
-6
No files found.
include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_async.hpp
View file @
cde3b677
...
...
@@ -309,6 +309,10 @@ struct BlockFmhaPipelineQRKSVSAsync
{
0
,
seqlen_k_start
},
// TODO: hdim split?
Policy
::
template
MakeVDramTileDistribution
<
Problem
>());
using
v_tile_type
=
decltype
(
load_tile
(
v_dram_window
));
statically_indexed_array
<
v_tile_type
,
NumVLdsBuffers
>
v_tiles
;
index_t
i_total_loops
=
0
;
do
...
...
@@ -467,12 +471,8 @@ struct BlockFmhaPipelineQRKSVSAsync
const
auto
bias_tile
=
load_tile
(
bias_dram_window
);
// load bias tile
using
v_tile_type
=
decltype
(
load_tile
(
v_dram_window
));
statically_indexed_array
<
v_tile_type
,
NumVLdsBuffers
>
v_tiles
;
static_for
<
0
,
NumVLdsBuffers
,
1
>
{}([
&
](
auto
i_k1
)
{
v_tiles
[
i_k1
]
=
load_tile
(
v_dram_window
);
static_for
<
0
,
NumVLdsBuffers
,
1
>
{}([
&
](
auto
i_buf
)
{
v_tiles
[
i_buf
]
=
load_tile
(
v_dram_window
);
move_tile_window
(
v_dram_window
,
{
0
,
kK1
});
});
...
...
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