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
bec35fbc
Commit
bec35fbc
authored
Nov 21, 2019
by
Chao Liu
Browse files
rename
parent
72d5b799
Changes
5
Show whitespace changes
Inline
Side-by-side
Showing
5 changed files
with
16 additions
and
16 deletions
+16
-16
composable_kernel/include/kernel_algorithm/gridwise_convolution_backward_data_implicit_gemm_v1r1_nchw_kcyx_nkhw_lds_double_buffer.hpp
...a_implicit_gemm_v1r1_nchw_kcyx_nkhw_lds_double_buffer.hpp
+3
-3
composable_kernel/include/kernel_algorithm/gridwise_convolution_backward_data_implicit_gemm_v1r2_nchw_kcyx_nkhw_lds_double_buffer.hpp
...a_implicit_gemm_v1r2_nchw_kcyx_nkhw_lds_double_buffer.hpp
+3
-3
driver/include/device_convolution_backward_data_implicit_gemm_v1r1_nchw_kcyx_nkhw.hpp
...ution_backward_data_implicit_gemm_v1r1_nchw_kcyx_nkhw.hpp
+3
-3
driver/include/device_convolution_backward_data_implicit_gemm_v1r2_nchw_kcyx_nkhw.hpp
...ution_backward_data_implicit_gemm_v1r2_nchw_kcyx_nkhw.hpp
+3
-3
driver/src/conv_bwd_data_driver.cpp
driver/src/conv_bwd_data_driver.cpp
+4
-4
No files found.
composable_kernel/include/kernel_algorithm/gridwise_convolution_backward_data_implicit_gemm_v
4r4
_nchw_kcyx_nkhw_lds_double_buffer.hpp
→
composable_kernel/include/kernel_algorithm/gridwise_convolution_backward_data_implicit_gemm_v
1r1
_nchw_kcyx_nkhw_lds_double_buffer.hpp
View file @
bec35fbc
#ifndef CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V
4R4
_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
#ifndef CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V
1R1
_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
#define CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V
4R4
_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
#define CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V
1R1
_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
#include "common_header.hpp"
#include "common_header.hpp"
#include "tensor_descriptor.hpp"
#include "tensor_descriptor.hpp"
...
@@ -41,7 +41,7 @@ template <index_t GridSize,
...
@@ -41,7 +41,7 @@ template <index_t GridSize,
typename
WeiBlockCopyClusterLengths_K_E
,
typename
WeiBlockCopyClusterLengths_K_E
,
index_t
WeiBlockCopyDataPerAccess_E
,
index_t
WeiBlockCopyDataPerAccess_E
,
index_t
InThreadCopyDataPerAccess_B
>
index_t
InThreadCopyDataPerAccess_B
>
struct
GridwiseConvolutionBackwardDataImplicitGemm_v
4r4
_nchw_kcyx_nkhw_lds_double_buffer
struct
GridwiseConvolutionBackwardDataImplicitGemm_v
1r1
_nchw_kcyx_nkhw_lds_double_buffer
{
{
__device__
void
Run
(
Float
*
const
__restrict__
p_in_global
,
__device__
void
Run
(
Float
*
const
__restrict__
p_in_global
,
const
Float
*
const
__restrict__
p_wei_global
,
const
Float
*
const
__restrict__
p_wei_global
,
...
...
composable_kernel/include/kernel_algorithm/gridwise_convolution_backward_data_implicit_gemm_v
4r5
_nchw_kcyx_nkhw_lds_double_buffer.hpp
→
composable_kernel/include/kernel_algorithm/gridwise_convolution_backward_data_implicit_gemm_v
1r2
_nchw_kcyx_nkhw_lds_double_buffer.hpp
View file @
bec35fbc
#ifndef CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V
4R5
_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
#ifndef CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V
1R2
_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
#define CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V
4R5
_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
#define CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V
1R2
_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
#include "common_header.hpp"
#include "common_header.hpp"
#include "tensor_descriptor.hpp"
#include "tensor_descriptor.hpp"
...
@@ -43,7 +43,7 @@ template <index_t GridSize,
...
@@ -43,7 +43,7 @@ template <index_t GridSize,
index_t
WeiBlockCopySrcDataPerRead_E
,
index_t
WeiBlockCopySrcDataPerRead_E
,
index_t
WeiBlockCopyDstDataPerWrite_C0
,
index_t
WeiBlockCopyDstDataPerWrite_C0
,
index_t
InThreadCopyDstDataPerWrite_B
>
index_t
InThreadCopyDstDataPerWrite_B
>
struct
GridwiseConvolutionBackwardDataImplicitGemm_v
4r5
_nchw_kcyx_nkhw_lds_double_buffer
struct
GridwiseConvolutionBackwardDataImplicitGemm_v
1r2
_nchw_kcyx_nkhw_lds_double_buffer
{
{
__device__
void
Run
(
Float
*
const
__restrict__
p_in_global
,
__device__
void
Run
(
Float
*
const
__restrict__
p_in_global
,
const
Float
*
const
__restrict__
p_wei_global
,
const
Float
*
const
__restrict__
p_wei_global
,
...
...
driver/include/device_convolution_backward_data_implicit_gemm_v
4r4
_nchw_kcyx_nkhw.hpp
→
driver/include/device_convolution_backward_data_implicit_gemm_v
1r1
_nchw_kcyx_nkhw.hpp
View file @
bec35fbc
...
@@ -3,7 +3,7 @@
...
@@ -3,7 +3,7 @@
#include "device.hpp"
#include "device.hpp"
#include "tensor.hpp"
#include "tensor.hpp"
#include "gridwise_operation_wrapper.hpp"
#include "gridwise_operation_wrapper.hpp"
#include "gridwise_convolution_backward_data_implicit_gemm_v
4r4
_nchw_kcyx_nkhw_lds_double_buffer.hpp"
#include "gridwise_convolution_backward_data_implicit_gemm_v
1r1
_nchw_kcyx_nkhw_lds_double_buffer.hpp"
template
<
typename
T
,
template
<
typename
T
,
typename
InDesc
,
typename
InDesc
,
...
@@ -13,7 +13,7 @@ template <typename T,
...
@@ -13,7 +13,7 @@ template <typename T,
typename
ConvDilations
,
typename
ConvDilations
,
typename
LeftPads
,
typename
LeftPads
,
typename
RightPads
>
typename
RightPads
>
void
device_convolution_backward_data_implicit_gemm_v
4r4
_nchw_kcyx_nkhw
(
InDesc
in_nchw_desc
,
void
device_convolution_backward_data_implicit_gemm_v
1r1
_nchw_kcyx_nkhw
(
InDesc
in_nchw_desc
,
Tensor
<
T
>&
in_nchw
,
Tensor
<
T
>&
in_nchw
,
WeiDesc
wei_kcyx_desc
,
WeiDesc
wei_kcyx_desc
,
const
Tensor
<
T
>&
wei_kcyx
,
const
Tensor
<
T
>&
wei_kcyx
,
...
@@ -85,7 +85,7 @@ void device_convolution_backward_data_implicit_gemm_v4r4_nchw_kcyx_nkhw(InDesc i
...
@@ -85,7 +85,7 @@ void device_convolution_backward_data_implicit_gemm_v4r4_nchw_kcyx_nkhw(InDesc i
printf
(
"%s: BlockSize %u, GridSize %u
\n
"
,
__func__
,
BlockSize
,
GridSize
);
printf
(
"%s: BlockSize %u, GridSize %u
\n
"
,
__func__
,
BlockSize
,
GridSize
);
constexpr
auto
gridwise_conv
=
constexpr
auto
gridwise_conv
=
GridwiseConvolutionBackwardDataImplicitGemm_v
4r4
_nchw_kcyx_nkhw_lds_double_buffer
<
GridwiseConvolutionBackwardDataImplicitGemm_v
1r1
_nchw_kcyx_nkhw_lds_double_buffer
<
GridSize
,
GridSize
,
BlockSize
,
BlockSize
,
T
,
T
,
...
...
driver/include/device_convolution_backward_data_implicit_gemm_v
4r5
_nchw_kcyx_nkhw.hpp
→
driver/include/device_convolution_backward_data_implicit_gemm_v
1r2
_nchw_kcyx_nkhw.hpp
View file @
bec35fbc
...
@@ -3,7 +3,7 @@
...
@@ -3,7 +3,7 @@
#include "device.hpp"
#include "device.hpp"
#include "tensor.hpp"
#include "tensor.hpp"
#include "gridwise_operation_wrapper.hpp"
#include "gridwise_operation_wrapper.hpp"
#include "gridwise_convolution_backward_data_implicit_gemm_v
4r5
_nchw_kcyx_nkhw_lds_double_buffer.hpp"
#include "gridwise_convolution_backward_data_implicit_gemm_v
1r2
_nchw_kcyx_nkhw_lds_double_buffer.hpp"
template
<
typename
T
,
template
<
typename
T
,
typename
InDesc
,
typename
InDesc
,
...
@@ -13,7 +13,7 @@ template <typename T,
...
@@ -13,7 +13,7 @@ template <typename T,
typename
ConvDilations
,
typename
ConvDilations
,
typename
LeftPads
,
typename
LeftPads
,
typename
RightPads
>
typename
RightPads
>
void
device_convolution_backward_data_implicit_gemm_v
4r5
_nchw_kcyx_nkhw
(
InDesc
in_nchw_desc
,
void
device_convolution_backward_data_implicit_gemm_v
1r2
_nchw_kcyx_nkhw
(
InDesc
in_nchw_desc
,
Tensor
<
T
>&
in_nchw
,
Tensor
<
T
>&
in_nchw
,
WeiDesc
wei_kcyx_desc
,
WeiDesc
wei_kcyx_desc
,
const
Tensor
<
T
>&
wei_kcyx
,
const
Tensor
<
T
>&
wei_kcyx
,
...
@@ -93,7 +93,7 @@ void device_convolution_backward_data_implicit_gemm_v4r5_nchw_kcyx_nkhw(InDesc i
...
@@ -93,7 +93,7 @@ void device_convolution_backward_data_implicit_gemm_v4r5_nchw_kcyx_nkhw(InDesc i
printf
(
"%s: BlockSize %u, GridSize %u
\n
"
,
__func__
,
BlockSize
,
GridSize
);
printf
(
"%s: BlockSize %u, GridSize %u
\n
"
,
__func__
,
BlockSize
,
GridSize
);
constexpr
auto
gridwise_conv
=
constexpr
auto
gridwise_conv
=
GridwiseConvolutionBackwardDataImplicitGemm_v
4r5
_nchw_kcyx_nkhw_lds_double_buffer
<
GridwiseConvolutionBackwardDataImplicitGemm_v
1r2
_nchw_kcyx_nkhw_lds_double_buffer
<
GridSize
,
GridSize
,
BlockSize
,
BlockSize
,
T
,
T
,
...
...
driver/src/conv_bwd_data_driver.cpp
View file @
bec35fbc
...
@@ -13,8 +13,8 @@
...
@@ -13,8 +13,8 @@
#include "device_tensor.hpp"
#include "device_tensor.hpp"
#include "conv_common.hpp"
#include "conv_common.hpp"
#include "host_conv_bwd_data.hpp"
#include "host_conv_bwd_data.hpp"
#include "device_convolution_backward_data_implicit_gemm_v
4r4
_nchw_kcyx_nkhw.hpp"
#include "device_convolution_backward_data_implicit_gemm_v
1r1
_nchw_kcyx_nkhw.hpp"
#include "device_convolution_backward_data_implicit_gemm_v
4r5
_nchw_kcyx_nkhw.hpp"
#include "device_convolution_backward_data_implicit_gemm_v
1r2
_nchw_kcyx_nkhw.hpp"
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
...
@@ -346,9 +346,9 @@ int main(int argc, char* argv[])
...
@@ -346,9 +346,9 @@ int main(int argc, char* argv[])
}
}
#if 0
#if 0
device_convolution_backward_data_implicit_gemm_v
4r4
_nchw_kcyx_nkhw
device_convolution_backward_data_implicit_gemm_v
1r1
_nchw_kcyx_nkhw
#else
#else
device_convolution_backward_data_implicit_gemm_v
4r5
_nchw_kcyx_nkhw
device_convolution_backward_data_implicit_gemm_v
1r2
_nchw_kcyx_nkhw
#endif
#endif
(
in_nchw_desc
,
(
in_nchw_desc
,
in_nchw_device
,
in_nchw_device
,
...
...
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