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
0e81cc18
Commit
0e81cc18
authored
Jul 18, 2022
by
Chao Liu
Browse files
clean
parent
d84da366
Changes
9
Show whitespace changes
Inline
Side-by-side
Showing
9 changed files
with
280 additions
and
594 deletions
+280
-594
example/09_convnd_fwd/convnd_fwd_common.hpp
example/09_convnd_fwd/convnd_fwd_common.hpp
+4
-69
example/17_convnd_bwd_data/convnd_bwd_data_common.hpp
example/17_convnd_bwd_data/convnd_bwd_data_common.hpp
+22
-87
example/20_convnd_bwd_weight/convnd_bwd_weight_common.hpp
example/20_convnd_bwd_weight/convnd_bwd_weight_common.hpp
+5
-70
library/include/ck/library/utility/convolution_host_tensor_descriptor_helper.hpp
...ary/utility/convolution_host_tensor_descriptor_helper.hpp
+124
-0
profiler/include/profile_conv_bwd_data_impl.hpp
profiler/include/profile_conv_bwd_data_impl.hpp
+21
-86
profiler/include/profile_conv_bwd_weight_impl.hpp
profiler/include/profile_conv_bwd_weight_impl.hpp
+21
-86
profiler/include/profile_conv_fwd_impl.hpp
profiler/include/profile_conv_fwd_impl.hpp
+28
-93
test/conv_util/conv_util.cpp
test/conv_util/conv_util.cpp
+38
-91
test/reference_conv_fwd/reference_conv_fwd.cpp
test/reference_conv_fwd/reference_conv_fwd.cpp
+17
-12
No files found.
example/09_convnd_fwd/convnd_fwd_common.hpp
View file @
0e81cc18
...
...
@@ -15,6 +15,7 @@
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_conv_fwd.hpp"
void
print_helper_msg
()
...
...
@@ -110,75 +111,9 @@ int run_conv_fwd(bool do_verification,
const
WeiElementOp
&
wei_element_op
,
const
OutElementOp
&
out_element_op
)
{
// make host tensor descritpor
auto
f_nhwc_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nhwc_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nhwc_lengths
.
insert
(
nhwc_lengths
.
begin
()
+
1
,
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nhwc_lengths
);
};
auto
f_nchw_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nchw_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nchw_lengths
.
insert
(
nchw_lengths
.
end
(),
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nchw_lengths
);
};
HostTensorDescriptor
in_desc
,
wei_desc
,
out_desc
;
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NWC
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NHWC
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWC
>
)
{
in_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
else
if
constexpr
(
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCW
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCHW
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCDHW
>
)
{
in_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KXC
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KYXC
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KZYXC
>
)
{
wei_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
else
if
constexpr
(
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCX
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCYX
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCZYX
>
)
{
wei_desc
=
f_nchw_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NWK
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NHWK
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWK
>
)
{
out_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
else
if
constexpr
(
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKW
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKHW
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
)
{
out_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
const
auto
in_desc
=
ck
::
utils
::
conv
::
get_input_host_tensor_descriptor
<
InLayout
>
(
conv_param
);
const
auto
wei_desc
=
ck
::
utils
::
conv
::
get_weight_host_tensor_descriptor
<
WeiLayout
>
(
conv_param
);
const
auto
out_desc
=
ck
::
utils
::
conv
::
get_output_host_tensor_descriptor
<
OutLayout
>
(
conv_param
);
Tensor
<
InDataType
>
in
(
in_desc
);
Tensor
<
WeiDataType
>
wei
(
wei_desc
);
...
...
example/17_convnd_bwd_data/convnd_bwd_data_common.hpp
View file @
0e81cc18
...
...
@@ -11,10 +11,11 @@
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/library/utility/check_err.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/utility/device_memory.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_conv_bwd_data.hpp"
ck
::
tensor_operation
::
device
::
ConvParams
...
...
@@ -105,80 +106,14 @@ template <ck::index_t NDimSpatial,
int
run_conv_bwd_data
(
bool
do_verification
,
int
init_method
,
bool
time_kernel
,
const
ck
::
tensor_operation
::
device
::
ConvParams
&
param
s
,
const
ck
::
tensor_operation
::
device
::
ConvParams
&
conv_
param
,
const
InElementOp
&
in_element_op
,
const
WeiElementOp
&
wei_element_op
,
const
OutElementOp
&
out_element_op
)
{
// make host tensor descritpor
auto
f_nhwc_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nhwc_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nhwc_lengths
.
insert
(
nhwc_lengths
.
begin
()
+
1
,
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nhwc_lengths
);
};
auto
f_nchw_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nchw_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nchw_lengths
.
insert
(
nchw_lengths
.
end
(),
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nchw_lengths
);
};
HostTensorDescriptor
in_desc
,
wei_desc
,
out_desc
;
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NWC
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NHWC
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWC
>
)
{
in_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
else
if
constexpr
(
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCW
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCHW
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCDHW
>
)
{
in_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KXC
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KYXC
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KZYXC
>
)
{
wei_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
else
if
constexpr
(
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCX
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCYX
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCZYX
>
)
{
wei_desc
=
f_nchw_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NWK
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NHWK
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWK
>
)
{
out_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
else
if
constexpr
(
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKW
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKHW
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
)
{
out_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
const
auto
in_desc
=
ck
::
utils
::
conv
::
get_input_host_tensor_descriptor
<
InLayout
>
(
conv_param
);
const
auto
wei_desc
=
ck
::
utils
::
conv
::
get_weight_host_tensor_descriptor
<
WeiLayout
>
(
conv_param
);
const
auto
out_desc
=
ck
::
utils
::
conv
::
get_output_host_tensor_descriptor
<
OutLayout
>
(
conv_param
);
Tensor
<
InDataType
>
in_host
(
in_desc
);
Tensor
<
InDataType
>
in_device
(
in_desc
);
...
...
@@ -217,16 +152,16 @@ int run_conv_bwd_data(bool do_verification,
auto
argument
=
conv
.
MakeArgument
(
static_cast
<
InDataType
*>
(
in_device_buf
.
GetDeviceBuffer
()),
static_cast
<
WeiDataType
*>
(
wei_device_buf
.
GetDeviceBuffer
()),
static_cast
<
OutDataType
*>
(
out_device_buf
.
GetDeviceBuffer
()),
param
s
.
N_
,
param
s
.
K_
,
param
s
.
C_
,
param
s
.
input_spatial_lengths_
,
param
s
.
filter_spatial_lengths_
,
param
s
.
GetOutputSpatialLengths
(),
param
s
.
conv_filter_strides_
,
param
s
.
conv_filter_dilations_
,
param
s
.
input_left_pads_
,
param
s
.
input_right_pads_
,
conv_
param
.
N_
,
conv_
param
.
K_
,
conv_
param
.
C_
,
conv_
param
.
input_spatial_lengths_
,
conv_
param
.
filter_spatial_lengths_
,
conv_
param
.
GetOutputSpatialLengths
(),
conv_
param
.
conv_filter_strides_
,
conv_
param
.
conv_filter_dilations_
,
conv_
param
.
input_left_pads_
,
conv_
param
.
input_right_pads_
,
in_element_op
,
wei_element_op
,
out_element_op
);
...
...
@@ -240,8 +175,8 @@ int run_conv_bwd_data(bool do_verification,
float
ave_time
=
invoker
.
Run
(
argument
,
StreamConfig
{
nullptr
,
time_kernel
});
std
::
size_t
flop
=
param
s
.
GetFlops
();
std
::
size_t
num_btype
=
param
s
.
GetByte
<
InDataType
,
WeiDataType
,
OutDataType
>
();
std
::
size_t
flop
=
conv_
param
.
GetFlops
();
std
::
size_t
num_btype
=
conv_
param
.
GetByte
<
InDataType
,
WeiDataType
,
OutDataType
>
();
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
ave_time
;
...
...
@@ -268,10 +203,10 @@ int run_conv_bwd_data(bool do_verification,
auto
ref_argument
=
ref_conv
.
MakeArgument
(
in_host
,
wei
,
out
,
param
s
.
conv_filter_strides_
,
param
s
.
conv_filter_dilations_
,
param
s
.
input_left_pads_
,
param
s
.
input_right_pads_
,
conv_
param
.
conv_filter_strides_
,
conv_
param
.
conv_filter_dilations_
,
conv_
param
.
input_left_pads_
,
conv_
param
.
input_right_pads_
,
in_element_op
,
wei_element_op
,
out_element_op
);
...
...
example/20_convnd_bwd_weight/convnd_bwd_weight_common.hpp
View file @
0e81cc18
...
...
@@ -11,10 +11,11 @@
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/library/utility/check_err.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/utility/device_memory.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_conv_bwd_weight.hpp"
void
print_helper_msg
()
...
...
@@ -111,75 +112,9 @@ int run_conv_bwd_weight(bool do_verification,
const
OutElementOp
&
out_element_op
,
ck
::
index_t
split_k
)
{
// make host tensor descritpor
auto
f_nhwc_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nhwc_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nhwc_lengths
.
insert
(
nhwc_lengths
.
begin
()
+
1
,
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nhwc_lengths
);
};
auto
f_nchw_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nchw_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nchw_lengths
.
insert
(
nchw_lengths
.
end
(),
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nchw_lengths
);
};
HostTensorDescriptor
in_desc
,
wei_desc
,
out_desc
;
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NWC
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NHWC
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWC
>
)
{
in_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
else
if
constexpr
(
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCW
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCHW
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCDHW
>
)
{
in_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KXC
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KYXC
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KZYXC
>
)
{
wei_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
else
if
constexpr
(
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCX
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCYX
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCZYX
>
)
{
wei_desc
=
f_nchw_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NWK
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NHWK
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWK
>
)
{
out_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
else
if
constexpr
(
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKW
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKHW
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
)
{
out_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
const
auto
in_desc
=
ck
::
utils
::
conv
::
get_input_host_tensor_descriptor
<
InLayout
>
(
conv_param
);
const
auto
wei_desc
=
ck
::
utils
::
conv
::
get_weight_host_tensor_descriptor
<
WeiLayout
>
(
conv_param
);
const
auto
out_desc
=
ck
::
utils
::
conv
::
get_output_host_tensor_descriptor
<
OutLayout
>
(
conv_param
);
Tensor
<
InDataType
>
in
(
in_desc
);
Tensor
<
WeiDataType
>
wei_host_result
(
wei_desc
);
...
...
library/include/ck/library/utility/convolution_host_tensor_descriptor_helper.hpp
0 → 100644
View file @
0e81cc18
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include "ck/ck.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
namespace
ck
{
namespace
utils
{
namespace
conv
{
template
<
typename
InLayout
>
HostTensorDescriptor
get_input_host_tensor_descriptor
(
const
ck
::
tensor_operation
::
device
::
ConvParams
&
param
)
{
if
constexpr
(
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NWC
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NHWC
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWC
>
)
{
std
::
vector
<
std
::
size_t
>
nhwc_lengths
{
static_cast
<
std
::
size_t
>
(
param
.
N_
),
static_cast
<
std
::
size_t
>
(
param
.
C_
)};
nhwc_lengths
.
insert
(
nhwc_lengths
.
begin
()
+
1
,
param
.
input_spatial_lengths_
.
begin
(),
param
.
input_spatial_lengths_
.
end
());
return
HostTensorDescriptor
(
nhwc_lengths
);
}
else
if
constexpr
(
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCW
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCHW
>
||
ck
::
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCDHW
>
)
{
std
::
vector
<
std
::
size_t
>
nchw_lengths
{
static_cast
<
std
::
size_t
>
(
param
.
N_
),
static_cast
<
std
::
size_t
>
(
param
.
C_
)};
nchw_lengths
.
insert
(
nchw_lengths
.
end
(),
param
.
input_spatial_lengths_
.
begin
(),
param
.
input_spatial_lengths_
.
end
());
return
HostTensorDescriptor
(
nchw_lengths
);
}
else
{
throw
std
::
runtime_error
(
"wrong! unsupported layout"
);
}
}
template
<
typename
WeiLayout
>
HostTensorDescriptor
get_weight_host_tensor_descriptor
(
const
ck
::
tensor_operation
::
device
::
ConvParams
&
param
)
{
if
constexpr
(
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KXC
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KYXC
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KZYXC
>
)
{
std
::
vector
<
std
::
size_t
>
kyxc_lengths
{
static_cast
<
std
::
size_t
>
(
param
.
K_
),
static_cast
<
std
::
size_t
>
(
param
.
C_
)};
kyxc_lengths
.
insert
(
kyxc_lengths
.
begin
()
+
1
,
param
.
filter_spatial_lengths_
.
begin
(),
param
.
filter_spatial_lengths_
.
end
());
return
HostTensorDescriptor
(
kyxc_lengths
);
}
else
if
constexpr
(
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCX
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCYX
>
||
ck
::
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCZYX
>
)
{
std
::
vector
<
std
::
size_t
>
kcyx_lengths
{
static_cast
<
std
::
size_t
>
(
param
.
K_
),
static_cast
<
std
::
size_t
>
(
param
.
C_
)};
kcyx_lengths
.
insert
(
kcyx_lengths
.
end
(),
param
.
filter_spatial_lengths_
.
begin
(),
param
.
filter_spatial_lengths_
.
end
());
return
HostTensorDescriptor
(
kcyx_lengths
);
}
else
{
throw
std
::
runtime_error
(
"wrong! unsupported layout"
);
}
}
template
<
typename
OutLayout
>
HostTensorDescriptor
get_output_host_tensor_descriptor
(
const
ck
::
tensor_operation
::
device
::
ConvParams
&
param
)
{
if
constexpr
(
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NWK
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NHWK
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWK
>
)
{
std
::
vector
<
std
::
size_t
>
nhwk_lengths
{
static_cast
<
std
::
size_t
>
(
param
.
N_
),
static_cast
<
std
::
size_t
>
(
param
.
K_
)};
nhwk_lengths
.
insert
(
nhwk_lengths
.
begin
()
+
1
,
param
.
output_spatial_lengths_
.
begin
(),
param
.
output_spatial_lengths_
.
end
());
return
HostTensorDescriptor
(
nhwk_lengths
);
}
else
if
constexpr
(
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKW
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKHW
>
||
ck
::
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
)
{
std
::
vector
<
std
::
size_t
>
nkhw_lengths
{
static_cast
<
std
::
size_t
>
(
param
.
N_
),
static_cast
<
std
::
size_t
>
(
param
.
K_
)};
nkhw_lengths
.
insert
(
nkhw_lengths
.
end
(),
param
.
output_spatial_lengths_
.
begin
(),
param
.
output_spatial_lengths_
.
end
());
return
HostTensorDescriptor
(
nkhw_lengths
);
}
else
{
throw
std
::
runtime_error
(
"wrong! unsupported layout"
);
}
}
}
// namespace conv
}
// namespace utils
}
// namespace ck
profiler/include/profile_conv_bwd_data_impl.hpp
View file @
0e81cc18
...
...
@@ -15,6 +15,7 @@
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_conv_bwd_data.hpp"
namespace
ck
{
...
...
@@ -57,7 +58,7 @@ bool profile_conv_bwd_data_impl(int do_verification,
int
init_method
,
bool
do_log
,
bool
time_kernel
,
const
ck
::
tensor_operation
::
device
::
ConvParams
&
param
s
)
const
ck
::
tensor_operation
::
device
::
ConvParams
&
conv_
param
)
{
using
InElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
WeiElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
...
...
@@ -67,75 +68,9 @@ bool profile_conv_bwd_data_impl(int do_verification,
const
auto
wei_element_op
=
WeiElementOp
{};
const
auto
out_element_op
=
OutElementOp
{};
// make host tensor descritpor
auto
f_nhwc_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nhwc_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nhwc_lengths
.
insert
(
nhwc_lengths
.
begin
()
+
1
,
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nhwc_lengths
);
};
auto
f_nchw_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nchw_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nchw_lengths
.
insert
(
nchw_lengths
.
end
(),
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nchw_lengths
);
};
HostTensorDescriptor
in_desc
,
wei_desc
,
out_desc
;
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NWC
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NHWC
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWC
>
)
{
in_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
else
if
constexpr
(
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCW
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCHW
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCDHW
>
)
{
in_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KXC
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KYXC
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KZYXC
>
)
{
wei_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
else
if
constexpr
(
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCX
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCYX
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCZYX
>
)
{
wei_desc
=
f_nchw_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NWK
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NHWK
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWK
>
)
{
out_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
else
if
constexpr
(
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKW
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKHW
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
)
{
out_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
const
auto
in_desc
=
ck
::
utils
::
conv
::
get_input_host_tensor_descriptor
<
InLayout
>
(
conv_param
);
const
auto
wei_desc
=
ck
::
utils
::
conv
::
get_weight_host_tensor_descriptor
<
WeiLayout
>
(
conv_param
);
const
auto
out_desc
=
ck
::
utils
::
conv
::
get_output_host_tensor_descriptor
<
OutLayout
>
(
conv_param
);
Tensor
<
InDataType
>
input_host_result
(
in_desc
);
Tensor
<
InDataType
>
input_device_result
(
in_desc
);
...
...
@@ -183,10 +118,10 @@ bool profile_conv_bwd_data_impl(int do_verification,
auto
ref_argument
=
ref_conv
.
MakeArgument
(
input_host_result
,
weight
,
output
,
param
s
.
conv_filter_strides_
,
param
s
.
conv_filter_dilations_
,
param
s
.
input_left_pads_
,
param
s
.
input_right_pads_
,
conv_
param
.
conv_filter_strides_
,
conv_
param
.
conv_filter_dilations_
,
conv_
param
.
input_left_pads_
,
conv_
param
.
input_right_pads_
,
InElementOp
{},
WeiElementOp
{},
OutElementOp
{});
...
...
@@ -224,16 +159,16 @@ bool profile_conv_bwd_data_impl(int do_verification,
op_ptr
->
MakeArgumentPointer
(
static_cast
<
InDataType
*>
(
in_device_buf
.
GetDeviceBuffer
()),
static_cast
<
WeiDataType
*>
(
wei_device_buf
.
GetDeviceBuffer
()),
static_cast
<
OutDataType
*>
(
out_device_buf
.
GetDeviceBuffer
()),
param
s
.
N_
,
param
s
.
K_
,
param
s
.
C_
,
param
s
.
input_spatial_lengths_
,
param
s
.
filter_spatial_lengths_
,
param
s
.
output_spatial_lengths_
,
param
s
.
conv_filter_strides_
,
param
s
.
conv_filter_dilations_
,
param
s
.
input_left_pads_
,
param
s
.
input_right_pads_
,
conv_
param
.
N_
,
conv_
param
.
K_
,
conv_
param
.
C_
,
conv_
param
.
input_spatial_lengths_
,
conv_
param
.
filter_spatial_lengths_
,
conv_
param
.
output_spatial_lengths_
,
conv_
param
.
conv_filter_strides_
,
conv_
param
.
conv_filter_dilations_
,
conv_
param
.
input_left_pads_
,
conv_
param
.
input_right_pads_
,
in_element_op
,
wei_element_op
,
out_element_op
);
...
...
@@ -251,8 +186,8 @@ bool profile_conv_bwd_data_impl(int do_verification,
float
avg_time
=
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
StreamConfig
{
nullptr
,
time_kernel
});
std
::
size_t
flop
=
param
s
.
GetFlops
();
std
::
size_t
num_btype
=
param
s
.
GetByte
<
InDataType
,
WeiDataType
,
OutDataType
>
();
std
::
size_t
flop
=
conv_
param
.
GetFlops
();
std
::
size_t
num_btype
=
conv_
param
.
GetByte
<
InDataType
,
WeiDataType
,
OutDataType
>
();
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
avg_time
;
float
gb_per_sec
=
num_btype
/
1.E6
/
avg_time
;
...
...
profiler/include/profile_conv_bwd_weight_impl.hpp
View file @
0e81cc18
...
...
@@ -20,6 +20,7 @@
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_conv_bwd_weight.hpp"
namespace
ck
{
...
...
@@ -62,7 +63,7 @@ bool profile_conv_bwd_weight_impl(int do_verification,
int
init_method
,
bool
do_log
,
bool
time_kernel
,
const
ck
::
tensor_operation
::
device
::
ConvParams
&
param
s
,
const
ck
::
tensor_operation
::
device
::
ConvParams
&
conv_
param
,
ck
::
index_t
split_k
)
{
using
InElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
...
...
@@ -73,75 +74,9 @@ bool profile_conv_bwd_weight_impl(int do_verification,
const
auto
wei_element_op
=
WeiElementOp
{};
const
auto
out_element_op
=
OutElementOp
{};
// make host tensor descritpor
auto
f_nhwc_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nhwc_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nhwc_lengths
.
insert
(
nhwc_lengths
.
begin
()
+
1
,
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nhwc_lengths
);
};
auto
f_nchw_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nchw_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nchw_lengths
.
insert
(
nchw_lengths
.
end
(),
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nchw_lengths
);
};
HostTensorDescriptor
in_desc
,
wei_desc
,
out_desc
;
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NWC
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NHWC
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWC
>
)
{
in_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
else
if
constexpr
(
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCW
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCHW
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCDHW
>
)
{
in_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KXC
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KYXC
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KZYXC
>
)
{
wei_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
else
if
constexpr
(
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCX
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCYX
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCZYX
>
)
{
wei_desc
=
f_nchw_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NWK
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NHWK
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWK
>
)
{
out_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
else
if
constexpr
(
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKW
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKHW
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
)
{
out_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
const
auto
in_desc
=
ck
::
utils
::
conv
::
get_input_host_tensor_descriptor
<
InLayout
>
(
conv_param
);
const
auto
wei_desc
=
ck
::
utils
::
conv
::
get_weight_host_tensor_descriptor
<
WeiLayout
>
(
conv_param
);
const
auto
out_desc
=
ck
::
utils
::
conv
::
get_output_host_tensor_descriptor
<
OutLayout
>
(
conv_param
);
Tensor
<
InDataType
>
input
(
in_desc
);
Tensor
<
WeiDataType
>
weight_host_result
(
wei_desc
);
...
...
@@ -189,10 +124,10 @@ bool profile_conv_bwd_weight_impl(int do_verification,
auto
ref_argument
=
ref_conv
.
MakeArgument
(
input
,
weight_host_result
,
output
,
param
s
.
conv_filter_strides_
,
param
s
.
conv_filter_dilations_
,
param
s
.
input_left_pads_
,
param
s
.
input_right_pads_
,
conv_
param
.
conv_filter_strides_
,
conv_
param
.
conv_filter_dilations_
,
conv_
param
.
input_left_pads_
,
conv_
param
.
input_right_pads_
,
in_element_op
,
wei_element_op
,
out_element_op
);
...
...
@@ -231,16 +166,16 @@ bool profile_conv_bwd_weight_impl(int do_verification,
op_ptr
->
MakeArgumentPointer
(
static_cast
<
InDataType
*>
(
in_device_buf
.
GetDeviceBuffer
()),
static_cast
<
WeiDataType
*>
(
wei_device_buf
.
GetDeviceBuffer
()),
static_cast
<
OutDataType
*>
(
out_device_buf
.
GetDeviceBuffer
()),
param
s
.
N_
,
param
s
.
K_
,
param
s
.
C_
,
param
s
.
input_spatial_lengths_
,
param
s
.
filter_spatial_lengths_
,
param
s
.
output_spatial_lengths_
,
param
s
.
conv_filter_strides_
,
param
s
.
conv_filter_dilations_
,
param
s
.
input_left_pads_
,
param
s
.
input_right_pads_
,
conv_
param
.
N_
,
conv_
param
.
K_
,
conv_
param
.
C_
,
conv_
param
.
input_spatial_lengths_
,
conv_
param
.
filter_spatial_lengths_
,
conv_
param
.
output_spatial_lengths_
,
conv_
param
.
conv_filter_strides_
,
conv_
param
.
conv_filter_dilations_
,
conv_
param
.
input_left_pads_
,
conv_
param
.
input_right_pads_
,
in_element_op
,
wei_element_op
,
out_element_op
,
...
...
@@ -258,8 +193,8 @@ bool profile_conv_bwd_weight_impl(int do_verification,
float
avg_time
=
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
StreamConfig
{
nullptr
,
time_kernel
});
std
::
size_t
flop
=
param
s
.
GetFlops
();
std
::
size_t
num_btype
=
param
s
.
GetByte
<
InDataType
,
WeiDataType
,
OutDataType
>
();
std
::
size_t
flop
=
conv_
param
.
GetFlops
();
std
::
size_t
num_btype
=
conv_
param
.
GetByte
<
InDataType
,
WeiDataType
,
OutDataType
>
();
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
avg_time
;
float
gb_per_sec
=
num_btype
/
1.E6
/
avg_time
;
...
...
profiler/include/profile_conv_fwd_impl.hpp
View file @
0e81cc18
...
...
@@ -19,6 +19,7 @@
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_conv_fwd.hpp"
namespace
ck
{
...
...
@@ -32,83 +33,15 @@ template <ck::index_t NDimSpatial,
typename
InDataType
,
typename
WeiDataType
,
typename
OutDataType
>
int
profile_conv_fwd_impl
(
int
do_verification
,
bool
profile_conv_fwd_impl
(
int
do_verification
,
int
init_method
,
bool
do_log
,
bool
time_kernel
,
const
ck
::
tensor_operation
::
device
::
ConvParams
&
param
s
)
const
ck
::
tensor_operation
::
device
::
ConvParams
&
conv_
param
)
{
bool
pass
=
true
;
// make host tensor descritpor
auto
f_nhwc_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nhwc_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nhwc_lengths
.
insert
(
nhwc_lengths
.
begin
()
+
1
,
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nhwc_lengths
);
};
auto
f_nchw_host_tensor_descriptor
=
[](
ck
::
index_t
n
,
ck
::
index_t
c
,
std
::
vector
<
ck
::
index_t
>
spatial_lengths
)
{
std
::
vector
<
std
::
size_t
>
nchw_lengths
{
static_cast
<
std
::
size_t
>
(
n
),
static_cast
<
std
::
size_t
>
(
c
)};
nchw_lengths
.
insert
(
nchw_lengths
.
end
(),
spatial_lengths
.
begin
(),
spatial_lengths
.
end
());
return
HostTensorDescriptor
(
nchw_lengths
);
};
HostTensorDescriptor
in_desc
,
wei_desc
,
out_desc
;
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NWC
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NHWC
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWC
>
)
{
in_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
else
if
constexpr
(
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCW
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCHW
>
||
is_same_v
<
InLayout
,
ck
::
tensor_layout
::
convolution
::
NCDHW
>
)
{
in_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
C_
,
params
.
input_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KXC
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KYXC
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KZYXC
>
)
{
wei_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
else
if
constexpr
(
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCX
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCYX
>
||
is_same_v
<
WeiLayout
,
ck
::
tensor_layout
::
convolution
::
KCZYX
>
)
{
wei_desc
=
f_nchw_host_tensor_descriptor
(
params
.
K_
,
params
.
C_
,
params
.
filter_spatial_lengths_
);
}
// FIXME: properly implement "make host descriptor" for different layout
if
constexpr
(
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NWK
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NHWK
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NDHWK
>
)
{
out_desc
=
f_nhwc_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
else
if
constexpr
(
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKW
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKHW
>
||
is_same_v
<
OutLayout
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
)
{
out_desc
=
f_nchw_host_tensor_descriptor
(
params
.
N_
,
params
.
K_
,
params
.
GetOutputSpatialLengths
());
}
const
auto
in_desc
=
ck
::
utils
::
conv
::
get_input_host_tensor_descriptor
<
InLayout
>
(
conv_param
);
const
auto
wei_desc
=
ck
::
utils
::
conv
::
get_weight_host_tensor_descriptor
<
WeiLayout
>
(
conv_param
);
const
auto
out_desc
=
ck
::
utils
::
conv
::
get_output_host_tensor_descriptor
<
OutLayout
>
(
conv_param
);
Tensor
<
InDataType
>
input
(
in_desc
);
Tensor
<
WeiDataType
>
weight
(
wei_desc
);
...
...
@@ -164,10 +97,10 @@ int profile_conv_fwd_impl(int do_verification,
auto
ref_argument
=
ref_conv
.
MakeArgument
(
input
,
weight
,
host_output
,
param
s
.
conv_filter_strides_
,
param
s
.
conv_filter_dilations_
,
param
s
.
input_left_pads_
,
param
s
.
input_right_pads_
,
conv_
param
.
conv_filter_strides_
,
conv_
param
.
conv_filter_dilations_
,
conv_
param
.
input_left_pads_
,
conv_
param
.
input_right_pads_
,
in_element_op
,
wei_element_op
,
out_element_op
);
...
...
@@ -201,22 +134,24 @@ int profile_conv_fwd_impl(int do_verification,
float
best_gb_per_sec
=
0
;
// profile device op instances
bool
pass
=
true
;
for
(
auto
&
op_ptr
:
op_ptrs
)
{
auto
argument_ptr
=
op_ptr
->
MakeArgumentPointer
(
static_cast
<
InDataType
*>
(
in_device_buf
.
GetDeviceBuffer
()),
static_cast
<
WeiDataType
*>
(
wei_device_buf
.
GetDeviceBuffer
()),
static_cast
<
OutDataType
*>
(
out_device_buf
.
GetDeviceBuffer
()),
param
s
.
N_
,
param
s
.
K_
,
param
s
.
C_
,
param
s
.
input_spatial_lengths_
,
param
s
.
filter_spatial_lengths_
,
param
s
.
GetOutputSpatialLengths
(),
param
s
.
conv_filter_strides_
,
param
s
.
conv_filter_dilations_
,
param
s
.
input_left_pads_
,
param
s
.
input_right_pads_
,
conv_
param
.
N_
,
conv_
param
.
K_
,
conv_
param
.
C_
,
conv_
param
.
input_spatial_lengths_
,
conv_
param
.
filter_spatial_lengths_
,
conv_
param
.
GetOutputSpatialLengths
(),
conv_
param
.
conv_filter_strides_
,
conv_
param
.
conv_filter_dilations_
,
conv_
param
.
input_left_pads_
,
conv_
param
.
input_right_pads_
,
in_element_op
,
wei_element_op
,
out_element_op
);
...
...
@@ -233,8 +168,8 @@ int profile_conv_fwd_impl(int do_verification,
float
avg_time
=
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
StreamConfig
{
nullptr
,
time_kernel
});
std
::
size_t
flop
=
param
s
.
GetFlops
();
std
::
size_t
num_btype
=
param
s
.
GetByte
<
InDataType
,
WeiDataType
,
OutDataType
>
();
std
::
size_t
flop
=
conv_
param
.
GetFlops
();
std
::
size_t
num_btype
=
conv_
param
.
GetByte
<
InDataType
,
WeiDataType
,
OutDataType
>
();
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
avg_time
;
...
...
@@ -278,7 +213,7 @@ int profile_conv_fwd_impl(int do_verification,
<<
"
\n
name: "
<<
best_op_name
<<
"
\n
avg_time: "
<<
best_avg_time
<<
"
\n
tflops: "
<<
best_tflops
<<
"
\n
GB/s: "
<<
best_gb_per_sec
<<
std
::
endl
;
return
0
;
return
pass
;
}
}
// namespace profiler
...
...
test/conv_util/conv_util.cpp
View file @
0e81cc18
...
...
@@ -35,48 +35,11 @@ class TestConvUtil : public ::testing::Test
// stride {2,2},
// dilations {1,1},
// padding {{1,1}, {1,1}}
ck
::
utils
::
conv
::
ConvParams
conv_params
;
ck
::
tensor_operation
::
device
::
ConvParams
conv_params
;
};
}
// namespace
TEST_F
(
TestConvUtil
,
ConvParamsGetOutputSpatialLengths2D
)
{
ck
::
utils
::
conv
::
ConvParams
conv_params
;
std
::
vector
<
ck
::
index_t
>
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
36
,
36
},
"Error: ConvParams 2D default constructor."
));
conv_params
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
71
,
71
},
"Error: ConvParams 2D stride {1,1}."
));
conv_params
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
conv_params
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
conv_params
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
37
,
37
},
"Error: ConvParams 2D padding left/right {2,2}."
));
conv_params
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
36
,
36
},
"Error: ConvParams 2D dilation {2,2}."
));
conv_params
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
};
conv_params
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
conv_params
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
conv_params
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
23
,
23
},
"Error: ConvParams 2D strides{3,3}, padding {1,1}, dilations {2,2}."
));
}
TEST_F
(
TestConvUtil
,
ConvParamsGetOutputSpatialLengths1D
)
{
SetNDParams
(
1
);
...
...
@@ -114,6 +77,43 @@ TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths1D)
"Error: ConvParams 1D strides{3}, padding {1}, dilations {2}."
));
}
TEST_F
(
TestConvUtil
,
ConvParamsGetOutputSpatialLengths2D
)
{
ck
::
tensor_operation
::
device
::
ConvParams
conv_params
;
std
::
vector
<
ck
::
index_t
>
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
36
,
36
},
"Error: ConvParams 2D default constructor."
));
conv_params
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
71
,
71
},
"Error: ConvParams 2D stride {1,1}."
));
conv_params
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
conv_params
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
conv_params
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
37
,
37
},
"Error: ConvParams 2D padding left/right {2,2}."
));
conv_params
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
36
,
36
},
"Error: ConvParams 2D dilation {2,2}."
));
conv_params
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
};
conv_params
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
conv_params
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
conv_params
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
out_spatial_len
=
conv_params
.
GetOutputSpatialLengths
();
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
out_spatial_len
,
std
::
vector
<
ck
::
index_t
>
{
23
,
23
},
"Error: ConvParams 2D strides{3,3}, padding {1,1}, dilations {2,2}."
));
}
TEST_F
(
TestConvUtil
,
ConvParamsGetOutputSpatialLengths3D
)
{
SetNDParams
(
3
);
...
...
@@ -152,56 +152,3 @@ TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths3D)
std
::
vector
<
ck
::
index_t
>
{
23
,
23
,
23
},
"Error: ConvParams 3D strides{3, 3, 3}, padding {1, 1, 1}, dilations {2, 2, 2}."
));
}
TEST
(
ConvUtil
,
GetHostTensorDescriptor
)
{
namespace
tl
=
ck
::
tensor_layout
::
convolution
;
std
::
vector
<
std
::
size_t
>
dims
{
2
,
3
,
4
,
5
};
HostTensorDescriptor
h
=
ck
::
utils
::
conv
::
get_host_tensor_descriptor
(
dims
,
tl
::
NHWC
{});
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetLengths
(),
{
2
,
3
,
4
,
5
},
"Error: wrong NHWC dimensions lengths!"
));
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetStrides
(),
{
3
*
4
*
5
,
1
,
3
*
5
,
3
},
"Error: wrong NHWC dimensions strides!"
));
h
=
ck
::
utils
::
conv
::
get_host_tensor_descriptor
(
dims
,
tl
::
NCHW
{});
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetLengths
(),
{
2
,
3
,
4
,
5
},
"Error: wrong NCHW dimensions lengths!"
));
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetStrides
(),
{
3
*
4
*
5
,
4
*
5
,
5
,
1
},
"Error: wrong NCHW dimensions strides!"
));
dims
=
std
::
vector
<
std
::
size_t
>
{
2
,
3
,
4
};
h
=
ck
::
utils
::
conv
::
get_host_tensor_descriptor
(
dims
,
tl
::
NWC
{});
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetLengths
(),
{
2
,
3
,
4
},
"Error: wrong NWC dimensions lengths!"
));
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetStrides
(),
{
3
*
4
,
1
,
3
},
"Error: wrong NWC dimensions strides!"
));
h
=
ck
::
utils
::
conv
::
get_host_tensor_descriptor
(
dims
,
tl
::
NCW
{});
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetLengths
(),
{
2
,
3
,
4
},
"Error: wrong NCW dimensions lengths!"
));
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetStrides
(),
{
3
*
4
,
4
,
1
},
"Error: wrong NCW dimensions strides!"
));
dims
=
std
::
vector
<
std
::
size_t
>
{
2
,
3
,
4
,
5
,
6
};
h
=
ck
::
utils
::
conv
::
get_host_tensor_descriptor
(
dims
,
tl
::
NDHWC
{});
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetLengths
(),
dims
,
"Error: wrong NDHWC dimensions lengths!"
));
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetStrides
(),
{
3
*
4
*
5
*
6
,
// N
1
,
// C
3
*
5
*
6
,
// D
3
*
6
,
// H
3
},
// W
"Error: wrong NDHWC dimensions strides!"
));
h
=
ck
::
utils
::
conv
::
get_host_tensor_descriptor
(
dims
,
tl
::
NCDHW
{});
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetLengths
(),
dims
,
"Error: wrong NCDHW dimensions lengths!"
));
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
h
.
GetStrides
(),
{
3
*
4
*
5
*
6
,
// N
4
*
5
*
6
,
// C
5
*
6
,
// D
6
,
// H
1
},
// W
"Error: wrong NCDHW dimensions strides!"
));
}
test/reference_conv_fwd/reference_conv_fwd.cpp
View file @
0e81cc18
...
...
@@ -15,14 +15,16 @@
#include "ck/library/utility/check_err.hpp"
#include "ck/library/utility/fill.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/convolution_parameter.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_conv_fwd.hpp"
namespace
{
using
InElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
WeiElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
template
<
ck
::
index_t
NDim
,
template
<
ck
::
index_t
NDim
Spatial
,
typename
InDataType
=
float
,
typename
WeiDataType
=
float
,
typename
OutDataType
=
float
,
...
...
@@ -32,7 +34,7 @@ template <ck::index_t NDim,
typename
FillInputOp
=
ck
::
utils
::
FillMonotonicSeq
<
InDataType
>,
typename
FillWeightsOp
=
ck
::
utils
::
FillConstant
<
WeiDataType
>>
Tensor
<
OutDataType
>
run_reference_convolution_forward
(
const
ck
::
utils
::
conv
::
ConvParams
&
params
,
run_reference_convolution_forward
(
const
ck
::
tensor_operation
::
device
::
ConvParams
&
params
,
const
FillInputOp
&
fill_input_op
=
FillInputOp
{},
const
FillWeightsOp
&
fill_weights_op
=
FillWeightsOp
{
0.5
f
})
{
...
...
@@ -65,13 +67,16 @@ run_reference_convolution_forward(const ck::utils::conv::ConvParams& params,
fill_weights_op
(
weights
.
begin
(),
weights
.
end
());
std
::
fill
(
host_output
.
begin
(),
host_output
.
end
(),
OutDataType
(
0.
f
));
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvFwd
<
InDataType
,
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvFwd
<
NDimSpatial
,
InLayout
,
WeiLayout
,
OutLayout
,
InDataType
,
WeiDataType
,
OutDataType
,
InElementOp
,
WeiElementOp
,
OutElementOp
,
NDim
>
();
OutElementOp
>
();
auto
ref_invoker
=
ref_conv
.
MakeInvoker
();
auto
ref_argument
=
ref_conv
.
MakeArgument
(
input
,
weights
,
...
...
@@ -92,7 +97,7 @@ run_reference_convolution_forward(const ck::utils::conv::ConvParams& params,
TEST
(
ReferenceConvolutionFWD
,
Conv2DNHWC
)
{
ck
::
utils
::
conv
::
ConvParams
params
;
ck
::
tensor_operation
::
device
::
ConvParams
params
;
params
.
N_
=
1
;
params
.
K_
=
1
;
params
.
C_
=
2
;
...
...
@@ -128,7 +133,7 @@ TEST(ReferenceConvolutionFWD, Conv2DNHWC)
TEST
(
ReferenceConvolutionFWD
,
Conv2DNHWCStridesDilationsPadding
)
{
ck
::
utils
::
conv
::
ConvParams
params
;
ck
::
tensor_operation
::
device
::
ConvParams
params
;
params
.
N_
=
1
;
params
.
K_
=
2
;
params
.
C_
=
2
;
...
...
@@ -154,7 +159,7 @@ TEST(ReferenceConvolutionFWD, Conv2DNHWCStridesDilationsPadding)
TEST
(
ReferenceConvolutionFWD
,
Conv1DNWC
)
{
ck
::
utils
::
conv
::
ConvParams
params
;
ck
::
tensor_operation
::
device
::
ConvParams
params
;
params
.
num_dim_spatial_
=
1
;
params
.
N_
=
1
;
params
.
K_
=
1
;
...
...
@@ -183,7 +188,7 @@ TEST(ReferenceConvolutionFWD, Conv1DNWC)
TEST
(
ReferenceConvolutionFWD
,
Conv1DNWCStridesDilationsPadding
)
{
ck
::
utils
::
conv
::
ConvParams
params
;
ck
::
tensor_operation
::
device
::
ConvParams
params
;
params
.
num_dim_spatial_
=
1
;
params
.
N_
=
1
;
params
.
K_
=
2
;
...
...
@@ -212,7 +217,7 @@ TEST(ReferenceConvolutionFWD, Conv1DNWCStridesDilationsPadding)
TEST
(
ReferenceConvolutionFWD
,
Conv1DNWCSameOutputSize
)
{
ck
::
utils
::
conv
::
ConvParams
params
;
ck
::
tensor_operation
::
device
::
ConvParams
params
;
params
.
num_dim_spatial_
=
1
;
params
.
N_
=
2
;
params
.
K_
=
16
;
...
...
@@ -306,7 +311,7 @@ TEST(ReferenceConvolutionFWD, Conv1DNWCSameOutputSize)
TEST
(
ReferenceConvolutionFWD
,
Conv3DNCDHW
)
{
ck
::
utils
::
conv
::
ConvParams
params
;
ck
::
tensor_operation
::
device
::
ConvParams
params
;
params
.
num_dim_spatial_
=
3
;
params
.
N_
=
1
;
params
.
K_
=
1
;
...
...
@@ -345,7 +350,7 @@ TEST(ReferenceConvolutionFWD, Conv3DNCDHW)
TEST
(
ReferenceConvolutionFWD
,
Conv3DNCDHWStridesDilations
)
{
ck
::
utils
::
conv
::
ConvParams
params
;
ck
::
tensor_operation
::
device
::
ConvParams
params
;
params
.
num_dim_spatial_
=
3
;
params
.
N_
=
1
;
params
.
K_
=
2
;
...
...
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