Skip to content

GitLab

Commit 870f565e authored by Shucai Xiao's avatar Shucai Xiao
Browse files

fake up to make the int8_convolution outputs the int32_t data type. With... 

fake up to make the int8_convolution outputs the int32_t data type. With MIOpen to support int32_t output shape, this tmp change will be removed.
parent dbd34685
Show whitespace changes
Inline Side-by-side
Showing with 34 additions and 23 deletions
src/include/migraphx/op/quant_convolution.hpp
View file @ 870f565e
......@@ -50,7 +50,7 @@ struct quant_convolution
{
MIGRAPHX_THROW("QUANT_CONVOLUTION: only accept input and weights of type int8_t");
}
t = shape::float_type;
t = shape::int32_type;
if(padding_mode == default_)
{
......
src/targets/cpu/lowering.cpp
View file @ 870f565e
......@@ -219,7 +219,7 @@ struct cpu_quant_convolution
argument compute(context&, shape output_shape, std::vector<argument> args) const
{
argument result{output_shape};
auto output = result.get<float>();
auto output = result.get<int32_t>();
visit_all(args[0], args[1])([&](auto input, auto weights) {
auto in = input.get_shape().lens();
auto in_h = in[2];
......@@ -240,7 +240,7 @@ struct cpu_quant_convolution
const auto start_y = j * op.stride[1] - op.padding[1];
const auto group_id = w / (wei_n / op.group);
float acc = 0;
int32_t acc = 0;
dfor(wei_c, wei_h, wei_w)([&](std::size_t k, std::size_t x, std::size_t y) {
const auto in_x = start_x + x;
const auto in_y = start_y + y;
......
src/targets/gpu/lowering.cpp
View file @ 870f565e
......@@ -167,10 +167,14 @@ struct miopen_apply
auto ws = conv.compile(ctx, ins->get_shape(), to_shapes(ins->inputs()));
auto workspace = insert_allocation(ins, ws, "workspace");
// add a temp float output to store the miopen convolution output
shape tmp_output_shape{shape::float_type, ins->get_shape().lens()};
auto tmp_output = insert_allocation(ins, tmp_output_shape, "tmp_out");
auto output = insert_allocation(ins, ins->get_shape());
return prog->replace_instruction(
ins, conv, ins->inputs().at(0), ins->inputs().at(1), workspace, output);
ins, conv, ins->inputs().at(0), ins->inputs().at(1), workspace, tmp_output, output);
});
}
......
src/targets/gpu/quant_convolution.cpp
View file @ 870f565e
#include <migraphx/gpu/quant_convolution.hpp>
#include <migraphx/gpu/device/convert.hpp>
#include <migraphx/gpu/context.hpp>
#include <migraphx/generate.hpp>
......@@ -8,7 +9,7 @@ namespace gpu {
shape miopen_quant_convolution::compute_shape(const std::vector<shape>& inputs) const
{
check_shapes{inputs, *this}.has(4).standard();
check_shapes{inputs, *this}.has(5).standard();
return op.compute_shape({inputs.at(0), inputs.at(1)});
}
argument miopen_quant_convolution::compute(context& ctx,
......@@ -19,7 +20,8 @@ argument miopen_quant_convolution::compute(context& ctx,
auto x_desc_vec4 = make_tensor(args[0].get_shape(), true);
auto w_desc = make_tensor(args[1].get_shape());
auto w_desc_vec4 = make_tensor(args[1].get_shape(), true);
auto y_desc = make_tensor(output_shape);
shape tmp_output_shape{shape::float_type, output_shape.lens()};
auto y_desc = make_tensor(tmp_output_shape);
float alpha = 1;
float beta = 0;
......@@ -67,7 +69,11 @@ argument miopen_quant_convolution::compute(context& ctx,
{
MIGRAPHX_THROW("QUANT_CONVOLUTION: run convolution forward failed");
}
return args[3];
// Add a conversion from float to int32_t
device::convert(ctx.get_stream().get(), args[4], args[3]);
return args[4];
}
shape miopen_quant_convolution::compile(context& ctx,
......@@ -77,7 +83,8 @@ shape miopen_quant_convolution::compile(context& ctx,
shape workspace_shape{};
auto x_desc = make_tensor(inputs[0], true);
auto w_desc = make_tensor(inputs[1], true);
auto y_desc = make_tensor(output_shape);
shape tmp_output_shape{shape::float_type, output_shape.lens()};
auto y_desc = make_tensor(tmp_output_shape);
std::size_t workspace_size = 0;
miopenConvolutionForwardGetWorkSpaceSize(ctx.get_stream().get_miopen(),
......@@ -90,7 +97,7 @@ shape miopen_quant_convolution::compile(context& ctx,
arg_vec4_x = to_gpu(generate_argument(pack_int8_shape(inputs[0])));
arg_vec4_w = to_gpu(generate_argument(pack_int8_shape(inputs[1])));
auto y = allocate_gpu(output_shape);
auto y = allocate_gpu(tmp_output_shape);
auto workspace = allocate_gpu(workspace_shape);
int algo_count = 1;
......
test/cpu_ops_test.cpp
View file @ 870f565e
......@@ -1317,7 +1317,7 @@ TEST_CASE(quant_conv2d_test)
p.compile(migraphx::cpu::target{});
auto result = p.eval({});
std::vector<float> s = {10197,
std::vector<int32_t> s = {10197,
10548,
11601,
11952,
......@@ -1334,7 +1334,7 @@ TEST_CASE(quant_conv2d_test)
81666,
82746};
std::vector<float> results_vector;
std::vector<int32_t> results_vector;
result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
EXPECT(migraphx::verify_range(results_vector, s));
}
......@@ -1357,14 +1357,14 @@ TEST_CASE(quant_conv2d_test_default_mode)
p.compile(migraphx::cpu::target{});
auto result = p.eval({});
std::vector<float> s = {
std::vector<int32_t> s = {
10197, 10548, 6939, 3420, 11601, 11952, 7839, 3852, 7383, 7590, 4953, 2421, 3480,
3570, 2316, 1125, 25506, 26586, 17874, 9009, 29826, 30906, 20718, 10413, 20505, 21198,
14187, 7119, 10527, 10860, 7257, 3636, 27045, 27396, 17739, 8604, 28449, 28800, 18639,
9036, 17319, 17526, 11289, 5445, 7800, 7890, 5052, 2421, 77346, 78426, 52002, 25857,
81666, 82746, 54846, 27261, 53769, 54462, 36075, 17919, 26511, 26844, 17769, 8820};
std::vector<float> results_vector;
std::vector<int32_t> results_vector;
result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
EXPECT(migraphx::verify_range(results_vector, s));
}
......@@ -1387,7 +1387,7 @@ TEST_CASE(quant_conv2d_test_valid_mode)
p.compile(migraphx::cpu::target{});
auto result = p.eval({});
std::vector<float> s = {10197,
std::vector<int32_t> s = {10197,
10548,
11601,
11952,
......@@ -1404,7 +1404,7 @@ TEST_CASE(quant_conv2d_test_valid_mode)
81666,
82746};
std::vector<float> results_vector;
std::vector<int32_t> results_vector;
result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
EXPECT(migraphx::verify_range(results_vector, s));
}
......@@ -1423,14 +1423,14 @@ TEST_CASE(quant_conv2d_padding_test)
p.add_instruction(migraphx::op::quant_convolution{{{1, 1}}, {{1, 1}}}, al, cl);
p.compile(migraphx::cpu::target{});
auto result = p.eval({});
std::vector<float> s = {
std::vector<int32_t> s = {
4521, 6753, 7014, 4635, 6858, 10197, 10548, 6939, 7830, 11601, 11952, 7839, 5007,
7383, 7590, 4953, 10515, 15987, 16734, 11277, 16821, 25506, 26586, 17874, 19737, 29826,
30906, 20718, 13593, 20505, 21198, 14187, 13161, 19281, 19542, 12699, 18522, 27045, 27396,
17739, 19494, 28449, 28800, 18639, 11919, 17319, 17526, 11289, 34707, 51843, 52590, 34893,
51813, 77346, 78426, 52002, 54729, 81666, 82746, 54846, 36057, 53769, 54462, 36075};
std::vector<float> results_vector;
std::vector<int32_t> results_vector;
result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
EXPECT(migraphx::verify_range(results_vector, s));
}
......@@ -1450,7 +1450,7 @@ TEST_CASE(quant_conv2d_padding_stride_test)
p.compile(migraphx::cpu::target{});
auto result = p.eval({});
std::vector<float> s = {4521,
std::vector<int32_t> s = {4521,
7014,
7830,
11952,
......@@ -1466,7 +1466,7 @@ TEST_CASE(quant_conv2d_padding_stride_test)
52590,
54729,
82746};
std::vector<float> results_vector;
std::vector<int32_t> results_vector;
result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
EXPECT(migraphx::verify_range(results_vector, s));
}
......
test/op_shape_test.cpp
View file @ 870f565e
......@@ -78,24 +78,24 @@ TEST_CASE(convolution_shape)
TEST_CASE(quant_convolution_shape)
{
migraphx::shape output{migraphx::shape::float_type, {4, 4, 1, 1}};
migraphx::shape output{migraphx::shape::int32_type, {4, 4, 1, 1}};
migraphx::shape input{migraphx::shape::int8_type, {4, 3, 3, 3}};
migraphx::shape weights{migraphx::shape::int8_type, {4, 3, 3, 3}};
expect_shape(output, migraphx::op::quant_convolution{}, input, weights);
throws_shape(migraphx::op::quant_convolution{}, input);
migraphx::shape input2{migraphx::shape::float_type, {3, 3}};
migraphx::shape input2{migraphx::shape::int32_type, {3, 3}};
migraphx::shape weights2{migraphx::shape::float_type, {3, 3}};
throws_shape(migraphx::op::quant_convolution{}, input2, weights2);
throws_shape(migraphx::op::quant_convolution{}, input2, weights);
migraphx::shape input3{migraphx::shape::float_type, {4, 3, 3, 3}};
migraphx::shape input3{migraphx::shape::int32_type, {4, 3, 3, 3}};
migraphx::shape weight3{migraphx::shape::float_type, {4, 3, 3, 3}};
throws_shape(migraphx::op::quant_convolution{}, input3, weights);
throws_shape(migraphx::op::quant_convolution{}, input, weight3);
throws_shape(migraphx::op::quant_convolution{}, input3, weight3);
migraphx::shape output_same_mode{migraphx::shape::float_type, {4, 4, 3, 3}};
migraphx::shape output_same_mode{migraphx::shape::int32_type, {4, 4, 3, 3}};
expect_shape(output_same_mode,
migraphx::op::quant_convolution{{{0, 0}}, {{1, 1}}, {{1, 1}}, migraphx::op::same},
input,
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment