Merge branch 'int8_miopen_call' into int8_quantize

src/include/migraphx/op/quant_convolution.hpp

@@ -95,7 +95,7 @@ struct quant_convolution
         }
         else
         {
-            MIGRAPHX_THROW("Invalid padding mode");
+            MIGRAPHX_THROW("QUANT_CONVOLUTION: invalid padding mode");
         }
     }
 };

src/targets/gpu/device/include/migraphx/gpu/device/tensor.hpp

@@ -61,23 +61,19 @@ struct hip_tensor_descriptor
     {
         std::copy(s.lens().begin(), s.lens().end(), lens);
         std::copy(s.strides().begin(), s.strides().end(), strides);
-        std::vector<std::size_t> vec_idx(s.lens().size());
-        std::iota(vec_idx.begin(), vec_idx.end(), 0);
-        std::sort(vec_idx.begin(), vec_idx.end(), [&](size_t i, size_t j) {
-            return strides[i] > strides[j];
-        });
-        std::copy(vec_idx.begin(), vec_idx.end(), indices);
     }
 
     __device__ __host__ hip_index<NDim> multi(size_t idx) const
     {
         hip_index<NDim> result{};
         size_t tidx = idx;
+
         for(size_t is = 0; is < NDim; is++)
         {
-            result[indices[is]] = tidx / strides[indices[is]];
-            tidx                = tidx % strides[indices[is]];
+            result[is] = tidx / strides[is];
+            tidx       = tidx % strides[is];
         }
+
         return result;
     }
 
@@ -90,7 +86,6 @@ struct hip_tensor_descriptor
     }
     size_t lens[NDim]    = {};
     size_t strides[NDim] = {};
-    size_t indices[NDim] = {};
 };
 
 } // namespace device

src/targets/gpu/device/pack.cpp

@@ -40,11 +40,15 @@ void pack_a(hipStream_t stream, const argument& result, const argument& arg)
 
 void pack_b(hipStream_t stream, const argument& result, const argument& arg)
 {
-    auto output_shape  = result.get_shape();
-    auto out_lens      = output_shape.lens();
-    auto dim_0         = output_shape.lens().size() - 2;
-    auto dim_1         = output_shape.lens().size() - 1;
-    std::size_t ldb    = output_shape.strides()[dim_1];
+    auto trans_shape = result.get_shape();
+    auto out_lens    = trans_shape.lens();
+    auto dim_0       = trans_shape.lens().size() - 2;
+    auto dim_1       = trans_shape.lens().size() - 1;
+    std::size_t ldb  = trans_shape.strides()[dim_1];
+
+    auto wrap_lens = out_lens;
+    std::swap(wrap_lens[dim_0], wrap_lens[dim_1]);
+    shape output_shape{trans_shape.type(), wrap_lens};
     std::size_t m_size = out_lens[dim_0] * out_lens[dim_1];
     visit_all(result, arg)([&](auto output, auto input) {
         std::size_t nelements = output_shape.elements();
@@ -55,8 +59,8 @@ void pack_b(hipStream_t stream, const argument& result, const argument& arg)
             gs_launch(stream, nelements)([=](auto ii) {
                 const size_t nb    = 4;
                 auto idx           = desc.multi(ii);
-                std::size_t i_n    = idx[dim_0];
-                std::size_t i_k    = idx[dim_1];
+                std::size_t i_n    = idx[dim_1];
+                std::size_t i_k    = idx[dim_0];
                 std::size_t offset = ii / m_size * m_size;
                 out_ptr[i_k % nb + (i_n + (i_k / nb) * ldb) * nb + offset] =
                     in_ptr[i_n + i_k * ldb + offset];

src/targets/gpu/quant_convolution.cpp

@@ -99,7 +99,7 @@ shape miopen_quant_convolution::compile(context& ctx,
                                                         x_desc.get(),
                                                         arg_vec4_x.implicit(),
                                                         w_desc.get(),
-                                                        arg_vec4_x.implicit(),
+                                                        arg_vec4_w.implicit(),
                                                         cd.get(),
                                                         y_desc.get(),
                                                         y.implicit(),

test/gpu/miopen.cpp

@@ -1263,11 +1263,12 @@ struct quant_dot_3args_2 : verify_program<quant_dot_3args_2>
         migraphx::shape m1_shape{migraphx::shape::int8_type, {8, 2}};
         migraphx::shape m2_shape{migraphx::shape::int8_type, {8, 7}};
         migraphx::shape m3_shape{migraphx::shape::int32_type, {2, 7}};
+        std::vector<int> m3_data(2 * 7, 1);
 
         auto l1  = p.add_parameter("a", m1_shape);
         auto tl1 = p.add_instruction(migraphx::op::transpose{{1, 0}}, l1);
         auto l2  = p.add_parameter("b", m2_shape);
-        auto l3  = p.add_parameter("c", m3_shape);
+        auto l3  = p.add_literal(m3_shape, m3_data);
         p.add_instruction(migraphx::op::quant_dot{1, 3}, tl1, l2, l3);
         return p;
     }

test/op_shape_test.cpp

@@ -104,11 +104,6 @@ TEST_CASE(quant_convolution_shape)
                  migraphx::op::quant_convolution{{{0, 0}}, {{1, 1}}, {{1, 1}}, migraphx::op::valid},
                  input,
                  weights);
-    throws_shape(
-        migraphx::op::quant_convolution{
-            {{0, 0}}, {{1, 1}}, {{1, 1}}, migraphx::op::padding_mode_t(9999)},
-        input,
-        weights);
 }
 
 TEST_CASE(transpose_shape)