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
c6683eea
Commit
c6683eea
authored
Mar 15, 2023
by
rocking
Browse files
Remove CShuffleDataType in dlops
Let acc and CShuffleDataType be the same in xdlops
parent
25e3772e
Changes
11
Hide whitespace changes
Inline
Side-by-side
Showing
11 changed files
with
19 additions
and
23 deletions
+19
-23
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_bias_relu_perchannel_quantization_int8.cpp
.../conv2d_fwd_dl_bias_relu_perchannel_quantization_int8.cpp
+0
-1
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_bias_relu_perlayer_quantization_int8.cpp
...on/conv2d_fwd_dl_bias_relu_perlayer_quantization_int8.cpp
+5
-6
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_perchannel_quantization_int8.cpp
...antization/conv2d_fwd_dl_perchannel_quantization_int8.cpp
+0
-1
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_perlayer_quantization_int8.cpp
...quantization/conv2d_fwd_dl_perlayer_quantization_int8.cpp
+4
-5
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_bias_relu_perchannel_quantization_int8.cpp
...conv2d_fwd_xdl_bias_relu_perchannel_quantization_int8.cpp
+1
-1
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_bias_relu_perlayer_quantization_int8.cpp
...n/conv2d_fwd_xdl_bias_relu_perlayer_quantization_int8.cpp
+1
-1
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_perchannel_quantization_int8.cpp
...ntization/conv2d_fwd_xdl_perchannel_quantization_int8.cpp
+1
-1
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_perlayer_quantization_int8.cpp
...uantization/conv2d_fwd_xdl_perlayer_quantization_int8.cpp
+1
-1
example/40_conv2d_fwd_quantization/run_conv2d_fwd_bias_relu_perchannel_quantization_example.inc
..._conv2d_fwd_bias_relu_perchannel_quantization_example.inc
+2
-2
example/40_conv2d_fwd_quantization/run_conv2d_fwd_bias_relu_perlayer_quantization_example.inc
...un_conv2d_fwd_bias_relu_perlayer_quantization_example.inc
+2
-2
example/40_conv2d_fwd_quantization/run_conv2d_fwd_perchannel_quantization_example.inc
...zation/run_conv2d_fwd_perchannel_quantization_example.inc
+2
-2
No files found.
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_bias_relu_perchannel_quantization_int8.cpp
View file @
c6683eea
...
@@ -9,7 +9,6 @@ using WeiDataType = int8_t;
...
@@ -9,7 +9,6 @@ using WeiDataType = int8_t;
using
BiasDataType
=
int32_t
;
using
BiasDataType
=
int32_t
;
using
RequantScaleDataType
=
float
;
using
RequantScaleDataType
=
float
;
using
AccDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
CShuffleDataType
=
int32_t
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_bias_relu_perlayer_quantization_int8.cpp
View file @
c6683eea
...
@@ -4,12 +4,11 @@
...
@@ -4,12 +4,11 @@
#include "common.hpp"
#include "common.hpp"
#include "ck/tensor_operation/gpu/device/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp"
#include "ck/tensor_operation/gpu/device/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp"
using
InDataType
=
int8_t
;
using
InDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
BiasDataType
=
int32_t
;
using
BiasDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
CShuffleDataType
=
int32_t
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
using
S
=
ck
::
Sequence
<
Is
...
>
;
using
S
=
ck
::
Sequence
<
Is
...
>
;
...
...
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_perchannel_quantization_int8.cpp
View file @
c6683eea
...
@@ -8,7 +8,6 @@ using InDataType = int8_t;
...
@@ -8,7 +8,6 @@ using InDataType = int8_t;
using
WeiDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
RequantScaleDataType
=
float
;
using
RequantScaleDataType
=
float
;
using
AccDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
CShuffleDataType
=
int32_t
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_perlayer_quantization_int8.cpp
View file @
c6683eea
...
@@ -4,11 +4,10 @@
...
@@ -4,11 +4,10 @@
#include "common.hpp"
#include "common.hpp"
#include "ck/tensor_operation/gpu/device/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp"
#include "ck/tensor_operation/gpu/device/device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp"
using
InDataType
=
int8_t
;
using
InDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
AccDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
CShuffleDataType
=
int32_t
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
using
S
=
ck
::
Sequence
<
Is
...
>
;
using
S
=
ck
::
Sequence
<
Is
...
>
;
...
...
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_bias_relu_perchannel_quantization_int8.cpp
View file @
c6683eea
...
@@ -9,7 +9,7 @@ using WeiDataType = int8_t;
...
@@ -9,7 +9,7 @@ using WeiDataType = int8_t;
using
BiasDataType
=
int32_t
;
using
BiasDataType
=
int32_t
;
using
RequantScaleDataType
=
float
;
using
RequantScaleDataType
=
float
;
using
AccDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
CShuffleDataType
=
int32_t
;
using
CShuffleDataType
=
AccDataType
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_bias_relu_perlayer_quantization_int8.cpp
View file @
c6683eea
...
@@ -8,7 +8,7 @@ using InDataType = int8_t;
...
@@ -8,7 +8,7 @@ using InDataType = int8_t;
using
WeiDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
BiasDataType
=
int32_t
;
using
BiasDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
CShuffleDataType
=
int32_t
;
using
CShuffleDataType
=
AccDataType
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_perchannel_quantization_int8.cpp
View file @
c6683eea
...
@@ -8,7 +8,7 @@ using InDataType = int8_t;
...
@@ -8,7 +8,7 @@ using InDataType = int8_t;
using
WeiDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
RequantScaleDataType
=
float
;
using
RequantScaleDataType
=
float
;
using
AccDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
CShuffleDataType
=
int32_t
;
using
CShuffleDataType
=
AccDataType
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_perlayer_quantization_int8.cpp
View file @
c6683eea
...
@@ -7,7 +7,7 @@
...
@@ -7,7 +7,7 @@
using
InDataType
=
int8_t
;
using
InDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
WeiDataType
=
int8_t
;
using
AccDataType
=
int32_t
;
using
AccDataType
=
int32_t
;
using
CShuffleDataType
=
int32_t
;
using
CShuffleDataType
=
AccDataType
;
using
OutDataType
=
int8_t
;
using
OutDataType
=
int8_t
;
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
example/40_conv2d_fwd_quantization/run_conv2d_fwd_bias_relu_perchannel_quantization_example.inc
View file @
c6683eea
...
@@ -129,12 +129,12 @@ bool run_grouped_conv_fwd(bool do_verification,
...
@@ -129,12 +129,12 @@ bool run_grouped_conv_fwd(bool do_verification,
if
(
do_verification
)
if
(
do_verification
)
{
{
Tensor
<
CShuffle
DataType
>
c_host
(
out_g_n_k_wos_desc
);
Tensor
<
Acc
DataType
>
c_host
(
out_g_n_k_wos_desc
);
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvFwd
<
NDimSpatial
,
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvFwd
<
NDimSpatial
,
InDataType
,
InDataType
,
WeiDataType
,
WeiDataType
,
CShuffle
DataType
,
Acc
DataType
,
InElementOp
,
InElementOp
,
WeiElementOp
,
WeiElementOp
,
PassThrough
>
();
PassThrough
>
();
...
...
example/40_conv2d_fwd_quantization/run_conv2d_fwd_bias_relu_perlayer_quantization_example.inc
View file @
c6683eea
...
@@ -118,12 +118,12 @@ bool run_grouped_conv_fwd(bool do_verification,
...
@@ -118,12 +118,12 @@ bool run_grouped_conv_fwd(bool do_verification,
if
(
do_verification
)
if
(
do_verification
)
{
{
Tensor
<
CShuffle
DataType
>
c_host
(
out_g_n_k_wos_desc
);
Tensor
<
Acc
DataType
>
c_host
(
out_g_n_k_wos_desc
);
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvFwd
<
NDimSpatial
,
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvFwd
<
NDimSpatial
,
InDataType
,
InDataType
,
WeiDataType
,
WeiDataType
,
CShuffle
DataType
,
Acc
DataType
,
InElementOp
,
InElementOp
,
WeiElementOp
,
WeiElementOp
,
PassThrough
>
();
PassThrough
>
();
...
...
example/40_conv2d_fwd_quantization/run_conv2d_fwd_perchannel_quantization_example.inc
View file @
c6683eea
...
@@ -119,12 +119,12 @@ bool run_grouped_conv_fwd(bool do_verification,
...
@@ -119,12 +119,12 @@ bool run_grouped_conv_fwd(bool do_verification,
if
(
do_verification
)
if
(
do_verification
)
{
{
Tensor
<
CShuffle
DataType
>
c_host
(
out_g_n_k_wos_desc
);
Tensor
<
Acc
DataType
>
c_host
(
out_g_n_k_wos_desc
);
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvFwd
<
NDimSpatial
,
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvFwd
<
NDimSpatial
,
InDataType
,
InDataType
,
WeiDataType
,
WeiDataType
,
CShuffle
DataType
,
Acc
DataType
,
InElementOp
,
InElementOp
,
WeiElementOp
,
WeiElementOp
,
PassThrough
>
();
PassThrough
>
();
...
...
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