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
ea2c27ca
Commit
ea2c27ca
authored
Jul 12, 2023
by
ltqin
Browse files
add limit for CDE0BlockTransferSrcScalarPerVector
parent
fc50b641
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
4 additions
and
0 deletions
+4
-0
include/ck/tensor_operation/gpu/grid/gridwise_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle_v1.hpp
...tched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle_v1.hpp
+4
-0
No files found.
include/ck/tensor_operation/gpu/grid/gridwise_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle_v1.hpp
View file @
ea2c27ca
...
@@ -734,6 +734,10 @@ struct GridwiseBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
...
@@ -734,6 +734,10 @@ struct GridwiseBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
const
auto
wave_id
=
GetGemm0WaveIdx
();
const
auto
wave_id
=
GetGemm0WaveIdx
();
const
auto
wave_m_n_id
=
GetGemm0WaveMNIdx
(
wave_id
[
I2
]);
// I2: 0~63
const
auto
wave_m_n_id
=
GetGemm0WaveMNIdx
(
wave_id
[
I2
]);
// I2: 0~63
static_assert
(
CDE0BlockTransferSrcScalarPerVector
<=
n4
,
"vector load must be not greater than n4"
);
static_assert
(
n4
%
CDE0BlockTransferSrcScalarPerVector
==
0
);
auto
d0s_threadwise_copy
=
generate_tuple
(
auto
d0s_threadwise_copy
=
generate_tuple
(
[
&
](
auto
i
)
{
[
&
](
auto
i
)
{
return
ThreadwiseTensorSliceTransfer_v2
<
return
ThreadwiseTensorSliceTransfer_v2
<
...
...
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