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
6ebaa81e
Commit
6ebaa81e
authored
Jun 06, 2024
by
Adam Osewski
Browse files
Fix usage of CK_ENV
parent
129e58ae
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
15 additions
and
15 deletions
+15
-15
include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_splitk_xdl_cshuffle_tile_loop.hpp
...grouped_gemm_multiple_d_splitk_xdl_cshuffle_tile_loop.hpp
+2
-2
include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_splitk_cshuffle_v2.hpp
.../grid/gridwise_gemm_multiple_d_xdl_splitk_cshuffle_v2.hpp
+12
-12
profiler/include/profiler/profile_grouped_gemm_multiple_d_splitk_impl.hpp
.../profiler/profile_grouped_gemm_multiple_d_splitk_impl.hpp
+1
-1
No files found.
include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_splitk_xdl_cshuffle_tile_loop.hpp
View file @
6ebaa81e
...
@@ -787,7 +787,7 @@ struct DeviceGroupedGemmMultipleDSplitKXdlCShuffle
...
@@ -787,7 +787,7 @@ struct DeviceGroupedGemmMultipleDSplitKXdlCShuffle
if
((
ck
::
type_convert
<
ck
::
index_t
>
(
arg
.
gemm_kernel_args_
.
size
())
+
if
((
ck
::
type_convert
<
ck
::
index_t
>
(
arg
.
gemm_kernel_args_
.
size
())
+
arg
.
skipped_group_count_
)
!=
arg
.
group_count_
)
arg
.
skipped_group_count_
)
!=
arg
.
group_count_
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"The group count is not equal to sum of skipped groups "
std
::
cout
<<
"The group count is not equal to sum of skipped groups "
"and kernel args size!"
"and kernel args size!"
...
@@ -811,7 +811,7 @@ struct DeviceGroupedGemmMultipleDSplitKXdlCShuffle
...
@@ -811,7 +811,7 @@ struct DeviceGroupedGemmMultipleDSplitKXdlCShuffle
arg
.
K_BATCH
);
arg
.
K_BATCH
);
if
(
not
group_arg_valid
)
if
(
not
group_arg_valid
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"["
<<
__func__
<<
"] group id: "
<<
i
std
::
cout
<<
"["
<<
__func__
<<
"] group id: "
<<
i
<<
" has invalid GridwiseGemm settings!"
<<
std
::
endl
;
<<
" has invalid GridwiseGemm settings!"
<<
std
::
endl
;
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_splitk_cshuffle_v2.hpp
View file @
6ebaa81e
...
@@ -415,7 +415,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -415,7 +415,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
{
{
if
(
!
(
M
%
MPerBlock
==
0
))
if
(
!
(
M
%
MPerBlock
==
0
))
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"Arg M value is not a multiple of MPerBlock! M: "
<<
M
<<
" "
std
::
cout
<<
"Arg M value is not a multiple of MPerBlock! M: "
<<
M
<<
" "
<<
__FILE__
<<
":"
<<
__LINE__
<<
", in function: "
<<
__func__
<<
__FILE__
<<
":"
<<
__LINE__
<<
", in function: "
<<
__func__
...
@@ -429,7 +429,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -429,7 +429,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
{
{
if
(
!
(
N
%
NPerBlock
==
0
))
if
(
!
(
N
%
NPerBlock
==
0
))
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"Arg N value is not a multiple of NPerBlock! N: "
<<
N
<<
" "
std
::
cout
<<
"Arg N value is not a multiple of NPerBlock! N: "
<<
N
<<
" "
<<
__FILE__
<<
":"
<<
__LINE__
<<
", in function: "
<<
__func__
<<
__FILE__
<<
":"
<<
__LINE__
<<
", in function: "
<<
__func__
...
@@ -444,7 +444,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -444,7 +444,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
{
{
if
(
!
(
K
%
KPerBlock
==
0
))
if
(
!
(
K
%
KPerBlock
==
0
))
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"Arg K value is not a multiple of ! KPerBlock: "
<<
K
<<
" "
std
::
cout
<<
"Arg K value is not a multiple of ! KPerBlock: "
<<
K
<<
" "
<<
__FILE__
<<
":"
<<
__LINE__
<<
", in function: "
<<
__func__
<<
__FILE__
<<
":"
<<
__LINE__
<<
", in function: "
<<
__func__
...
@@ -458,7 +458,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -458,7 +458,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
{
{
if
(
K
%
ABlockTransferSrcScalarPerVector
!=
0
)
if
(
K
%
ABlockTransferSrcScalarPerVector
!=
0
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"Arg K ("
<<
K
std
::
cout
<<
"Arg K ("
<<
K
<<
") value is not a multiple of ABlockTransferSrcScalarPerVector ("
<<
") value is not a multiple of ABlockTransferSrcScalarPerVector ("
...
@@ -472,7 +472,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -472,7 +472,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
{
{
if
(
M
%
ABlockTransferSrcScalarPerVector
!=
0
)
if
(
M
%
ABlockTransferSrcScalarPerVector
!=
0
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"Arg M ("
<<
M
std
::
cout
<<
"Arg M ("
<<
M
<<
") value is not a multiple of ABlockTransferSrcScalarPerVector ("
<<
") value is not a multiple of ABlockTransferSrcScalarPerVector ("
...
@@ -487,7 +487,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -487,7 +487,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
{
{
if
(
N
%
BBlockTransferSrcScalarPerVector
!=
0
)
if
(
N
%
BBlockTransferSrcScalarPerVector
!=
0
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"Arg N ("
<<
N
std
::
cout
<<
"Arg N ("
<<
N
<<
") value is not a multiple of BBlockTransferSrcScalarPerVector ("
<<
") value is not a multiple of BBlockTransferSrcScalarPerVector ("
...
@@ -501,7 +501,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -501,7 +501,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
{
{
if
(
K
%
BBlockTransferSrcScalarPerVector
!=
0
)
if
(
K
%
BBlockTransferSrcScalarPerVector
!=
0
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"Arg K ("
<<
K
std
::
cout
<<
"Arg K ("
<<
K
<<
") value is not a multiple of BBlockTransferSrcScalarPerVector ("
<<
") value is not a multiple of BBlockTransferSrcScalarPerVector ("
...
@@ -516,7 +516,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -516,7 +516,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
{
{
if
(
N
%
CDEShuffleBlockTransferScalarPerVector_NPerBlock
!=
0
)
if
(
N
%
CDEShuffleBlockTransferScalarPerVector_NPerBlock
!=
0
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"Arg N ("
<<
N
std
::
cout
<<
"Arg N ("
<<
N
<<
") value is not a multiple of "
<<
") value is not a multiple of "
...
@@ -532,7 +532,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -532,7 +532,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
{
{
if
(
M
%
CDEShuffleBlockTransferScalarPerVector_NPerBlock
!=
0
)
if
(
M
%
CDEShuffleBlockTransferScalarPerVector_NPerBlock
!=
0
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"Arg M ("
<<
M
std
::
cout
<<
"Arg M ("
<<
M
<<
") value is not a multiple of "
<<
") value is not a multiple of "
...
@@ -550,7 +550,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -550,7 +550,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
if
(
k_batch_size
<
KPerBlock
)
if
(
k_batch_size
<
KPerBlock
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"The k-batch size ("
<<
k_batch_size
std
::
cout
<<
"The k-batch size ("
<<
k_batch_size
<<
") value is less than KPerBlock!
\n
"
<<
") value is less than KPerBlock!
\n
"
...
@@ -562,7 +562,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -562,7 +562,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
if
(
k_batch_size
%
KPerBlock
!=
0
)
if
(
k_batch_size
%
KPerBlock
!=
0
)
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"The k-batch size ("
<<
k_batch_size
std
::
cout
<<
"The k-batch size ("
<<
k_batch_size
<<
") value is not a multiple of KPerBlock!
\n
"
<<
") value is not a multiple of KPerBlock!
\n
"
...
@@ -582,7 +582,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
...
@@ -582,7 +582,7 @@ class GridwiseGemmMultipleD_xdl_splitk_cshuffle_v2
if
(
!
GridwiseGemmPipe
::
IsSupported
(
num_k_loop
))
if
(
!
GridwiseGemmPipe
::
IsSupported
(
num_k_loop
))
{
{
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"The number of k loops ("
<<
num_k_loop
std
::
cout
<<
"The number of k loops ("
<<
num_k_loop
<<
") value is not supported by GridwiseGemm Pipeline."
<<
") value is not supported by GridwiseGemm Pipeline."
...
...
profiler/include/profiler/profile_grouped_gemm_multiple_d_splitk_impl.hpp
View file @
6ebaa81e
...
@@ -84,7 +84,7 @@ bool profile_ggemm_multid_splitk(int do_verification,
...
@@ -84,7 +84,7 @@ bool profile_ggemm_multid_splitk(int do_verification,
c_m_n_host_results
.
push_back
(
c_m_n_host_results
.
push_back
(
Tensor
<
CDataType
>
(
f_host_tensor_descriptor
(
Ms
[
i
],
Ns
[
i
],
StrideCs
[
i
],
CLayout
{})));
Tensor
<
CDataType
>
(
f_host_tensor_descriptor
(
Ms
[
i
],
Ns
[
i
],
StrideCs
[
i
],
CLayout
{})));
if
(
ck
::
EnvIsEnabled
(
ENV
(
CK_LOGGING
)))
if
(
ck
::
EnvIsEnabled
(
CK_
ENV
(
CK_LOGGING
)))
{
{
std
::
cout
<<
"group: "
<<
i
<<
" a_m_k["
<<
i
<<
"]:"
<<
a_m_k
[
i
].
mDesc
<<
", b_k_n["
std
::
cout
<<
"group: "
<<
i
<<
" a_m_k["
<<
i
<<
"]:"
<<
a_m_k
[
i
].
mDesc
<<
", b_k_n["
<<
i
<<
"]:"
<<
b_k_n
[
i
].
mDesc
<<
", c_m_n_device_results["
<<
i
<<
i
<<
"]:"
<<
b_k_n
[
i
].
mDesc
<<
", c_m_n_device_results["
<<
i
...
...
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