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
a7a1e3c1
Commit
a7a1e3c1
authored
Nov 20, 2019
by
Chao Liu
Browse files
add bwd-data v4r5
parent
2f0f26d3
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
26 additions
and
19 deletions
+26
-19
driver/include/device_convolution_backward_data_implicit_gemm_v4r5_nchw_kcyx_nkhw.hpp
...ution_backward_data_implicit_gemm_v4r5_nchw_kcyx_nkhw.hpp
+8
-2
driver/src/conv_bwd_data_driver.cpp
driver/src/conv_bwd_data_driver.cpp
+18
-17
No files found.
driver/include/device_convolution_backward_data_implicit_gemm_v4r5_nchw_kcyx_nkhw.hpp
View file @
a7a1e3c1
...
...
@@ -78,8 +78,14 @@ void device_convolution_backward_data_implicit_gemm_v4r5_nchw_kcyx_nkhw(InDesc i
constexpr
index_t
InThreadCopyDstDataPerWrite_B
=
1
;
#endif
constexpr
index_t
E
=
C
*
Y
*
X
;
constexpr
index_t
B
=
(
N
*
Ho
*
Wo
);
constexpr
index_t
C0
=
GemmMPerThreadSubC
;
constexpr
index_t
N0
=
GemmNPerThreadSubC
;
constexpr
index_t
C1
=
C
/
C0
;
constexpr
index_t
N1
=
N
/
N0
;
constexpr
index_t
E
=
C1
*
Y
*
X
;
constexpr
index_t
B
=
(
N1
*
Ho
*
Wo
);
constexpr
index_t
GridSize
=
((
E
+
EPerBlock
-
1
)
/
EPerBlock
)
*
((
B
+
BPerBlock
-
1
)
/
BPerBlock
);
...
...
driver/src/conv_bwd_data_driver.cpp
View file @
a7a1e3c1
...
...
@@ -21,19 +21,19 @@ int main(int argc, char* argv[])
using
namespace
ck
;
#if 0
constexpr index_t N =
2
;
constexpr index_t C =
8
;
constexpr index_t HI =
8
;
constexpr index_t WI =
8
;
constexpr index_t K =
128
;
constexpr index_t Y =
4
;
constexpr index_t X =
4
;
constexpr index_t N =
128
;
constexpr index_t C =
256
;
constexpr index_t HI =
35
;
constexpr index_t WI =
35
;
constexpr index_t K =
384
;
constexpr index_t Y =
3
;
constexpr index_t X =
3
;
using ConvStrides = Sequence<
1
,
1
>;
using ConvStrides = Sequence<
2
,
2
>;
using ConvDilations = Sequence<1, 1>;
using LeftPads = Sequence<
1
,
1
>;
using RightPads = Sequence<
2
,
2
>;
using LeftPads = Sequence<
0
,
0
>;
using RightPads = Sequence<
0
,
0
>;
#elif
0
// 3x3, 34x34
constexpr
index_t
N
=
64
;
...
...
@@ -49,7 +49,7 @@ int main(int argc, char* argv[])
using
LeftPads
=
Sequence
<
0
,
0
>
;
using
RightPads
=
Sequence
<
0
,
0
>
;
#elif
0
#elif
1
// 1x1 filter, 8x8 image
// cudnn@V100 68%, ck@V100 72%, ck@P100 52%, ck@VII 42%
constexpr
index_t
N
=
64
;
...
...
@@ -97,7 +97,7 @@ int main(int argc, char* argv[])
using
LeftPads
=
Sequence
<
0
,
0
>
;
using
RightPads
=
Sequence
<
0
,
0
>
;
#elif
1
#elif
0
// 1x1 filter, 8x8 image
// cudnn@V100 83%, ck@V100 75%, ck@P100 78%, ck@VII 65%
constexpr
index_t
N
=
128
;
...
...
@@ -241,7 +241,7 @@ int main(int argc, char* argv[])
using
LeftPads
=
Sequence
<
0
,
0
>
;
using
RightPads
=
Sequence
<
0
,
0
>
;
#elif
0
#elif
1
// 3x3 filter, 2x2 stride, 35x35 input, 17x17 output
// cudnn@V100 90%, ck@V100 93%, ck@P100 83%, ck@VII 81%
constexpr
index_t
N
=
128
;
...
...
@@ -287,7 +287,7 @@ int main(int argc, char* argv[])
using
LeftPads
=
Sequence
<
3
,
0
>
;
using
RightPads
=
Sequence
<
3
,
0
>
;
#elif
1
#elif
0
// 1x7 filter, 0x3 pad, 17x17 input
constexpr
index_t
N
=
128
;
constexpr
index_t
C
=
128
;
...
...
@@ -375,9 +375,10 @@ int main(int argc, char* argv[])
check_error
(
in_nchw_host
,
in_nchw_device
);
#if 0
LogRange(std::cout << "col_eb : ", col_eb.mData, ",") << std::endl;
LogRange(std::cout << "img_nchw_host : ", img_nchw_host.mData, ",") << std::endl;
LogRange(std::cout << "img_nchw_device : ", img_nchw_device.mData, ",") << std::endl;
LogRange(std::cout << "out_nkhw : ", out_nkhw.mData, ",") << std::endl;
LogRange(std::cout << "wei_kcyx : ", wei_kcyx.mData, ",") << std::endl;
LogRange(std::cout << "in_nchw_host : ", in_nchw_host.mData, ",") << std::endl;
LogRange(std::cout << "in_nchw_device : ", in_nchw_device.mData, ",") << std::endl;
#endif
}
}
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