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
8e8ae66d
Commit
8e8ae66d
authored
Jul 18, 2022
by
Chao Liu
Browse files
clean
parent
0e81cc18
Changes
27
Show whitespace changes
Inline
Side-by-side
Showing
7 changed files
with
121 additions
and
139 deletions
+121
-139
profiler/src/profile_conv_bwd_weight.cpp
profiler/src/profile_conv_bwd_weight.cpp
+11
-12
profiler/src/profile_conv_fwd.cpp
profiler/src/profile_conv_fwd.cpp
+11
-12
test/conv_util/conv_util.cpp
test/conv_util/conv_util.cpp
+2
-2
test/convnd_bwd_data/convnd_bwd_data.cpp
test/convnd_bwd_data/convnd_bwd_data.cpp
+1
-1
test/convnd_bwd_weight/convnd_bwd_weight.cpp
test/convnd_bwd_weight/convnd_bwd_weight.cpp
+1
-1
test/convnd_fwd/convnd_fwd.cpp
test/convnd_fwd/convnd_fwd.cpp
+1
-1
test/reference_conv_fwd/reference_conv_fwd.cpp
test/reference_conv_fwd/reference_conv_fwd.cpp
+94
-110
No files found.
profiler/src/profile_conv_bwd_weight.cpp
View file @
8e8ae66d
...
...
@@ -49,8 +49,7 @@ static void print_helper_msg()
<<
std
::
endl
;
}
ck
::
tensor_operation
::
device
::
ConvParams
parse_conv_params
(
int
num_dim_spatial
,
int
arg_idx
,
char
*
const
argv
[])
ck
::
utils
::
conv
::
ConvParam
parse_conv_params
(
int
num_dim_spatial
,
int
arg_idx
,
char
*
const
argv
[])
{
const
ck
::
index_t
N
=
std
::
stoi
(
argv
[
arg_idx
++
]);
const
ck
::
index_t
K
=
std
::
stoi
(
argv
[
arg_idx
++
]);
...
...
@@ -93,7 +92,7 @@ parse_conv_params(int num_dim_spatial, int arg_idx, char* const argv[])
input_right_pads
[
i
]
=
std
::
stoi
(
argv
[
arg_idx
++
]);
}
return
ck
::
tensor_operation
::
device
::
ConvParam
s
{
num_dim_spatial
,
return
ck
::
utils
::
conv
::
ConvParam
{
num_dim_spatial
,
N
,
K
,
C
,
...
...
profiler/src/profile_conv_fwd.cpp
View file @
8e8ae66d
...
...
@@ -51,8 +51,7 @@ static void print_helper_msg()
<<
std
::
endl
;
}
ck
::
tensor_operation
::
device
::
ConvParams
parse_conv_params
(
int
num_dim_spatial
,
int
arg_idx
,
char
*
const
argv
[])
ck
::
utils
::
conv
::
ConvParam
parse_conv_params
(
int
num_dim_spatial
,
int
arg_idx
,
char
*
const
argv
[])
{
const
ck
::
index_t
N
=
std
::
stoi
(
argv
[
arg_idx
++
]);
const
ck
::
index_t
K
=
std
::
stoi
(
argv
[
arg_idx
++
]);
...
...
@@ -95,7 +94,7 @@ parse_conv_params(int num_dim_spatial, int arg_idx, char* const argv[])
input_right_pads
[
i
]
=
std
::
stoi
(
argv
[
arg_idx
++
]);
}
return
ck
::
tensor_operation
::
device
::
ConvParam
s
{
num_dim_spatial
,
return
ck
::
utils
::
conv
::
ConvParam
{
num_dim_spatial
,
N
,
K
,
C
,
...
...
test/conv_util/conv_util.cpp
View file @
8e8ae66d
...
...
@@ -35,7 +35,7 @@ class TestConvUtil : public ::testing::Test
// stride {2,2},
// dilations {1,1},
// padding {{1,1}, {1,1}}
ck
::
tensor_operation
::
device
::
ConvParam
s
conv_params
;
ck
::
utils
::
conv
::
ConvParam
conv_params
;
};
}
// namespace
...
...
@@ -79,7 +79,7 @@ TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths1D)
TEST_F
(
TestConvUtil
,
ConvParamsGetOutputSpatialLengths2D
)
{
ck
::
tensor_operation
::
device
::
ConvParam
s
conv_params
;
ck
::
utils
::
conv
::
ConvParam
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
},
...
...
test/convnd_bwd_data/convnd_bwd_data.cpp
View file @
8e8ae66d
...
...
@@ -13,7 +13,7 @@ int main()
{
bool
pass
=
true
;
std
::
vector
<
ck
::
tensor_operation
::
device
::
ConvParam
s
>
params
;
std
::
vector
<
ck
::
utils
::
conv
::
ConvParam
>
params
;
// check 1d
params
.
push_back
({
1
,
128
,
128
,
256
,
{
1
},
{
14
},
{
2
},
{
1
},
{
0
},
{
0
}});
...
...
test/convnd_bwd_weight/convnd_bwd_weight.cpp
View file @
8e8ae66d
...
...
@@ -13,7 +13,7 @@ int main()
{
bool
pass
=
true
;
std
::
vector
<
ck
::
tensor_operation
::
device
::
ConvParam
s
>
params
;
std
::
vector
<
ck
::
utils
::
conv
::
ConvParam
>
params
;
// check 1d
params
.
push_back
({
1
,
128
,
256
,
256
,
{
1
},
{
7
},
{
2
},
{
1
},
{
0
},
{
0
}});
...
...
test/convnd_fwd/convnd_fwd.cpp
View file @
8e8ae66d
...
...
@@ -13,7 +13,7 @@ int main()
{
bool
pass
=
true
;
std
::
vector
<
ck
::
tensor_operation
::
device
::
ConvParam
s
>
params
;
std
::
vector
<
ck
::
utils
::
conv
::
ConvParam
>
params
;
// check 1d
params
.
push_back
({
1
,
128
,
128
,
256
,
{
1
},
{
14
},
{
2
},
{
1
},
{
0
},
{
0
}});
...
...
test/reference_conv_fwd/reference_conv_fwd.cpp
View file @
8e8ae66d
...
...
@@ -16,6 +16,7 @@
#include "ck/library/utility/fill.hpp"
#include "ck/library/utility/host_tensor.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
{
...
...
@@ -34,34 +35,17 @@ template <ck::index_t NDimSpatial,
typename
FillInputOp
=
ck
::
utils
::
FillMonotonicSeq
<
InDataType
>,
typename
FillWeightsOp
=
ck
::
utils
::
FillConstant
<
WeiDataType
>>
Tensor
<
OutDataType
>
run_reference_convolution_forward
(
const
ck
::
tensor_operation
::
device
::
ConvParam
s
&
param
s
,
run_reference_convolution_forward
(
const
ck
::
utils
::
conv
::
ConvParam
&
conv_
param
,
const
FillInputOp
&
fill_input_op
=
FillInputOp
{},
const
FillWeightsOp
&
fill_weights_op
=
FillWeightsOp
{
0.5
f
})
{
std
::
vector
<
std
::
size_t
>
input_dims
{
static_cast
<
std
::
size_t
>
(
params
.
N_
),
static_cast
<
std
::
size_t
>
(
params
.
C_
)};
input_dims
.
insert
(
std
::
end
(
input_dims
),
std
::
begin
(
params
.
input_spatial_lengths_
),
std
::
end
(
params
.
input_spatial_lengths_
));
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
);
std
::
vector
<
std
::
size_t
>
filter_dims
{
static_cast
<
std
::
size_t
>
(
params
.
K_
),
static_cast
<
std
::
size_t
>
(
params
.
C_
)};
filter_dims
.
insert
(
std
::
end
(
filter_dims
),
std
::
begin
(
params
.
filter_spatial_lengths_
),
std
::
end
(
params
.
filter_spatial_lengths_
));
const
std
::
vector
<
ck
::
index_t
>&
output_spatial_lengths
=
params
.
GetOutputSpatialLengths
();
std
::
vector
<
std
::
size_t
>
output_dims
{
static_cast
<
std
::
size_t
>
(
params
.
N_
),
static_cast
<
std
::
size_t
>
(
params
.
K_
)};
output_dims
.
insert
(
std
::
end
(
output_dims
),
std
::
begin
(
output_spatial_lengths
),
std
::
end
(
output_spatial_lengths
));
Tensor
<
InDataType
>
input
(
ck
::
utils
::
conv
::
get_host_tensor_descriptor
(
input_dims
,
InLayout
{}));
Tensor
<
WeiDataType
>
weights
(
ck
::
utils
::
conv
::
get_host_tensor_descriptor
(
filter_dims
,
WeiLayout
{}));
Tensor
<
OutDataType
>
host_output
(
ck
::
utils
::
conv
::
get_host_tensor_descriptor
(
output_dims
,
OutLayout
{}));
Tensor
<
InDataType
>
input
(
in_desc
);
Tensor
<
WeiDataType
>
weights
(
wei_desc
);
Tensor
<
OutDataType
>
host_output
(
out_desc
);
fill_input_op
(
input
.
begin
(),
input
.
end
());
fill_weights_op
(
weights
.
begin
(),
weights
.
end
());
...
...
@@ -81,10 +65,10 @@ run_reference_convolution_forward(const ck::tensor_operation::device::ConvParams
auto
ref_argument
=
ref_conv
.
MakeArgument
(
input
,
weights
,
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_
,
InElementOp
{},
WeiElementOp
{},
OutElementOp
{});
...
...
@@ -97,18 +81,18 @@ run_reference_convolution_forward(const ck::tensor_operation::device::ConvParams
TEST
(
ReferenceConvolutionFWD
,
Conv2DNHWC
)
{
ck
::
tensor_operation
::
device
::
ConvParam
s
param
s
;
param
s
.
N_
=
1
;
param
s
.
K_
=
1
;
param
s
.
C_
=
2
;
param
s
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
};
param
s
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
6
,
6
};
param
s
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
param
s
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
param
s
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
};
param
s
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
};
ck
::
utils
::
conv
::
ConvParam
conv_
param
;
conv_
param
.
N_
=
1
;
conv_
param
.
K_
=
1
;
conv_
param
.
C_
=
2
;
conv_
param
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
};
conv_
param
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
6
,
6
};
conv_
param
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
conv_
param
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
conv_
param
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
};
conv_
param
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
};
auto
out_tensor
=
run_reference_convolution_forward
<
2
>
(
param
s
);
auto
out_tensor
=
run_reference_convolution_forward
<
2
>
(
conv_
param
);
std
::
vector
<
std
::
size_t
>
ref_dims
{
1
,
1
,
4
,
4
};
std
::
vector
<
float
>
ref_data
{
130.5
,
148.5
,
...
...
@@ -133,18 +117,18 @@ TEST(ReferenceConvolutionFWD, Conv2DNHWC)
TEST
(
ReferenceConvolutionFWD
,
Conv2DNHWCStridesDilationsPadding
)
{
ck
::
tensor_operation
::
device
::
ConvParam
s
param
s
;
param
s
.
N_
=
1
;
param
s
.
K_
=
2
;
param
s
.
C_
=
2
;
param
s
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
};
param
s
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
12
,
12
};
param
s
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
param
s
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
param
s
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
param
s
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
ck
::
utils
::
conv
::
ConvParam
conv_
param
;
conv_
param
.
N_
=
1
;
conv_
param
.
K_
=
2
;
conv_
param
.
C_
=
2
;
conv_
param
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
};
conv_
param
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
12
,
12
};
conv_
param
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
conv_
param
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
2
,
2
};
conv_
param
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
conv_
param
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
};
auto
out_tensor
=
run_reference_convolution_forward
<
2
>
(
param
s
);
auto
out_tensor
=
run_reference_convolution_forward
<
2
>
(
conv_
param
);
std
::
vector
<
std
::
size_t
>
ref_dims
=
std
::
vector
<
std
::
size_t
>
{
1
,
2
,
5
,
5
};
std
::
vector
<
float
>
ref_data
{
210.
,
210.
,
327.
,
327.
,
351.
,
351.
,
375.
,
375.
,
399.
,
399.
,
...
...
@@ -159,17 +143,17 @@ TEST(ReferenceConvolutionFWD, Conv2DNHWCStridesDilationsPadding)
TEST
(
ReferenceConvolutionFWD
,
Conv1DNWC
)
{
ck
::
tensor_operation
::
device
::
ConvParam
s
param
s
;
param
s
.
num_dim_spatial_
=
1
;
param
s
.
N_
=
1
;
param
s
.
K_
=
1
;
param
s
.
C_
=
2
;
param
s
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
};
param
s
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
6
};
param
s
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
param
s
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
param
s
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
};
param
s
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
};
ck
::
utils
::
conv
::
ConvParam
conv_
param
;
conv_
param
.
num_dim_spatial_
=
1
;
conv_
param
.
N_
=
1
;
conv_
param
.
K_
=
1
;
conv_
param
.
C_
=
2
;
conv_
param
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
};
conv_
param
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
6
};
conv_
param
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
conv_
param
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
conv_
param
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
};
conv_
param
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
};
auto
out_tensor
=
run_reference_convolution_forward
<
1
,
...
...
@@ -178,7 +162,7 @@ TEST(ReferenceConvolutionFWD, Conv1DNWC)
float
,
ck
::
tensor_layout
::
convolution
::
NWC
,
ck
::
tensor_layout
::
convolution
::
KXC
,
ck
::
tensor_layout
::
convolution
::
NWK
>
(
param
s
);
ck
::
tensor_layout
::
convolution
::
NWK
>
(
conv_
param
);
std
::
vector
<
std
::
size_t
>
ref_dims
{
1
,
1
,
4
};
std
::
vector
<
float
>
ref_data
{
7.5
,
13.5
,
19.5
,
25.5
};
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
...
...
@@ -188,17 +172,17 @@ TEST(ReferenceConvolutionFWD, Conv1DNWC)
TEST
(
ReferenceConvolutionFWD
,
Conv1DNWCStridesDilationsPadding
)
{
ck
::
tensor_operation
::
device
::
ConvParam
s
param
s
;
param
s
.
num_dim_spatial_
=
1
;
param
s
.
N_
=
1
;
param
s
.
K_
=
2
;
param
s
.
C_
=
2
;
param
s
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
};
param
s
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
12
};
param
s
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
2
};
param
s
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
2
};
param
s
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
param
s
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
ck
::
utils
::
conv
::
ConvParam
conv_
param
;
conv_
param
.
num_dim_spatial_
=
1
;
conv_
param
.
N_
=
1
;
conv_
param
.
K_
=
2
;
conv_
param
.
C_
=
2
;
conv_
param
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
};
conv_
param
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
12
};
conv_
param
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
2
};
conv_
param
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
2
};
conv_
param
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
conv_
param
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
auto
out_tensor
=
run_reference_convolution_forward
<
1
,
...
...
@@ -207,7 +191,7 @@ TEST(ReferenceConvolutionFWD, Conv1DNWCStridesDilationsPadding)
float
,
ck
::
tensor_layout
::
convolution
::
NWC
,
ck
::
tensor_layout
::
convolution
::
KXC
,
ck
::
tensor_layout
::
convolution
::
NWK
>
(
param
s
);
ck
::
tensor_layout
::
convolution
::
NWK
>
(
conv_
param
);
std
::
vector
<
std
::
size_t
>
ref_dims
{
1
,
2
,
5
};
std
::
vector
<
float
>
ref_data
{
9.
,
9.
,
19.5
,
19.5
,
31.5
,
31.5
,
43.5
,
43.5
,
55.5
,
55.5
};
EXPECT_TRUE
(
ck
::
utils
::
check_err
(
...
...
@@ -217,17 +201,17 @@ TEST(ReferenceConvolutionFWD, Conv1DNWCStridesDilationsPadding)
TEST
(
ReferenceConvolutionFWD
,
Conv1DNWCSameOutputSize
)
{
ck
::
tensor_operation
::
device
::
ConvParam
s
param
s
;
param
s
.
num_dim_spatial_
=
1
;
param
s
.
N_
=
2
;
param
s
.
K_
=
16
;
param
s
.
C_
=
4
;
param
s
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
};
param
s
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
16
};
param
s
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
param
s
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
param
s
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
param
s
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
ck
::
utils
::
conv
::
ConvParam
conv_
param
;
conv_
param
.
num_dim_spatial_
=
1
;
conv_
param
.
N_
=
2
;
conv_
param
.
K_
=
16
;
conv_
param
.
C_
=
4
;
conv_
param
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
};
conv_
param
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
16
};
conv_
param
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
conv_
param
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
conv_
param
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
conv_
param
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
1
};
auto
out_tensor2
=
run_reference_convolution_forward
<
1
,
float
,
...
...
@@ -236,7 +220,7 @@ TEST(ReferenceConvolutionFWD, Conv1DNWCSameOutputSize)
ck
::
tensor_layout
::
convolution
::
NWC
,
ck
::
tensor_layout
::
convolution
::
KXC
,
ck
::
tensor_layout
::
convolution
::
NWK
>
(
param
s
,
ck
::
utils
::
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
conv_
param
,
ck
::
utils
::
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
std
::
vector
<
std
::
size_t
>
ref_dims
{
2
,
16
,
16
};
std
::
vector
<
float
>
ref_data
{
...
...
@@ -311,17 +295,17 @@ TEST(ReferenceConvolutionFWD, Conv1DNWCSameOutputSize)
TEST
(
ReferenceConvolutionFWD
,
Conv3DNCDHW
)
{
ck
::
tensor_operation
::
device
::
ConvParam
s
param
s
;
param
s
.
num_dim_spatial_
=
3
;
param
s
.
N_
=
1
;
param
s
.
K_
=
1
;
param
s
.
C_
=
2
;
param
s
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
,
3
};
param
s
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
6
,
6
,
6
};
param
s
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
,
1
};
param
s
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
,
1
};
param
s
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
,
0
};
param
s
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
,
0
};
ck
::
utils
::
conv
::
ConvParam
conv_
param
;
conv_
param
.
num_dim_spatial_
=
3
;
conv_
param
.
N_
=
1
;
conv_
param
.
K_
=
1
;
conv_
param
.
C_
=
2
;
conv_
param
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
,
3
};
conv_
param
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
6
,
6
,
6
};
conv_
param
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
,
1
};
conv_
param
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
,
1
};
conv_
param
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
,
0
};
conv_
param
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
,
0
};
auto
out_tensor
=
run_reference_convolution_forward
<
3
,
float
,
...
...
@@ -330,7 +314,7 @@ TEST(ReferenceConvolutionFWD, Conv3DNCDHW)
ck
::
tensor_layout
::
convolution
::
NCDHW
,
ck
::
tensor_layout
::
convolution
::
KCZYX
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
(
param
s
,
ck
::
utils
::
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
conv_
param
,
ck
::
utils
::
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
std
::
vector
<
std
::
size_t
>
ref_dims
{
1
,
1
,
4
,
4
,
4
};
std
::
vector
<
float
>
ref_data
{
407.7
,
410.40002
,
413.09998
,
415.80002
,
423.90002
,
426.6
,
429.30002
,
432.
,
...
...
@@ -350,17 +334,17 @@ TEST(ReferenceConvolutionFWD, Conv3DNCDHW)
TEST
(
ReferenceConvolutionFWD
,
Conv3DNCDHWStridesDilations
)
{
ck
::
tensor_operation
::
device
::
ConvParam
s
param
s
;
param
s
.
num_dim_spatial_
=
3
;
param
s
.
N_
=
1
;
param
s
.
K_
=
2
;
param
s
.
C_
=
2
;
param
s
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
,
3
};
param
s
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
12
,
12
,
12
};
param
s
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
,
3
};
param
s
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
,
1
};
param
s
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
,
0
};
param
s
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
,
0
};
ck
::
utils
::
conv
::
ConvParam
conv_
param
;
conv_
param
.
num_dim_spatial_
=
3
;
conv_
param
.
N_
=
1
;
conv_
param
.
K_
=
2
;
conv_
param
.
C_
=
2
;
conv_
param
.
filter_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
,
3
};
conv_
param
.
input_spatial_lengths_
=
std
::
vector
<
ck
::
index_t
>
{
12
,
12
,
12
};
conv_
param
.
conv_filter_strides_
=
std
::
vector
<
ck
::
index_t
>
{
3
,
3
,
3
};
conv_
param
.
conv_filter_dilations_
=
std
::
vector
<
ck
::
index_t
>
{
1
,
1
,
1
};
conv_
param
.
input_left_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
,
0
};
conv_
param
.
input_right_pads_
=
std
::
vector
<
ck
::
index_t
>
{
0
,
0
,
0
};
auto
out_tensor
=
run_reference_convolution_forward
<
3
,
float
,
...
...
@@ -369,7 +353,7 @@ TEST(ReferenceConvolutionFWD, Conv3DNCDHWStridesDilations)
ck
::
tensor_layout
::
convolution
::
NCDHW
,
ck
::
tensor_layout
::
convolution
::
KCZYX
,
ck
::
tensor_layout
::
convolution
::
NKDHW
>
(
param
s
,
ck
::
utils
::
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
conv_
param
,
ck
::
utils
::
FillMonotonicSeq
<
float
>
{
0.
f
,
0.1
f
});
std
::
vector
<
std
::
size_t
>
ref_dims
{
1
,
2
,
4
,
4
,
4
};
std
::
vector
<
float
>
ref_data
{
2756.7002
,
2764.7998
,
2772.9001
,
2781.
,
2853.9001
,
2862.
,
2870.1
,
2878.2002
,
...
...
Prev
1
2
Next
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