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
c6e072a6
Commit
c6e072a6
authored
Apr 15, 2021
by
Jing Zhang
Browse files
finished vec output
parent
c1159e3c
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
2 additions
and
15 deletions
+2
-15
composable_kernel/include/driver/driver_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw_outpad.hpp
...tion_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw_outpad.hpp
+0
-6
composable_kernel/include/tensor_operation/gridwise_dynamic_gemm_v2.hpp
...nel/include/tensor_operation/gridwise_dynamic_gemm_v2.hpp
+2
-5
driver/src/conv_driver.cpp
driver/src/conv_driver.cpp
+0
-4
No files found.
composable_kernel/include/driver/driver_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw_outpad.hpp
View file @
c6e072a6
...
@@ -31,7 +31,6 @@ struct DriverDynamicConvolutionForwardImplicitGemm_v5r1_nchw_kcyx_nkhw_outpad
...
@@ -31,7 +31,6 @@ struct DriverDynamicConvolutionForwardImplicitGemm_v5r1_nchw_kcyx_nkhw_outpad
{
{
template
<
typename
...
Wei
,
template
<
typename
...
Wei
,
typename
...
In
,
typename
...
In
,
typename
...
Add
,
typename
...
Out
,
typename
...
Out
,
typename
ConvStrides
,
typename
ConvStrides
,
typename
ConvDilations
,
typename
ConvDilations
,
...
@@ -80,9 +79,6 @@ struct DriverDynamicConvolutionForwardImplicitGemm_v5r1_nchw_kcyx_nkhw_outpad
...
@@ -80,9 +79,6 @@ struct DriverDynamicConvolutionForwardImplicitGemm_v5r1_nchw_kcyx_nkhw_outpad
const
auto
OutRightPadH
=
Hop
-
Ho
;
const
auto
OutRightPadH
=
Hop
-
Ho
;
const
auto
OutRightPadW
=
Wop
-
Wo
;
const
auto
OutRightPadW
=
Wop
-
Wo
;
const
auto
AddRightPadH
=
2
*
OutRightPadH
;
const
auto
AddRightPadW
=
2
*
OutRightPadW
;
const
auto
InLeftPadH
=
in_left_pads
[
I0
];
const
auto
InLeftPadH
=
in_left_pads
[
I0
];
const
auto
InLeftPadW
=
in_left_pads
[
I1
];
const
auto
InLeftPadW
=
in_left_pads
[
I1
];
...
@@ -93,8 +89,6 @@ struct DriverDynamicConvolutionForwardImplicitGemm_v5r1_nchw_kcyx_nkhw_outpad
...
@@ -93,8 +89,6 @@ struct DriverDynamicConvolutionForwardImplicitGemm_v5r1_nchw_kcyx_nkhw_outpad
<<
std
::
endl
;
<<
std
::
endl
;
std
::
cerr
<<
"InRightPadH = "
<<
InRightPadH
<<
" InRightPadW = "
<<
InRightPadW
std
::
cerr
<<
"InRightPadH = "
<<
InRightPadH
<<
" InRightPadW = "
<<
InRightPadW
<<
std
::
endl
;
<<
std
::
endl
;
std
::
cerr
<<
"AddRightPadH = "
<<
AddRightPadH
<<
" AddRightPadW = "
<<
AddRightPadW
<<
std
::
endl
;
// weight tensor
// weight tensor
const
auto
wei_e_k_global_desc
=
transform_dynamic_tensor_descriptor
(
const
auto
wei_e_k_global_desc
=
transform_dynamic_tensor_descriptor
(
...
...
composable_kernel/include/tensor_operation/gridwise_dynamic_gemm_v2.hpp
View file @
c6e072a6
...
@@ -386,11 +386,8 @@ struct GridwiseDynamicGemm_km_kn_mn_v2
...
@@ -386,11 +386,8 @@ struct GridwiseDynamicGemm_km_kn_mn_v2
static_for
<
0
,
CThreadTransferDstScalarPerVector
,
1
>
{}([
&
](
auto
i
)
{
static_for
<
0
,
CThreadTransferDstScalarPerVector
,
1
>
{}([
&
](
auto
i
)
{
t
.
template
AsType
<
int8_t
>()(
i
)
=
t
.
template
AsType
<
int8_t
>()(
i
)
=
p_c_thread
[
c_k_n_ho_wo_thread_desc_vec
.
CalculateOffset
(
p_c_thread
[
c_k_n_ho_wo_thread_desc_vec
.
CalculateOffset
(
make_tuple
(
make_tuple
(
k_i
*
CThreadTransferDstScalarPerVector
+
i
,
k_i
*
CThreadTransferDstScalarPerVector
+
i
,
0
,
h_i
,
w_i
))];
0
,
h_i
/
2
,
w_i
/
2
))];
});
});
// d_vec.template AsType<FloatC>()(
// d_vec.template AsType<FloatC>()(
...
...
driver/src/conv_driver.cpp
View file @
c6e072a6
...
@@ -625,12 +625,9 @@ int main(int argc, char* argv[])
...
@@ -625,12 +625,9 @@ int main(int argc, char* argv[])
constexpr
auto
Ho
=
out_nkhw_desc
.
GetLength
(
Number
<
2
>
{});
constexpr
auto
Ho
=
out_nkhw_desc
.
GetLength
(
Number
<
2
>
{});
constexpr
auto
Wo
=
out_nkhw_desc
.
GetLength
(
Number
<
3
>
{});
constexpr
auto
Wo
=
out_nkhw_desc
.
GetLength
(
Number
<
3
>
{});
// auto add_nkhw_desc = make_native_tensor_descriptor_packed(Sequence<N, K, Ho * 2, Wo * 2>{});
ostream_tensor_descriptor
(
in_nchw_desc
,
std
::
cout
<<
"in_nchw_desc: "
);
ostream_tensor_descriptor
(
in_nchw_desc
,
std
::
cout
<<
"in_nchw_desc: "
);
ostream_tensor_descriptor
(
wei_kcyx_desc
,
std
::
cout
<<
"wei_kcyx_desc: "
);
ostream_tensor_descriptor
(
wei_kcyx_desc
,
std
::
cout
<<
"wei_kcyx_desc: "
);
ostream_tensor_descriptor
(
out_nkhw_desc
,
std
::
cout
<<
"out_nkhw_desc: "
);
ostream_tensor_descriptor
(
out_nkhw_desc
,
std
::
cout
<<
"out_nkhw_desc: "
);
// ostream_tensor_descriptor(add_nkhw_desc, std::cout << "add_nkhw_desc: ");
print_array
(
"LeftPads"
,
to_multi_index
(
LeftPads
{}));
print_array
(
"LeftPads"
,
to_multi_index
(
LeftPads
{}));
print_array
(
"RightPads"
,
to_multi_index
(
RightPads
{}));
print_array
(
"RightPads"
,
to_multi_index
(
RightPads
{}));
...
@@ -661,7 +658,6 @@ int main(int argc, char* argv[])
...
@@ -661,7 +658,6 @@ int main(int argc, char* argv[])
Tensor
<
in_data_t
>
in_nchw
(
make_HostTensorDescriptor
(
in_nchw_desc
));
Tensor
<
in_data_t
>
in_nchw
(
make_HostTensorDescriptor
(
in_nchw_desc
));
Tensor
<
in_data_t
>
wei_kcyx
(
make_HostTensorDescriptor
(
wei_kcyx_desc
));
Tensor
<
in_data_t
>
wei_kcyx
(
make_HostTensorDescriptor
(
wei_kcyx_desc
));
Tensor
<
out_data_t
>
out_nkhw_host
(
make_HostTensorDescriptor
(
out_nkhw_desc
));
Tensor
<
out_data_t
>
out_nkhw_host
(
make_HostTensorDescriptor
(
out_nkhw_desc
));
Tensor
<
out_data_t
>
out_nkhw_device
(
make_HostTensorDescriptor
(
out_nkhw_desc
));
Tensor
<
out_data_t
>
out_nkhw_device
(
make_HostTensorDescriptor
(
out_nkhw_desc
));
...
...
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