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
Commits
1377186f
Commit
1377186f
authored
Jun 08, 2022
by
wangshaojie6
Browse files
4 stage prefetch
parent
64fbf5a2
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
56 additions
and
37 deletions
+56
-37
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_skip_all_lds_v1.hpp
...eration/gpu/grid/gridwise_gemm_xdlops_skip_all_lds_v1.hpp
+56
-37
No files found.
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_skip_all_lds_v1.hpp
View file @
1377186f
...
...
@@ -567,16 +567,17 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_all_lds_v1
// preload data to regiester and LDS
{
// Read
a_threadwise_copy
.
Run
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_grid_buf
,
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_0
);
b_threadwise_copy
.
Run
(
b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
b_grid_buf
,
b_thread_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
b_thread_buf_0
);
a_threadwise_copy
.
Run
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_grid_buf
,
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_0
);
// Move
a_threadwise_copy
.
MoveSrcSliceWindow
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
...
...
@@ -585,16 +586,16 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_all_lds_v1
b_thread_slice_copy_step
);
// Read
a_threadwise_copy
.
Run
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_grid_buf
,
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_1
);
b_threadwise_copy
.
Run
(
b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
b_grid_buf
,
b_thread_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
b_thread_buf_1
);
a_threadwise_copy
.
Run
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_grid_buf
,
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_1
);
// Move
a_threadwise_copy
.
MoveSrcSliceWindow
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
...
...
@@ -603,16 +604,17 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_all_lds_v1
b_thread_slice_copy_step
);
// Read
a_threadwise_copy
.
Run
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_grid_buf
,
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_2
);
b_threadwise_copy
.
Run
(
b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
b_grid_buf
,
b_thread_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
b_thread_buf_2
);
a_threadwise_copy
.
Run
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_grid_buf
,
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_2
);
// Move
a_threadwise_copy
.
MoveSrcSliceWindow
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
...
...
@@ -620,6 +622,20 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_all_lds_v1
b_threadwise_copy
.
MoveSrcSliceWindow
(
b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
b_thread_slice_copy_step
);
b_threadwise_copy
.
Run
(
b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
b_grid_buf
,
b_thread_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
b_thread_buf_3
);
a_threadwise_copy
.
Run
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_grid_buf
,
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_3
);
// Initialize C
c_thread_buf
.
Clear
();
// a data write to lds
...
...
@@ -632,17 +648,7 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_all_lds_v1
index_t
i
=
0
;
do
{
static_for
<
0
,
MultiK0
,
4
>
{}([
&
](
auto
)
{
b_threadwise_copy
.
Run
(
b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
b_grid_buf
,
b_thread_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
b_thread_buf_3
);
a_threadwise_copy
.
Run
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_grid_buf
,
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_3
);
{
blockwise_gemm
.
Run
(
a_thread_buf_0
,
b_thread_buf_0
,
c_thread_buf
);
b_threadwise_copy
.
MoveSrcSliceWindow
(
b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
...
...
@@ -703,15 +709,6 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_all_lds_v1
b_thread_slice_copy_step
);
a_threadwise_copy
.
MoveSrcSliceWindow
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_thread_slice_copy_step
);
});
i
+=
1
;
}
while
(
i
<
(
K0BlockMainLoop
-
1
));
}
// tail
{
static_for
<
0
,
MultiK0
,
4
>
{}([
&
](
auto
i
)
{
b_threadwise_copy
.
Run
(
b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
b_grid_buf
,
...
...
@@ -723,6 +720,17 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_all_lds_v1
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_3
);
}
i
+=
1
;
}
while
(
i
<
(
K0BlockMainLoop
-
1
));
}
// tail
{
static_for
<
0
,
MultiK0
,
4
>
{}([
&
](
auto
i
)
{
blockwise_gemm
.
Run
(
a_thread_buf_0
,
b_thread_buf_0
,
c_thread_buf
);
...
...
@@ -786,6 +794,17 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdlops_skip_all_lds_v1
b_thread_slice_copy_step
);
a_threadwise_copy
.
MoveSrcSliceWindow
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_thread_slice_copy_step
);
b_threadwise_copy
.
Run
(
b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
b_grid_buf
,
b_thread_desc_k0_k1_k2_n0_n1_n2_n3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
b_thread_buf_3
);
a_threadwise_copy
.
Run
(
a_grid_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
a_grid_buf
,
a_thread_desc_k0_k1_k2_m0_m1_m2_m3_k3
,
make_tuple
(
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
,
I0
),
a_thread_buf_3
);
}
});
}
...
...
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