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
4c949e20
Commit
4c949e20
authored
Nov 13, 2024
by
linsun12
Browse files
throw debug msg so those show inside MIOpen
parent
d4a1bcc1
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
37 additions
and
0 deletions
+37
-0
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
...mpl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
+30
-0
include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_abd_xdl_cshuffle.hpp
...tion/gpu/grid/gridwise_gemm_multiple_abd_xdl_cshuffle.hpp
+6
-0
profiler/include/profiler/profile_grouped_conv_fwd_impl.hpp
profiler/include/profiler/profile_grouped_conv_fwd_impl.hpp
+1
-0
No files found.
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
View file @
4c949e20
...
@@ -1101,14 +1101,17 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1101,14 +1101,17 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
// check device
// check device
if
(
get_device_name
()
==
"gfx908"
)
if
(
get_device_name
()
==
"gfx908"
)
{
{
printf
(
"device is gfx908
\n
"
);
// FIXME: re-enable fp64 when SWDEV-335738 is fixed
// FIXME: re-enable fp64 when SWDEV-335738 is fixed
if
constexpr
(
!
(
is_same_v
<
AccDataType
,
float
>
||
is_same_v
<
AccDataType
,
int32_t
>
))
if
constexpr
(
!
(
is_same_v
<
AccDataType
,
float
>
||
is_same_v
<
AccDataType
,
int32_t
>
))
{
{
printf
(
"accDataType is wrong
\n
"
);
return
false
;
return
false
;
}
}
}
}
if
(
!
ck
::
is_xdl_supported
())
if
(
!
ck
::
is_xdl_supported
())
{
{
printf
(
"not xdl supported
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -1116,6 +1119,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1116,6 +1119,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
constexpr
(
ConvForwardSpecialization
==
if
constexpr
(
ConvForwardSpecialization
==
ConvolutionForwardSpecialization
::
Filter1x1Stride1Pad0
)
ConvolutionForwardSpecialization
::
Filter1x1Stride1Pad0
)
{
{
printf
(
"convforwardspecialization is Filter1x1Stride1Pad0
\n
"
);
// check if it's 1x1, stride=1 conv
// check if it's 1x1, stride=1 conv
for
(
index_t
i
=
0
;
i
<
NDimSpatial
;
++
i
)
for
(
index_t
i
=
0
;
i
<
NDimSpatial
;
++
i
)
{
{
...
@@ -1126,6 +1130,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1126,6 +1130,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
!
(
X
==
1
&&
ConvStride
==
1
&&
LeftPad
==
0
&&
RightPad
==
0
))
if
(
!
(
X
==
1
&&
ConvStride
==
1
&&
LeftPad
==
0
&&
RightPad
==
0
))
{
{
printf
(
"argument doesn't support Filter1x1Stride1Pad0
\n
"
);
return
false
;
return
false
;
}
}
}
}
...
@@ -1133,6 +1138,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1133,6 +1138,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
else
if
constexpr
(
ConvForwardSpecialization
==
else
if
constexpr
(
ConvForwardSpecialization
==
ConvolutionForwardSpecialization
::
Filter1x1Pad0
)
ConvolutionForwardSpecialization
::
Filter1x1Pad0
)
{
{
printf
(
"convforwardspecialization is Filter1x1Pad0
\n
"
);
// check if it's 1x1 conv
// check if it's 1x1 conv
for
(
index_t
i
=
0
;
i
<
NDimSpatial
;
++
i
)
for
(
index_t
i
=
0
;
i
<
NDimSpatial
;
++
i
)
{
{
...
@@ -1142,14 +1148,17 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1142,14 +1148,17 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
!
(
X
==
1
&&
LeftPad
==
0
&&
RightPad
==
0
))
if
(
!
(
X
==
1
&&
LeftPad
==
0
&&
RightPad
==
0
))
{
{
printf
(
"argument doesn't support Filter1x1Pad0
\n
"
);
return
false
;
return
false
;
}
}
}
}
}
}
else
if
constexpr
(
ConvForwardSpecialization
==
ConvolutionForwardSpecialization
::
Filter3x3
)
else
if
constexpr
(
ConvForwardSpecialization
==
ConvolutionForwardSpecialization
::
Filter3x3
)
{
{
printf
(
"convforwardspecialization is Filter3x3
\n
"
);
if
(
C
!=
1
)
if
(
C
!=
1
)
{
{
printf
(
"channel != 1
\n
"
);
return
false
;
return
false
;
}
}
for
(
index_t
i
=
0
;
i
<
NDimSpatial
;
++
i
)
for
(
index_t
i
=
0
;
i
<
NDimSpatial
;
++
i
)
...
@@ -1158,28 +1167,34 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1158,28 +1167,34 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
filter_spatial_dim
!=
I3
)
if
(
filter_spatial_dim
!=
I3
)
{
{
printf
(
"filter spatial dim != 3
\n
"
);
return
false
;
return
false
;
}
}
}
}
if
constexpr
(
!
is_NSpatialGC_GKSpatial_NSpatialGK
<
ALayout
,
BLayout
,
ELayout
>
())
if
constexpr
(
!
is_NSpatialGC_GKSpatial_NSpatialGK
<
ALayout
,
BLayout
,
ELayout
>
())
{
{
printf
(
"!is_NSpatialGC_GKSpatial_NSpatialGK<ALayout, BLayout, ELayout>
\n
"
);
return
false
;
return
false
;
}
}
}
}
if
constexpr
(
NumGroupsToMerge
>
1
)
if
constexpr
(
NumGroupsToMerge
>
1
)
{
{
printf
(
"number of groups to merge is > 1
\n
"
);
if
(
!
(
C
==
1
))
if
(
!
(
C
==
1
))
{
{
printf
(
"channel != 1
\n
"
);
return
false
;
return
false
;
}
}
if
(
G
%
NumGroupsToMerge
!=
0
)
if
(
G
%
NumGroupsToMerge
!=
0
)
{
{
printf
(
"number of groups is wrong
\n
"
);
return
false
;
return
false
;
}
}
if
constexpr
(
!
(
is_NSpatialGC_GKSpatial_NSpatialGK
<
ALayout
,
BLayout
,
ELayout
>
()
||
if
constexpr
(
!
(
is_NSpatialGC_GKSpatial_NSpatialGK
<
ALayout
,
BLayout
,
ELayout
>
()
||
is_NGCSpatial_GKSpatial_NGKSpatial
<
ALayout
,
BLayout
,
ELayout
>
()))
is_NGCSpatial_GKSpatial_NGKSpatial
<
ALayout
,
BLayout
,
ELayout
>
()))
{
{
printf
(
"layout is wrong
\n
"
);
return
false
;
return
false
;
}
}
}
}
...
@@ -1202,12 +1217,14 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1202,12 +1217,14 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
is_NGCSpatial_GKSpatial_NGKSpatial
<
ALayout
,
BLayout
,
ELayout
>
())
&&
is_NGCSpatial_GKSpatial_NGKSpatial
<
ALayout
,
BLayout
,
ELayout
>
())
&&
G
%
ABlockTransferSrcScalarPerVector
==
0
))
G
%
ABlockTransferSrcScalarPerVector
==
0
))
{
{
printf
(
"vector access of A is wrong 0
\n
"
);
return
false
;
return
false
;
}
}
}
}
}
}
else
else
{
{
printf
(
"vector access of B is wrong 1
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -1222,11 +1239,13 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1222,11 +1239,13 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
{
{
if
(
!
(
BBlockTransferSrcVectorDim
==
2
&&
C
%
BBlockTransferSrcScalarPerVector
==
0
))
if
(
!
(
BBlockTransferSrcVectorDim
==
2
&&
C
%
BBlockTransferSrcScalarPerVector
==
0
))
{
{
printf
(
"vector access of B is wrong 0
\n
"
);
return
false
;
return
false
;
}
}
}
}
else
else
{
{
printf
(
"vector access of B is wrong 1
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -1245,6 +1264,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1245,6 +1264,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
{
{
if
(
!
(
K
%
CDEBlockTransferScalarPerVector_NPerBlock
==
0
))
if
(
!
(
K
%
CDEBlockTransferScalarPerVector_NPerBlock
==
0
))
{
{
printf
(
"vector access of CDE is wrong 0
\n
"
);
valid
=
false
;
valid
=
false
;
}
}
...
@@ -1254,6 +1274,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1254,6 +1274,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
arg
.
ds_g_n_k_wos_lengths_
[
i
][
0
]
!=
arg
.
e_g_n_k_wos_lengths_
[
0
]
||
if
(
arg
.
ds_g_n_k_wos_lengths_
[
i
][
0
]
!=
arg
.
e_g_n_k_wos_lengths_
[
0
]
||
arg
.
ds_g_n_k_wos_lengths_
[
i
][
2
]
!=
arg
.
e_g_n_k_wos_lengths_
[
2
])
arg
.
ds_g_n_k_wos_lengths_
[
i
][
2
]
!=
arg
.
e_g_n_k_wos_lengths_
[
2
])
{
{
printf
(
"G and K must be the same shape
\n
"
);
valid
=
false
;
valid
=
false
;
}
}
}
}
...
@@ -1264,6 +1285,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1264,6 +1285,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
{
{
if
(
arg
.
ds_g_n_k_wos_lengths_
[
i
][
d
]
!=
arg
.
e_g_n_k_wos_lengths_
[
d
])
if
(
arg
.
ds_g_n_k_wos_lengths_
[
i
][
d
]
!=
arg
.
e_g_n_k_wos_lengths_
[
d
])
{
{
printf
(
"E and D must be the same shape
\n
"
);
valid
=
false
;
valid
=
false
;
}
}
}
}
...
@@ -1271,6 +1293,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1271,6 +1293,7 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
}
}
else
else
{
{
printf
(
"CDE layout is wrong
\n
"
);
valid
=
false
;
valid
=
false
;
}
}
});
});
...
@@ -1280,11 +1303,13 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1280,11 +1303,13 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
{
{
if
((
G
*
C
)
%
CDEBlockTransferScalarPerVector_NPerBlock
!=
0
)
if
((
G
*
C
)
%
CDEBlockTransferScalarPerVector_NPerBlock
!=
0
)
{
{
printf
(
"vectorization is wrong 0
\n
"
);
return
false
;
return
false
;
}
}
if
((
G
*
K
)
%
CDEBlockTransferScalarPerVector_NPerBlock
!=
0
)
if
((
G
*
K
)
%
CDEBlockTransferScalarPerVector_NPerBlock
!=
0
)
{
{
printf
(
"vectorization is wrong 1
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -1295,17 +1320,20 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1295,17 +1320,20 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
input_spatial_acum
%
CDEBlockTransferScalarPerVector_NPerBlock
!=
0
)
if
(
input_spatial_acum
%
CDEBlockTransferScalarPerVector_NPerBlock
!=
0
)
{
{
printf
(
"vectorization is wrong 2
\n
"
);
return
false
;
return
false
;
}
}
if
(
output_spatial_acum
%
CDEBlockTransferScalarPerVector_NPerBlock
!=
0
)
if
(
output_spatial_acum
%
CDEBlockTransferScalarPerVector_NPerBlock
!=
0
)
{
{
printf
(
"vectorization is wrong 3
\n
"
);
return
false
;
return
false
;
}
}
}
}
if
(
!
valid
)
if
(
!
valid
)
{
{
printf
(
"layout is wrong
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -1319,11 +1347,13 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -1319,11 +1347,13 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
{
{
if
(
!
(
K
%
CDEBlockTransferScalarPerVector_NPerBlock
==
0
))
if
(
!
(
K
%
CDEBlockTransferScalarPerVector_NPerBlock
==
0
))
{
{
printf
(
"vector access of E is wrong
\n
"
);
return
false
;
return
false
;
}
}
}
}
else
else
{
{
printf
(
"layout is wrong
\n
"
);
return
false
;
return
false
;
}
}
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_abd_xdl_cshuffle.hpp
View file @
4c949e20
...
@@ -324,6 +324,7 @@ struct GridwiseGemmMultipleABD_xdl_cshuffle
...
@@ -324,6 +324,7 @@ struct GridwiseGemmMultipleABD_xdl_cshuffle
// check consistency of desc
// check consistency of desc
if
(
!
(
M
==
e_grid_desc_m_n
.
GetLength
(
I0
)
&&
N
==
e_grid_desc_m_n
.
GetLength
(
I1
)
&&
AK
==
BK
))
if
(
!
(
M
==
e_grid_desc_m_n
.
GetLength
(
I0
)
&&
N
==
e_grid_desc_m_n
.
GetLength
(
I1
)
&&
AK
==
BK
))
{
{
printf
(
"consistency of desc is wrong
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -353,12 +354,14 @@ struct GridwiseGemmMultipleABD_xdl_cshuffle
...
@@ -353,12 +354,14 @@ struct GridwiseGemmMultipleABD_xdl_cshuffle
if
(
!
valid
)
if
(
!
valid
)
{
{
printf
(
"gridwise validity error 0
\n
"
);
return
false
;
return
false
;
}
}
// check tile size
// check tile size
if
(
!
(
M
%
MPerBlock
==
0
&&
N
%
NPerBlock
==
0
&&
AK
%
KPerBlock
==
0
))
if
(
!
(
M
%
MPerBlock
==
0
&&
N
%
NPerBlock
==
0
&&
AK
%
KPerBlock
==
0
))
{
{
printf
(
"gridwise validity error 1
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -367,12 +370,14 @@ struct GridwiseGemmMultipleABD_xdl_cshuffle
...
@@ -367,12 +370,14 @@ struct GridwiseGemmMultipleABD_xdl_cshuffle
if
(
!
GridwiseGemmPipe
::
IsSupported
(
num_k_loop
))
if
(
!
GridwiseGemmPipe
::
IsSupported
(
num_k_loop
))
{
{
printf
(
"gridwise validity error 2
\n
"
);
return
false
;
return
false
;
}
}
// check block-to-E-tile
// check block-to-E-tile
if
(
!
block_2_etile_map
.
CheckValidity
(
e_grid_desc_m_n
))
if
(
!
block_2_etile_map
.
CheckValidity
(
e_grid_desc_m_n
))
{
{
printf
(
"gridwise validity error 3
\n
"
);
return
false
;
return
false
;
}
}
...
@@ -381,6 +386,7 @@ struct GridwiseGemmMultipleABD_xdl_cshuffle
...
@@ -381,6 +386,7 @@ struct GridwiseGemmMultipleABD_xdl_cshuffle
if
(
!
(
e_grid_desc_m_n
.
GetElementSpaceSize
()
*
sizeof
(
EDataType
)
<=
TwoGB
))
if
(
!
(
e_grid_desc_m_n
.
GetElementSpaceSize
()
*
sizeof
(
EDataType
)
<=
TwoGB
))
{
{
printf
(
"gridwise validity error 4
\n
"
);
return
false
;
return
false
;
}
}
...
...
profiler/include/profiler/profile_grouped_conv_fwd_impl.hpp
View file @
4c949e20
...
@@ -159,6 +159,7 @@ bool profile_grouped_conv_fwd_impl(int do_verification,
...
@@ -159,6 +159,7 @@ bool profile_grouped_conv_fwd_impl(int do_verification,
out_device_buf
.
SetZero
();
out_device_buf
.
SetZero
();
std
::
string
op_name
=
op_ptr
->
GetTypeString
();
std
::
string
op_name
=
op_ptr
->
GetTypeString
();
std
::
cout
<<
op_name
<<
" supports this argument!!!!!!!!!!!!!!"
<<
std
::
endl
;
auto
invoker_ptr
=
op_ptr
->
MakeInvokerPointer
();
auto
invoker_ptr
=
op_ptr
->
MakeInvokerPointer
();
...
...
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