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
76871a6f
Commit
76871a6f
authored
Jan 30, 2025
by
Qianfeng Zhang
Browse files
Move clear_tile(s_acc) for better interleaving
parent
a979d030
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
6 additions
and
2 deletions
+6
-2
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
-2
No files found.
include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_async.hpp
View file @
76871a6f
...
@@ -377,8 +377,6 @@ struct BlockFmhaPipelineQRKSVSAsync
...
@@ -377,8 +377,6 @@ struct BlockFmhaPipelineQRKSVSAsync
}
}
else
// executed by intermediate and last iteration
else
// executed by intermediate and last iteration
{
{
clear_tile
(
s_acc
);
// initialize C
if
(
i_total_loops
<
num_total_loop
-
1
)
if
(
i_total_loops
<
num_total_loop
-
1
)
{
{
move_tile_window
(
k_dram_window
,
{
kN0
,
0
});
move_tile_window
(
k_dram_window
,
{
kN0
,
0
});
...
@@ -390,6 +388,9 @@ struct BlockFmhaPipelineQRKSVSAsync
...
@@ -390,6 +388,9 @@ struct BlockFmhaPipelineQRKSVSAsync
sequence
<
((
i_k0
%
NumKLdsBuffers
)
+
1
)
*
kN0
,
kK0
>
{});
sequence
<
((
i_k0
%
NumKLdsBuffers
)
+
1
)
*
kN0
,
kK0
>
{});
store_tile
(
k_lds_window_tmp
,
k_tiles
[
number
<
i_k0
>
{}]);
store_tile
(
k_lds_window_tmp
,
k_tiles
[
number
<
i_k0
>
{}]);
if
constexpr
(
i_k0
==
0
)
clear_tile
(
s_acc
);
k_tiles
[
number
<
i_k0
>
{}]
=
load_tile
(
k_dram_window
);
k_tiles
[
number
<
i_k0
>
{}]
=
load_tile
(
k_dram_window
);
if
constexpr
(
i_k0
<
k0_loops
-
1
)
if
constexpr
(
i_k0
<
k0_loops
-
1
)
move_tile_window
(
k_dram_window
,
{
0
,
kK0
});
move_tile_window
(
k_dram_window
,
{
0
,
kK0
});
...
@@ -413,6 +414,9 @@ struct BlockFmhaPipelineQRKSVSAsync
...
@@ -413,6 +414,9 @@ struct BlockFmhaPipelineQRKSVSAsync
sequence
<
((
i_k0
%
NumKLdsBuffers
)
+
1
)
*
kN0
,
kK0
>
{});
sequence
<
((
i_k0
%
NumKLdsBuffers
)
+
1
)
*
kN0
,
kK0
>
{});
store_tile
(
k_lds_window_tmp
,
k_tiles
[
number
<
i_k0
>
{}]);
store_tile
(
k_lds_window_tmp
,
k_tiles
[
number
<
i_k0
>
{}]);
if
constexpr
(
i_k0
==
0
)
clear_tile
(
s_acc
);
block_sync_lds
();
block_sync_lds
();
// execute last unroll of gemm_0
// execute last unroll of gemm_0
gemm_0
(
s_acc
,
gemm_0
(
s_acc
,
...
...
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