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
68aa4da9
Commit
68aa4da9
authored
Nov 13, 2024
by
Astha
Browse files
added some debug prints to pinpoint issues in MIOpen
parent
83074f4c
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
22 additions
and
0 deletions
+22
-0
include/ck/tensor_operation/gpu/device/impl/codegen_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
...gen_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
+22
-0
No files found.
include/ck/tensor_operation/gpu/device/impl/codegen_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
View file @
68aa4da9
...
@@ -9,6 +9,7 @@
...
@@ -9,6 +9,7 @@
#include <iterator>
#include <iterator>
#include <numeric>
#include <numeric>
#include <sstream>
#include <sstream>
#include <stdio.h>
#include "ck/host_utility/device_prop.hpp"
#include "ck/host_utility/device_prop.hpp"
#include "ck/host_utility/kernel_launch.hpp"
#include "ck/host_utility/kernel_launch.hpp"
...
@@ -262,8 +263,14 @@ __global__ void
...
@@ -262,8 +263,14 @@ __global__ void
}
// namespace
}
// namespace
#ifdef CK_CODE_GEN_RTC
template
<
typename
T
>
template
<
typename
T
>
using
is_tuple
=
decltype
(
ck
::
declval
<
T
&>
().
IsTuple
());
using
is_tuple
=
decltype
(
ck
::
declval
<
T
&>
().
IsTuple
());
#else
template
<
typename
T
>
using
is_tuple
=
decltype
(
std
::
declval
<
T
&>
().
IsTuple
());
#endif
//
//
// @brief Device Convolution operation.
// @brief Device Convolution operation.
...
@@ -746,6 +753,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -746,6 +753,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
!
(
X
==
1
&&
ConvStride
==
1
&&
LeftPad
==
0
&&
RightPad
==
0
))
if
(
!
(
X
==
1
&&
ConvStride
==
1
&&
LeftPad
==
0
&&
RightPad
==
0
))
{
{
printf
(
"Check 1"
);
return
false
;
return
false
;
}
}
}
}
...
@@ -762,6 +770,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -762,6 +770,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
!
(
X
==
1
&&
LeftPad
==
0
&&
RightPad
==
0
))
if
(
!
(
X
==
1
&&
LeftPad
==
0
&&
RightPad
==
0
))
{
{
printf
(
"Check 2"
);
return
false
;
return
false
;
}
}
}
}
...
@@ -779,11 +788,13 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -779,11 +788,13 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
!
(
ABlockTransferSrcVectorDim
==
2
&&
C
%
ABlockTransferSrcScalarPerVector
==
0
))
if
(
!
(
ABlockTransferSrcVectorDim
==
2
&&
C
%
ABlockTransferSrcScalarPerVector
==
0
))
{
{
printf
(
"Check 3"
);
return
false
;
return
false
;
}
}
}
}
else
else
{
{
printf
(
"Check 4"
);
return
false
;
return
false
;
}
}
...
@@ -800,11 +811,13 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -800,11 +811,13 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
!
(
BBlockTransferSrcVectorDim
==
2
&&
C
%
BBlockTransferSrcScalarPerVector
==
0
))
if
(
!
(
BBlockTransferSrcVectorDim
==
2
&&
C
%
BBlockTransferSrcScalarPerVector
==
0
))
{
{
printf
(
"Check 5"
);
return
false
;
return
false
;
}
}
}
}
else
else
{
{
printf
(
"Check 6"
);
return
false
;
return
false
;
}
}
...
@@ -825,6 +838,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -825,6 +838,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
!
(
K
%
CDEBlockTransferScalarPerVector_NPerBlock
==
0
))
if
(
!
(
K
%
CDEBlockTransferScalarPerVector_NPerBlock
==
0
))
{
{
printf
(
"Check 7"
);
valid
=
false
;
valid
=
false
;
}
}
...
@@ -834,6 +848,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -834,6 +848,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_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
(
"Check 8"
);
valid
=
false
;
valid
=
false
;
}
}
}
}
...
@@ -844,6 +859,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -844,6 +859,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_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
(
"Check 9"
);
valid
=
false
;
valid
=
false
;
}
}
}
}
...
@@ -851,12 +867,14 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -851,12 +867,14 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
}
}
else
else
{
{
printf
(
"Check 10"
);
valid
=
false
;
valid
=
false
;
}
}
});
});
if
(
!
valid
)
if
(
!
valid
)
{
{
printf
(
"Check 11"
);
return
false
;
return
false
;
}
}
...
@@ -871,11 +889,13 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -871,11 +889,13 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if
(
!
(
K
%
CDEBlockTransferScalarPerVector_NPerBlock
==
0
))
if
(
!
(
K
%
CDEBlockTransferScalarPerVector_NPerBlock
==
0
))
{
{
printf
(
"Check 12"
);
return
false
;
return
false
;
}
}
}
}
else
else
{
{
printf
(
"Check 13"
);
return
false
;
return
false
;
}
}
...
@@ -887,6 +907,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -887,6 +907,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
generate_tuple
([
&
](
auto
)
{
return
arg
.
a_grid_desc_m_k_
;
},
Number
<
NumATensor
>
{});
generate_tuple
([
&
](
auto
)
{
return
arg
.
a_grid_desc_m_k_
;
},
Number
<
NumATensor
>
{});
const
auto
bs_grid_desc_bk0_n_bk1
=
const
auto
bs_grid_desc_bk0_n_bk1
=
generate_tuple
([
&
](
auto
)
{
return
arg
.
b_grid_desc_n_k_
;
},
Number
<
NumBTensor
>
{});
generate_tuple
([
&
](
auto
)
{
return
arg
.
b_grid_desc_n_k_
;
},
Number
<
NumBTensor
>
{});
printf
(
"Check 14"
);
return
GridwiseGemm
::
CheckValidity
(
as_grid_desc_ak0_m_ak1
,
return
GridwiseGemm
::
CheckValidity
(
as_grid_desc_ak0_m_ak1
,
bs_grid_desc_bk0_n_bk1
,
bs_grid_desc_bk0_n_bk1
,
arg
.
ds_grid_desc_m_n_
,
arg
.
ds_grid_desc_m_n_
,
...
@@ -895,6 +916,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
...
@@ -895,6 +916,7 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
}
}
else
else
{
{
printf
(
"Check 15"
);
return
GridwiseGemm
::
CheckValidity
(
arg
.
a_grid_desc_m_k_
,
return
GridwiseGemm
::
CheckValidity
(
arg
.
a_grid_desc_m_k_
,
arg
.
b_grid_desc_n_k_
,
arg
.
b_grid_desc_n_k_
,
arg
.
ds_grid_desc_m_n_
,
arg
.
ds_grid_desc_m_n_
,
...
...
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