fixed a race condition problem in quant_gemm.

src/targets/cpu/lowering.cpp

@@ -9,8 +9,6 @@
 #include <migraphx/cpu/gemm.hpp>
 #include <unordered_map>
 #include <utility>
-#include <fstream>
-#include <iomanip>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
@@ -565,8 +563,7 @@ struct cpu_quant_gemm
         }
 
         // 2 input arguments
-        int32_t beta = 0;
-        migemm(result, arg_0, arg_1, op.alpha, beta);
+        migemm(result, arg_0, arg_1, op.alpha, int32_t{0});
 
         return result;
     }

src/targets/gpu/device/pack.cpp

@@ -69,6 +69,11 @@ void pack_b(hipStream_t stream, const argument& result, const argument& arg)
     });
 }
 
+void sync_stream(hipStream_t stream)
+{
+    hipStreamSynchronize(stream);
+}
+
 } // namespace device
 } // namespace gpu
 } // namespace MIGRAPHX_INLINE_NS

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

@@ -13,6 +13,9 @@ namespace device {
 void pack_a(hipStream_t stream, const argument& result, const argument& arg);
 
 void pack_b(hipStream_t stream, const argument& result, const argument& arg);
+
+void sync_stream(hipStream_t stream);
+
 } // namespace device
 } // namespace gpu
 } // namespace MIGRAPHX_INLINE_NS

src/targets/gpu/quant_gemm.cpp

@@ -2,6 +2,8 @@
 #include <migraphx/gpu/device/pack.hpp>
 #include <migraphx/gpu/context.hpp>
 #include <migraphx/generate.hpp>
+#include <fstream>
+#include <iomanip>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
@@ -86,6 +88,7 @@ argument miopen_quant_gemm::compute(context& ctx,
     {
         device::pack_b(ctx.get_stream().get(), args[arg_num - 3], args[0]);
     }
+    device::sync_stream(ctx.get_stream().get());
 
     bool is_3inputs = (arg_num == 6);
     int32_t beta    = 0;