clang format

src/targets/gpu/include/migraphx/gpu/miopen.hpp

@@ -49,6 +49,8 @@ inline tensor_descriptor make_tensor(const migraphx::shape& s, bool pack = false
         d = miopenFloat;
     else if(s.type() == shape::half_type)
         d = miopenHalf;
+    else if(s.type() == shape::int32_type)
+        d = miopenInt32;
     else if(s.type() == shape::int8_type)
     {
         if(pack)

src/targets/gpu/lowering.cpp

@@ -177,14 +177,10 @@ 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, tmp_output, output);
+                ins, conv, ins->inputs().at(0), ins->inputs().at(1), workspace, output);
         });
     }
 

src/targets/gpu/quant_convolution.cpp

@@ -9,7 +9,7 @@ namespace gpu {
 
 shape miopen_quant_convolution::compute_shape(const std::vector<shape>& inputs) const
 {
-    check_shapes{inputs, *this}.has(5).standard();
+    check_shapes{inputs, *this}.has(4).standard();
     return op.compute_shape({inputs.at(0), inputs.at(1)});
 }
 argument miopen_quant_convolution::compute(context& ctx,
@@ -20,8 +20,7 @@ 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);
-    shape tmp_output_shape{shape::float_type, output_shape.lens()};
-    auto y_desc = make_tensor(tmp_output_shape);
+    auto y_desc = make_tensor(output_shape);
 
     float alpha = 1;
     float beta  = 0;
@@ -70,10 +69,7 @@ argument miopen_quant_convolution::compute(context& ctx,
         MIGRAPHX_THROW("QUANT_CONVOLUTION: run convolution forward failed");
     }
 
-    // Add a conversion from float to int32_t
-    device::convert(ctx.get_stream().get(), args[4], args[3]);
-
-    return args[4];
+    return args[3];
 }
 
 shape miopen_quant_convolution::compile(context& ctx,
@@ -83,8 +79,7 @@ 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);
-    shape tmp_output_shape{shape::float_type, output_shape.lens()};
-    auto y_desc = make_tensor(tmp_output_shape);
+    auto y_desc = make_tensor(output_shape);
 
     std::size_t workspace_size = 0;
     miopenConvolutionForwardGetWorkSpaceSize(ctx.get_stream().get_miopen(),
@@ -97,7 +92,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(tmp_output_shape);
+    auto y         = allocate_gpu(output_shape);
     auto workspace = allocate_gpu(workspace_shape);
 
     int algo_count = 1;