more changes to match the gemm gpu implementation

src/include/migraphx/generate.hpp

@@ -82,8 +82,6 @@ std::vector<T> generate_tensor_data(const migraphx::shape& s, unsigned long seed
 {
     std::vector<T> result(s.elements());
     std::generate(result.begin(), result.end(), xorshf96_generator<T>{seed});
-    // divide a value to avoid integer overflow
-    std::transform(result.begin(), result.end(), result.begin(), [](auto i) { return i / 32; });
     // std::generate(result.begin(), result.end(), [&]{ return seed % 7; });
     // std::generate(result.begin(), result.end(), []{ return 1; });
     return result;

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

@@ -24,6 +24,7 @@ struct rocblas_quant_gemm
     shape compute_shape(const std::vector<shape>& inputs) const;
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+    void batch_not_transposed(const std::vector<std::size_t>& strides) const;    
     std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
     {
         return shapes.size() - 1;

src/targets/gpu/quant_gemm.cpp

@@ -56,10 +56,27 @@ shape rocblas_quant_gemm::compute_shape(const std::vector<shape>& inputs) const
     std::vector<shape> in_shapes(inputs);
     in_shapes.pop_back();
     check_shapes{in_shapes}.not_broadcasted();
+    batch_not_transposed(inputs[0].strides());
+    batch_not_transposed(inputs[1].strides());
 
     return op.compute_shape(in_shapes);
 }
 
+void rocblas_quant_gemm::batch_not_transposed(const std::vector<std::size_t>& strides) const
+{
+    if(strides.size() <= 2)
+        return;
+    auto dim_0       = strides.size() - 2;
+    auto matrix_size = std::max(strides[dim_0], strides[dim_0 + 1]);
+    std::vector<std::size_t> batch(strides.begin(), strides.begin() + dim_0);
+    if(std::adjacent_find(batch.begin(), batch.end(), [&](auto i, auto j) {
+           return (i < j or i < matrix_size or j < matrix_size);
+       }) != batch.end())
+    {
+        MIGRAPHX_THROW("DOT: batch size {" + to_string_range(strides) + "} is transposed!");
+    }
+}
+
 argument rocblas_quant_gemm::compute(context& ctx,
                                      const shape& output_shape,
                                      const std::vector<argument>& args) const