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
83f9da8c
Commit
83f9da8c
authored
Apr 13, 2023
by
rocking
Browse files
Revise layout of group conv quantization client example
parent
898e1d22
Changes
7
Hide whitespace changes
Inline
Side-by-side
Showing
7 changed files
with
113 additions
and
89 deletions
+113
-89
client_example/07_grouped_convnd_fwd/grouped_conv2d_fwd.cpp
client_example/07_grouped_convnd_fwd/grouped_conv2d_fwd.cpp
+1
-1
client_example/09_quantization/conv2d_fwd_bias_relu_perchannel_quantization.cpp
...tization/conv2d_fwd_bias_relu_perchannel_quantization.cpp
+23
-19
client_example/09_quantization/conv2d_fwd_bias_relu_perlayer_quantization.cpp
...antization/conv2d_fwd_bias_relu_perlayer_quantization.cpp
+15
-11
client_example/09_quantization/conv2d_fwd_bias_tanh_perchannel_quantization.cpp
...tization/conv2d_fwd_bias_tanh_perchannel_quantization.cpp
+16
-12
client_example/09_quantization/conv2d_fwd_bias_tanh_perlayer_quantization.cpp
...antization/conv2d_fwd_bias_tanh_perlayer_quantization.cpp
+22
-18
client_example/09_quantization/conv2d_fwd_perchannel_quantization.cpp
...le/09_quantization/conv2d_fwd_perchannel_quantization.cpp
+22
-18
client_example/09_quantization/conv2d_fwd_perlayer_quantization.cpp
...mple/09_quantization/conv2d_fwd_perlayer_quantization.cpp
+14
-10
No files found.
client_example/07_grouped_convnd_fwd/grouped_conv2d_fwd.cpp
View file @
83f9da8c
...
@@ -52,7 +52,7 @@ struct SimpleDeviceMem
...
@@ -52,7 +52,7 @@ struct SimpleDeviceMem
int
main
()
int
main
()
{
{
// We have NHWGC/GKYXC/NHWGK in memory space
// We have NHWGC/GKYXC/NHWGK
(x, weight, y)
in memory space
// However, CK's API only accept length and stride with order of GNCHW/GKCYX/GNCHW
// However, CK's API only accept length and stride with order of GNCHW/GKCYX/GNCHW
// Hence, we need to adjust the order of stride
// Hence, we need to adjust the order of stride
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
...
...
client_example/09_quantization/conv2d_fwd_bias_relu_perchannel_quantization.cpp
View file @
83f9da8c
...
@@ -17,26 +17,26 @@ using BiasDataType = int32_t;
...
@@ -17,26 +17,26 @@ using BiasDataType = int32_t;
using
RequantScaleDataType
=
float
;
using
RequantScaleDataType
=
float
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWC
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
C
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
BiasLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
BiasLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
RequantScaleLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
RequantScaleLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWK
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
K
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
ActivationOp
=
ck
::
tensor_operation
::
element_wise
::
Relu
;
using
ActivationOp
=
ck
::
tensor_operation
::
element_wise
::
Relu
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Add_Activation_Mul2_Clamp
<
ActivationOp
>
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Add_Activation_Mul2_Clamp
<
ActivationOp
>
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
G
=
1
;
static
constexpr
ck
::
index_t
G
=
4
;
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
K
=
64
;
// output channel
static
constexpr
ck
::
index_t
K
=
32
;
// output channel
static
constexpr
ck
::
index_t
C
=
192
;
// input channel
static
constexpr
ck
::
index_t
C
=
64
;
// input channel
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
static
constexpr
ck
::
index_t
Wi
=
71
;
// input W
static
constexpr
ck
::
index_t
Wi
=
71
;
// input W
static
constexpr
ck
::
index_t
Ho
=
36
;
// output H
static
constexpr
ck
::
index_t
Ho
=
36
;
// output H
static
constexpr
ck
::
index_t
Wo
=
36
;
// output W
static
constexpr
ck
::
index_t
Wo
=
36
;
// output W
struct
SimpleDeviceMem
struct
SimpleDeviceMem
{
{
SimpleDeviceMem
()
=
delete
;
SimpleDeviceMem
()
=
delete
;
...
@@ -55,8 +55,11 @@ struct SimpleDeviceMem
...
@@ -55,8 +55,11 @@ struct SimpleDeviceMem
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
// We have NHWGC/GKYXC/NHWGK (x, weight, y) in memory space
// However, CK's API only accept length and stride with order of GNCHW/GKCYX/GNCHW
// Hence, we need to adjust the order of stride
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
N
*
Hi
*
Wi
*
C
,
Hi
*
Wi
*
C
,
1
,
Wi
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
bias_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
bias_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
...
@@ -64,17 +67,18 @@ int main(int argc, char* argv[])
...
@@ -64,17 +67,18 @@ int main(int argc, char* argv[])
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
N
*
Ho
*
Wo
*
K
,
Ho
*
Wo
*
K
,
1
,
Wo
*
K
,
K
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
C
,
Ho
*
Wo
*
G
*
C
,
1
,
Wo
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
C
);
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
G
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
G
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
bias
(
sizeof
(
BiasDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
bias
(
sizeof
(
BiasDataType
)
*
G
*
K
);
SimpleDeviceMem
requant_scale
(
sizeof
(
RequantScaleDataType
)
*
K
);
SimpleDeviceMem
requant_scale
(
sizeof
(
RequantScaleDataType
)
*
G
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
G
*
K
);
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
NumDimSpatial
,
...
...
client_example/09_quantization/conv2d_fwd_bias_relu_perlayer_quantization.cpp
View file @
83f9da8c
...
@@ -16,19 +16,19 @@ using WeiDataType = int8_t;
...
@@ -16,19 +16,19 @@ using WeiDataType = int8_t;
using
BiasDataType
=
int32_t
;
using
BiasDataType
=
int32_t
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWC
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
C
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
BiasLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
BiasLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWK
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
K
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
ActivationOp
=
ck
::
tensor_operation
::
element_wise
::
Relu
;
using
ActivationOp
=
ck
::
tensor_operation
::
element_wise
::
Relu
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Add_Activation_Mul_Clamp
<
ActivationOp
>
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Add_Activation_Mul_Clamp
<
ActivationOp
>
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
G
=
1
;
static
constexpr
ck
::
index_t
G
=
4
;
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
K
=
64
;
// output channel
static
constexpr
ck
::
index_t
K
=
32
;
// output channel
static
constexpr
ck
::
index_t
C
=
192
;
// input channel
static
constexpr
ck
::
index_t
C
=
64
;
// input channel
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
...
@@ -55,23 +55,27 @@ struct SimpleDeviceMem
...
@@ -55,23 +55,27 @@ struct SimpleDeviceMem
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
// We have NHWGC/GKYXC/NHWGK (x, weight, y) in memory space
// However, CK's API only accept length and stride with order of GNCHW/GKCYX/GNCHW
// Hence, we need to adjust the order of stride
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
N
*
Hi
*
Wi
*
C
,
Hi
*
Wi
*
C
,
1
,
Wi
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
bias_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
bias_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
bias_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
bias_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
N
*
Ho
*
Wo
*
K
,
Ho
*
Wo
*
K
,
1
,
Wo
*
K
,
K
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
C
,
Ho
*
Wo
*
G
*
C
,
1
,
Wo
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
C
);
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
G
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
G
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
bias
(
sizeof
(
BiasDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
bias
(
sizeof
(
BiasDataType
)
*
G
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
G
*
K
);
using
DeviceOp
=
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
...
...
client_example/09_quantization/conv2d_fwd_bias_tanh_perchannel_quantization.cpp
View file @
83f9da8c
...
@@ -17,21 +17,21 @@ using BiasDataType = int32_t;
...
@@ -17,21 +17,21 @@ using BiasDataType = int32_t;
using
RequantScaleDataType
=
float
;
using
RequantScaleDataType
=
float
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWC
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
C
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
BiasLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
BiasLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
RequantScaleLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
RequantScaleLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWK
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
K
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
ActivationOp
=
ck
::
tensor_operation
::
element_wise
::
TanH
;
using
ActivationOp
=
ck
::
tensor_operation
::
element_wise
::
TanH
;
using
OutElementOp
=
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Add_Mul2_Activation_Mul_Clamp
<
ActivationOp
>
;
ck
::
tensor_operation
::
element_wise
::
Add_Mul2_Activation_Mul_Clamp
<
ActivationOp
>
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
G
=
1
;
static
constexpr
ck
::
index_t
G
=
4
;
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
K
=
64
;
// output channel
static
constexpr
ck
::
index_t
K
=
32
;
// output channel
static
constexpr
ck
::
index_t
C
=
192
;
// input channel
static
constexpr
ck
::
index_t
C
=
64
;
// input channel
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
...
@@ -58,8 +58,11 @@ struct SimpleDeviceMem
...
@@ -58,8 +58,11 @@ struct SimpleDeviceMem
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
// We have NHWGC/GKYXC/NHWGK (x, weight, y) in memory space
// However, CK's API only accept length and stride with order of GNCHW/GKCYX/GNCHW
// Hence, we need to adjust the order of stride
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
N
*
Hi
*
Wi
*
C
,
Hi
*
Wi
*
C
,
1
,
Wi
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
bias_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
bias_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
...
@@ -67,17 +70,18 @@ int main(int argc, char* argv[])
...
@@ -67,17 +70,18 @@ int main(int argc, char* argv[])
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
N
*
Ho
*
Wo
*
K
,
Ho
*
Wo
*
K
,
1
,
Wo
*
K
,
K
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
C
,
Ho
*
Wo
*
G
*
C
,
1
,
Wo
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
C
);
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
G
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
G
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
bias
(
sizeof
(
BiasDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
bias
(
sizeof
(
BiasDataType
)
*
G
*
K
);
SimpleDeviceMem
requant_scale
(
sizeof
(
RequantScaleDataType
)
*
K
);
SimpleDeviceMem
requant_scale
(
sizeof
(
RequantScaleDataType
)
*
G
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
G
*
K
);
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
NumDimSpatial
,
...
...
client_example/09_quantization/conv2d_fwd_bias_tanh_perlayer_quantization.cpp
View file @
83f9da8c
...
@@ -16,25 +16,25 @@ using WeiDataType = int8_t;
...
@@ -16,25 +16,25 @@ using WeiDataType = int8_t;
using
BiasDataType
=
int32_t
;
using
BiasDataType
=
int32_t
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWC
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
C
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
BiasLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
BiasLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWK
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
K
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
ActivationOp
=
ck
::
tensor_operation
::
element_wise
::
TanH
;
using
ActivationOp
=
ck
::
tensor_operation
::
element_wise
::
TanH
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Add_Mul_Activation_Mul_Clamp
<
ActivationOp
>
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Add_Mul_Activation_Mul_Clamp
<
ActivationOp
>
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
G
=
1
;
static
constexpr
ck
::
index_t
G
=
4
;
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
K
=
64
;
// output channel
static
constexpr
ck
::
index_t
K
=
32
;
// output channel
static
constexpr
ck
::
index_t
C
=
192
;
// input channel
static
constexpr
ck
::
index_t
C
=
64
;
// input channel
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
static
constexpr
ck
::
index_t
Wi
=
71
;
// input W
static
constexpr
ck
::
index_t
Wi
=
71
;
// input W
static
constexpr
ck
::
index_t
Ho
=
36
;
// output H
static
constexpr
ck
::
index_t
Ho
=
36
;
// output H
static
constexpr
ck
::
index_t
Wo
=
36
;
// output W
static
constexpr
ck
::
index_t
Wo
=
36
;
// output W
static
constexpr
float
sacc
=
0.5
f
;
// scale of acc
static
constexpr
float
sacc
=
0.5
f
;
// scale of acc
static
constexpr
float
sz_inv
=
0.5
f
;
// inverse of scale_z
static
constexpr
float
sz_inv
=
0.5
f
;
// inverse of scale_z
...
@@ -56,23 +56,27 @@ struct SimpleDeviceMem
...
@@ -56,23 +56,27 @@ struct SimpleDeviceMem
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
// We have NHWGC/GKYXC/NHWGK (x, weight, y) in memory space
// However, CK's API only accept length and stride with order of GNCHW/GKCYX/GNCHW
// Hence, we need to adjust the order of stride
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
N
*
Hi
*
Wi
*
C
,
Hi
*
Wi
*
C
,
1
,
Wi
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
bias_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
bias_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
bias_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
bias_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
N
*
Ho
*
Wo
*
K
,
Ho
*
Wo
*
K
,
1
,
Wo
*
K
,
K
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
C
,
Ho
*
Wo
*
G
*
C
,
1
,
Wo
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
C
);
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
G
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
G
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
bias
(
sizeof
(
BiasDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
bias
(
sizeof
(
BiasDataType
)
*
G
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
G
*
K
);
using
DeviceOp
=
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
...
...
client_example/09_quantization/conv2d_fwd_perchannel_quantization.cpp
View file @
83f9da8c
...
@@ -16,25 +16,25 @@ using WeiDataType = int8_t;
...
@@ -16,25 +16,25 @@ using WeiDataType = int8_t;
using
RequantScaleDataType
=
float
;
using
RequantScaleDataType
=
float
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWC
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
C
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
RequantScaleLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
RequantScaleLayout
=
ck
::
tensor_layout
::
convolution
::
G_K
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWK
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
K
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
ActivationOp
=
PassThrough
;
using
ActivationOp
=
PassThrough
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Activation_Mul2_Clamp
<
ActivationOp
>
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Activation_Mul2_Clamp
<
ActivationOp
>
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
G
=
1
;
static
constexpr
ck
::
index_t
G
=
4
;
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
K
=
64
;
// output channel
static
constexpr
ck
::
index_t
K
=
32
;
// output channel
static
constexpr
ck
::
index_t
C
=
192
;
// input channel
static
constexpr
ck
::
index_t
C
=
64
;
// input channel
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
static
constexpr
ck
::
index_t
Wi
=
71
;
// input W
static
constexpr
ck
::
index_t
Wi
=
71
;
// input W
static
constexpr
ck
::
index_t
Ho
=
36
;
// output H
static
constexpr
ck
::
index_t
Ho
=
36
;
// output H
static
constexpr
ck
::
index_t
Wo
=
36
;
// output W
static
constexpr
ck
::
index_t
Wo
=
36
;
// output W
struct
SimpleDeviceMem
struct
SimpleDeviceMem
{
{
...
@@ -54,23 +54,27 @@ struct SimpleDeviceMem
...
@@ -54,23 +54,27 @@ struct SimpleDeviceMem
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
// We have NHWGC/GKYXC/NHWGK (x, weight, y) in memory space
// However, CK's API only accept length and stride with order of GNCHW/GKCYX/GNCHW
// Hence, we need to adjust the order of stride
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
N
*
Hi
*
Wi
*
C
,
Hi
*
Wi
*
C
,
1
,
Wi
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
requant_scale_strides
{
K
,
0
,
1
,
0
,
0
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
N
*
Ho
*
Wo
*
K
,
Ho
*
Wo
*
K
,
1
,
Wo
*
K
,
K
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
C
,
Ho
*
Wo
*
G
*
C
,
1
,
Wo
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
C
);
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
G
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
G
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
requant_scale
(
sizeof
(
RequantScaleDataType
)
*
K
);
SimpleDeviceMem
requant_scale
(
sizeof
(
RequantScaleDataType
)
*
G
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
G
*
K
);
using
DeviceOp
=
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
...
...
client_example/09_quantization/conv2d_fwd_perlayer_quantization.cpp
View file @
83f9da8c
...
@@ -15,18 +15,18 @@ using InDataType = int8_t;
...
@@ -15,18 +15,18 @@ using InDataType = int8_t;
using
WeiDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWC
;
using
InLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
C
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
WeiLayout
=
ck
::
tensor_layout
::
convolution
::
GKYXC
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
G
NHWK
;
using
OutLayout
=
ck
::
tensor_layout
::
convolution
::
NHW
G
K
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
ActivationOp
=
PassThrough
;
using
ActivationOp
=
PassThrough
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Activation_Mul_Clamp
<
ActivationOp
>
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
Activation_Mul_Clamp
<
ActivationOp
>
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
NumDimSpatial
=
2
;
static
constexpr
ck
::
index_t
G
=
1
;
static
constexpr
ck
::
index_t
G
=
4
;
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
N
=
4
;
// batch size
static
constexpr
ck
::
index_t
K
=
64
;
// output channel
static
constexpr
ck
::
index_t
K
=
32
;
// output channel
static
constexpr
ck
::
index_t
C
=
192
;
// input channel
static
constexpr
ck
::
index_t
C
=
64
;
// input channel
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
Y
=
3
;
// filter H
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
X
=
3
;
// filter W
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
static
constexpr
ck
::
index_t
Hi
=
71
;
// input H
...
@@ -53,20 +53,24 @@ struct SimpleDeviceMem
...
@@ -53,20 +53,24 @@ struct SimpleDeviceMem
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
{
{
// We have NHWGC/GKYXC/NHWGK (x, weight, y) in memory space
// However, CK's API only accept length and stride with order of GNCHW/GKCYX/GNCHW
// Hence, we need to adjust the order of stride
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_lengths
{
G
,
N
,
C
,
Hi
,
Wi
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
N
*
Hi
*
Wi
*
C
,
Hi
*
Wi
*
C
,
1
,
Wi
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
in_strides
{
C
,
Hi
*
Wi
*
G
*
C
,
1
,
Wi
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_lengths
{
G
,
K
,
C
,
Y
,
X
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
weight_strides
{
K
*
Y
*
X
*
C
,
Y
*
X
*
C
,
1
,
X
*
C
,
C
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_lengths
{
G
,
N
,
K
,
Ho
,
Wo
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
N
*
Ho
*
Wo
*
K
,
Ho
*
Wo
*
K
,
1
,
Wo
*
K
,
K
};
std
::
array
<
ck
::
index_t
,
5
>
out_strides
{
C
,
Ho
*
Wo
*
G
*
C
,
1
,
Wo
*
G
*
C
,
G
*
C
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_left_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
in_right_pad
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_strides
{
2
,
2
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
std
::
array
<
ck
::
index_t
,
2
>
conv_dilations
{
1
,
1
};
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
C
);
SimpleDeviceMem
in
(
sizeof
(
InDataType
)
*
N
*
Hi
*
Wi
*
G
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
wei
(
sizeof
(
WeiDataType
)
*
G
*
K
*
Y
*
X
*
C
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
K
);
SimpleDeviceMem
out
(
sizeof
(
OutDataType
)
*
N
*
Ho
*
Wo
*
G
*
K
);
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceGroupedConvFwdMultipleD
<
NumDimSpatial
,
InLayout
,
InLayout
,
...
...
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