Repository: alibaba/BladeDISC (language: python)

Issue:
Initial quantization support in BladeDISC
We're now planing to start support quantization in BladeDISC. The basic idea is:
- introduce a fake_quant op. The op will be used to pass quantization related info (e.g. scale/zero_point) from frontend to BladeDISC.
- The input of BladeDISC is still a FP32 graph, with some fake_quant ops to annotate which set of ops need to be quantized.
-  It's the responsibility of BladeDISC to replace the fake_quant op and the corresponding FP32 ops (e.g. Gemm/Conv) to the quantized versions. 

The task list is shown as follow:
- [x] TensorFlow/PyTorch fake_quant custom ops and their corresponding implementation in TF/PyTorch Dialect.
- [x] mhlo_disc fake_quantize/quantize/de-quantize IR definition.
- [x] mhlo_disc QCONV/QMATMUL IR defintion.
- [x] TF custom fake-quant -> disc fake quant
- [ ] Torch custom fake-quant -> disc fake quant
- [x] Fake_quantize + Conv + Fake_quantize  -> QConv pattern.
- [x] Fake_quantize + GEMM + Fake_quantize  -> QGEMM pattern.
- [x] Fake_quantize lowering.
- [x] Quantize op lowering patterns.
- [x] de-quantize op lowering patterns.
- [x] Placement rule for Qconv/Qmatmul.
- [ ] (low priority) QConv/QMatmul layout optimization.
- [x] QConv shape inference implementation.
- [x] QMatmul shape inference implementation.
- [x] QConv lowering to library call implementation.
- [x] QMatmul lowering to library call implementation.
- [x] QConv kernel and e2e UT for arm.
- [x] QMatMul kernel and e2e UT for arm.
- [x] QConv kernel and e2e UT for cuda.
- [x] QMatMul kernel and e2e UT for cuda.
- [ ] QConv kernel and e2e UT for x86.
- [ ] QMatMul kernel and e2e UT for x86.
- [ ] e2e demo model for ARM
- [ ] e2e demo model for X86
- [ ] e2e demo model for CUDA



Discussion / hints:

Fix patch:
diff --git a/tao_compiler/mlir/disc/IR/lhlo_disc_ops.cc b/tao_compiler/mlir/disc/IR/lhlo_disc_ops.cc
index a187a57562c..43d915a372f 100644
--- a/tao_compiler/mlir/disc/IR/lhlo_disc_ops.cc
+++ b/tao_compiler/mlir/disc/IR/lhlo_disc_ops.cc
@@ -39,8 +39,8 @@ LmhloDiscDialect::LmhloDiscDialect(MLIRContext* context)
 
 LogicalResult CustomCallOp::verify() {
   CustomCallOp op = *this;
-  if (op.target_arg_mapping()) {
-    lmhlo::CustomCallTargetArgMappingAttr mapping = *op.target_arg_mapping();
+  if (op.getTargetArgMapping()) {
+    lmhlo::CustomCallTargetArgMappingAttr mapping = *op.getTargetArgMapping();
     auto verify_mapping = [&](int64_t target_num, size_t op_num,
                               ::llvm::ArrayRef<int64_t> mapping,
                               StringRef kind) -> LogicalResult {
@@ -73,9 +73,9 @@ LogicalResult CustomCallOp::verify() {
       }
       return success();
     };
-    if (failed(verify_mapping(mapping.getNumArgs(), op.args().size(),
+    if (failed(verify_mapping(mapping.getNumArgs(), op.getArgs().size(),
                               mapping.getArgsToTargetArgs(), "args")) ||
-        failed(verify_mapping(mapping.getNumResults(), op.output().size(),
+        failed(verify_mapping(mapping.getNumResults(), op.getOutput().size(),
                               mapping.getResultsToTargetResults(), "results")))
       return failure();
   }
diff --git a/tao_compiler/mlir/disc/IR/lhlo_disc_ops.td b/tao_compiler/mlir/disc/IR/lhlo_disc_ops.td
index de57261ac38..44748969dea 100644
--- a/tao_compiler/mlir/disc/IR/lhlo_disc_ops.td
+++ b/tao_compiler/mlir/disc/IR/lhlo_disc_ops.td
@@ -28,6 +28,7 @@ include "mlir-hlo/Dialect/lhlo/IR/lhlo_structured_interface.td"
 def LHLODISC_Dialect : Dialect {
   let name = "lmhlo_disc";
   let cppNamespace = "::mlir::lmhlo_disc";
+  let emitAccessorPrefix = kEmitAccessorPrefix_Prefixed;
 }
 
 def LHLO_I32Tensor : MemRefOf<[SignlessIntOfWidths<[32]>]>;
diff --git a/tao_compiler/mlir/disc/IR/rng_uniform_custom_call_op.cc b/tao_compiler/mlir/disc/IR/rng_uniform_custom_call_op.cc
index 61f6f9aca55..450f6d7f6ed 100644
--- a/tao_compiler/mlir/disc/IR/rng_uniform_custom_call_op.cc
+++ b/tao_compiler/mlir/disc/IR/rng_uniform_custom_call_op.cc
@@ -99,7 +99,7 @@ LogicalResult lowerToLibraryCallImpl<RngUniformBackendConfig>(
   // seed & seed2
   llvm::Expected<RngUniformBackendConfig> backend_config =
       llvm::json::parse<RngUniformBackendConfig>(
-          op.backend_config().cast<StringAttr>());
+          op.getBackendConfig().cast<StringAttr>());
   if (auto e = backend_config.takeError()) {
     return op.emitOpError()
            << "Problem with parsing rng_uniform backend_config: "
diff --git a/tao_compiler/mlir/disc/IR/sparse_custom_call_op.cc b/tao_compiler/mlir/disc/IR/sparse_custom_call_op.cc
index b06b5d1efbe..31075ba519c 100644
--- a/tao_compiler/mlir/disc/IR/sparse_custom_call_op.cc
+++ b/tao_compiler/mlir/disc/IR/sparse_custom_call_op.cc
@@ -76,7 +76,7 @@ LogicalResult lowerToLibraryCallSpargeGemmImpl(CustomCallOp op,
   newOperands.push_back(op.getOperand(2));
 
   Location loc = op.getLoc();
-  auto config = op.backend_config().cast<DictionaryAttr>();
+  auto config = op.getBackendConfig().cast<DictionaryAttr>();
   int64_t lhs_contracting_dimensions =
       config.getAs<IntegerAttr>("lhs_contracting_dimensions").getInt();
   int64_t rhs_contracting_dimensions =
diff --git a/tao_compiler/mlir/disc/IR/topk_custom_call_op.cc b/tao_compiler/mlir/disc/IR/topk_custom_call_op.cc
index ee552de391d..445df2532d0 100644
--- a/tao_compiler/mlir/disc/IR/topk_custom_call_op.cc
+++ b/tao_compiler/mlir/disc/IR/topk_custom_call_op.cc
@@ -134,7 +134,7 @@ LogicalResult lowerToLibraryCallImpl<TopKBackendConfig>(
   // dimension
   llvm::Expected<TopKBackendConfig> backend_config =
       llvm::json::parse<TopKBackendConfig>(
-          op.backend_config().cast<StringAttr>());
+          op.getBackendConfig().cast<StringAttr>());
   if (auto e = backend_config.takeError()) {
     return op.emitOpError() << "Problem with parsing topk backend_config: "
                             << llvm::toString(std::move(e));
diff --git a/tao_compiler/mlir/disc/transforms/disc_lower_to_library_call.cc b/tao_compiler/mlir/disc/transforms/disc_lower_to_library_call.cc
index cab9d74282d..254f26f52b7 100644
--- a/tao_compiler/mlir/disc/transforms/disc_lower_to_library_call.cc
+++ b/tao_compiler/mlir/disc/transforms/disc_lower_to_library_call.cc
@@ -65,6 +65,7 @@ using lmhlo::ReshapeOp;
 using lmhlo_disc::CustomCallOp;
 using lmhlo_disc::D2HOp;
 using lmhlo_disc::H2DOp;
+using lmhlo_disc::QuantizedDynamicConvOp;
 
 // Suppose that the first argument of the function is the ctx value
 Value GetContextValueFromFunctionArguments(Operation* op) {
@@ -262,7 +263,11 @@ Value GetConvMetadata(OpTy op, PatternRewriter& rewriter) {
   //   - weight_is_const : indicate whether the weight is const.
   Location loc = op.getLoc();
   Type field_type = rewriter.getI32Type();
-  int rank = op.getOutput().getType().template dyn_cast<ShapedType>().getRank();
+  int numOperands = op->getNumOperands();
+  int rank = op->getOperand(numOperands - 1)
+                 .getType()
+                 .template dyn_cast<ShapedType>()
+                 .getRank();
   int num_spatial_dims = rank - 2;
   int num_metadata_fields = rank * 3 + (rank - 2) * 2 + 1;
   Value metadata_value = rewriter.create<memref::AllocaOp>(
@@ -364,15 +369,19 @@ struct ConvConverter : public OpRewritePattern<ConvolutionOp> {
   }
 };
 
-struct DynamicConvConverter : public OpRewritePattern<DynamicConvOp> {
-  using OpRewritePattern<DynamicConvOp>::OpRewritePattern;
+template <typename OpTy>
+struct DynamicConvLikeConverter : public OpRewritePattern<OpTy> {
+  DynamicConvLikeConverter(MLIRContext* context, StringRef target)
+      : OpRewritePattern<OpTy>::OpRewritePattern(context) {
+    this->target_ = target;
+  }
 
-  LogicalResult matchAndRewrite(DynamicConvOp op,
+  LogicalResult matchAndRewrite(OpTy op,
                                 PatternRewriter& rewriter) const override {
     Value ctx = GetContextValueFromFunctionArguments(op);
     if (!ctx) {
-      op->emitOpError()
-          << "the first argument of the function is not ral context type.";
+      return op->emitOpError()
+             << "the first argument of the function is not ral context type.";
     }
     Value stream_handle = GetDefaultStreamHandle(op, rewriter);
     SmallVector<Value, 4> newOperands{stream_handle};
@@ -384,12 +393,15 @@ struct DynamicConvConverter : public OpRewritePattern<DynamicConvOp> {
     // input & kernel & output layouts, strides and dilation
     newOperands.push_back(GetConvMetadata(op, rewriter));
 
-    bool on_gpu = placement_utils::isGpuMemRef(op->getOperand(3));
+    bool on_gpu = placement_utils::isGpuMemRef(op->getOperand(0));
     rewriter.replaceOpWithNewOp<DispatchOp>(op, llvm::None, ctx, newOperands,
-                                            "ral_conv", false,
+                                            target_, false,
                                             on_gpu ? "gpu" : "cpu");
     return success();
   }
+
+ private:
+  StringRef target_;
 };
 
 // Return null is not a copy-removable copy-like ops, otherwise return
@@ -575,7 +587,7 @@ struct CustomCallOpConvertor : public OpRewritePattern<CustomCallOp> {
     }
     Value stream_handle = GetDefaultStreamHandle(op, rewriter);
 
-    StringRef target_name = op.call_target_name();
+    StringRef target_name = op.getCallTargetName();
     auto lower_func =
         mhlo_disc::CustomCallRegistry::Global().FindLowerToLibraryCallFunc(
             target_name.str());
@@ -610,7 +622,6 @@ struct DiscLowerToLibraryCallPass
       DotGeneralOpConvertor,
       RecvInputOpConvertor,
       ConvConverter,
-      DynamicConvConverter,
       SendOutputOpConvertor
     >(context);
     // clang-format on
@@ -618,6 +629,10 @@ struct DiscLowerToLibraryCallPass
     // GPU copy related ops
     patterns.insert<GpuCopyOpConvertor<H2DOp>>(context, "h2d");
     patterns.insert<GpuCopyOpConvertor<D2HOp>>(context, "d2h");
+    patterns.insert<DynamicConvLikeConverter<DynamicConvOp>>(context,
+                                                             "ral_conv");
+    patterns.insert<DynamicConvLikeConverter<QuantizedDynamicConvOp>>(
+        context, "ral_qconv");
 
     if (failed(applyPatternsAndFoldGreedily(func, std::move(patterns)))) {
       func.emitError("applyPatternsAndFoldGreedily does not converge");
diff --git a/tao_compiler/mlir/disc/transforms/disc_to_llvm.cc b/tao_compiler/mlir/disc/transforms/disc_to_llvm.cc
index 82a4e9d87fe..7ef7e382569 100644
--- a/tao_compiler/mlir/disc/transforms/disc_to_llvm.cc
+++ b/tao_compiler/mlir/disc/transforms/disc_to_llvm.cc
@@ -1133,9 +1133,9 @@ LogicalResult PrintfToLLVMPattern::matchAndRewrite(
 
   // Get a symbol reference to the printf function, inserting it if necessary.
   auto printfRef = getOrInsertPrintf(rewriter, parentModule);
-  std::string key = llvm::Twine("DiscPrintf", op.format()).str();
+  std::string key = llvm::Twine("DiscPrintf", op.getFormat()).str();
   Value formatSpecCst =
-      getOrCreateGlobalString(loc, rewriter, key, op.format(), parentModule);
+      getOrCreateGlobalString(loc, rewriter, key, op.getFormat(), parentModule);
   SmallVector<Value, 4> val_to_print{formatSpecCst};
   for (Value operand : adaptor.getOperands()) {
     val_to_print.push_back(operand);
diff --git a/tao_compiler/mlir/disc/transforms/lhlo_legalize_roots_to_loops.cc b/tao_compiler/mlir/disc/transforms/lhlo_legalize_roots_to_loops.cc
index e751b80b4e2..98fa83fbc22 100644
--- a/tao_compiler/mlir/disc/transforms/lhlo_legalize_roots_to_loops.cc
+++ b/tao_compiler/mlir/disc/transforms/lhlo_legalize_roots_to_loops.cc
@@ -3694,11 +3694,11 @@ LogicalResult lowerWithScheduleSparseReshapeOpCPU(
   Value zero = b.create<arith::ConstantIndexOp>(loc, 0);
   Value one = b.create<arith::ConstantIndexOp>(loc, 1);
   Value num_values =
-      b.create<memref::DimOp>(loc, sparse_reshape_op.input_indices(), 0);
+      b.create<memref::DimOp>(loc, sparse_reshape_op.getInputIndices(), 0);
   Value origin_rank =
-      b.create<memref::DimOp>(loc, sparse_reshape_op.input_shape(), 0);
+      b.create<memref::DimOp>(loc, sparse_reshape_op.getInputShape(), 0);
   Value new_rank =
-      b.create<memref::DimOp>(loc, sparse_reshape_op.new_shape(), 0);
+      b.create<memref::DimOp>(loc, sparse_reshape_op.getNewShape(), 0);
 
   Value shape_accum;
   // TODO: How to emit a reverse loop?
@@ -3716,7 +3716,7 @@ LogicalResult lowerWithScheduleSparseReshapeOpCPU(
 
     Value reverse_end = b.create<arith::SubIOp>(loc, new_rank, one);
     Value init_value = b.create<memref::LoadOp>(
-        loc, sparse_reshape_op.new_shape(), reverse_end);
+        loc, sparse_reshape_op.getNewShape(), reverse_end);
     SmallVector<Value, 1> init_values({init_value});
     auto for_op = b.create<scf::ForOp>(loc, /* lowerBound */ one,
                                        /* upperBound */ new_rank,
@@ -3731,7 +3731,7 @@ LogicalResult lowerWithScheduleSparseReshapeOpCPU(
     b.create<memref::StoreOp>(loc, acc_init, shape_accum, i_reverse);
     Value yield_value = b.create<arith::MulIOp>(
         loc, acc_init,
-        b.create<memref::LoadOp>(loc, sparse_reshape_op.new_shape(),
+        b.create<memref::LoadOp>(loc, sparse_reshape_op.getNewShape(),
                                  i_reverse));
     b.create<scf::YieldOp>(loc, ValueRange({yield_value}));
 
@@ -3758,7 +3758,7 @@ LogicalResult lowerWithScheduleSparseReshapeOpCPU(
     // SparseReshapeOp
     SmallVector<Value, 2> load_index({n, zero});
     Value init_value = b.create<memref::LoadOp>(
-        loc, sparse_reshape_op.input_indices(), load_index);
+        loc, sparse_reshape_op.getInputIndices(), load_index);
 
     auto for_op =
         b.create<scf::ForOp>(loc, /* lowerBound */ one,
@@ -3773,9 +3773,9 @@ LogicalResult lowerWithScheduleSparseReshapeOpCPU(
     // multidim_index[n][i]
     load_index[1] = i;
     Value index_i = b.create<memref::LoadOp>(
-        loc, sparse_reshape_op.input_indices(), load_index);
+        loc, sparse_reshape_op.getInputIndices(), load_index);
     Value shape_i =
-        b.create<memref::LoadOp>(loc, sparse_reshape_op.input_shape(), i);
+        b.create<memref::LoadOp>(loc, sparse_reshape_op.getInputShape(), i);
 
     Value yield_value = b.create<arith::AddIOp>(
         loc, b.create<arith::MulIOp>(loc, carry_value, shape_i), index_i);
@@ -3807,13 +3807,13 @@ LogicalResult lowerWithScheduleSparseReshapeOpCPU(
     Value yield_value = b.create<arith::RemUIOp>(loc, carry_value, divisor);
 
     b.create<memref::StoreOp>(loc, multidim_index,
-                              sparse_reshape_op.output_indices(),
+                              sparse_reshape_op.getOutputIndices(),
                               ValueRange({n, i}));
     b.create<scf::YieldOp>(loc, yield_value);
     b.setInsertionPointToEnd(for_op_n.getBody());
 
     b.create<memref::StoreOp>(loc, for_op.getResult(0),
-                              sparse_reshape_op.output_indices(),
+                              sparse_reshape_op.getOutputIndices(),
                               ValueRange({n, reverse_end}));
   }
 
@@ -3832,9 +3832,9 @@ LogicalResult lowerWithScheduleSparseReshapeOpCPU(
 
     Value i = for_op_out_shape.getInductionVar();
     auto shape_val =
-        b.create<memref::LoadOp>(loc, sparse_reshape_op.new_shape(), i);
-    b.create<memref::StoreOp>(loc, shape_val, sparse_reshape_op.output_shape(),
-                              ValueRange({i}));
+        b.create<memref::LoadOp>(loc, sparse_reshape_op.getNewShape(), i);
+    b.create<memref::StoreOp>(
+        loc, shape_val, sparse_reshape_op.getOutputShape(), ValueRange({i}));
     b.create<scf::YieldOp>(loc, ValueRange({}));
     b.setInsertionPointAfter(for_op_n.getOperation());
   }
@@ -3877,14 +3877,14 @@ LogicalResult lowerWithScheduleSparseFillEmptyRowsOpCPU(
       loc, b.getIntegerAttr(b.getIntegerType(64), 1));
 
   // input
-  Value indices = sparse_fill_empty_rows_op.indices();
-  Value values = sparse_fill_empty_rows_op.values();
-  Value dense_shape = sparse_fill_empty_rows_op.dense_shape();
+  Value indices = sparse_fill_empty_rows_op.getIndices();
+  Value values = sparse_fill_empty_rows_op.getValues();
+  Value dense_shape = sparse_fill_empty_rows_op.getDenseShape();
   // output
-  Value output_indices = sparse_fill_empty_rows_op.output_indices();
-  Value output_values = sparse_fill_empty_rows_op.output_values();
-  Value empty_row_indicator = sparse_fill_empty_rows_op.empty_row_indicator();
-  Value reverse_index_map = sparse_fill_empty_rows_op.reverse_index_map();
+  Value output_indices = sparse_fill_empty_rows_op.getOutputIndices();
+  Value output_values = sparse_fill_empty_rows_op.getOutputValues();
+  Value empty_row_indicator = sparse_fill_empty_rows_op.getEmptyRowIndicator();
+  Value reverse_index_map = sparse_fill_empty_rows_op.getReverseIndexMap();
 
   Value num_rows = b.create<arith::IndexCastOp>(
       loc, b.getIndexType(), b.create<memref::LoadOp>(loc, dense_shape, zero));
@@ -3997,7 +3997,7 @@ LogicalResult lowerWithScheduleSparseFillEmptyRowsOpCPU(
   // store N_full
   b.create<memref::StoreOp>(
       loc, b.create<arith::IndexCastOp>(loc, b.getIntegerType(64), N_full),
-      sparse_fill_empty_rows_op.output_elements(), zero);
+      sparse_fill_empty_rows_op.getOutputElements(), zero);
 
   // row_start_idx = 0
   // for (i = 0; i < num_rows; ++i) {
@@ -4042,7 +4042,7 @@ LogicalResult lowerWithScheduleSparseFillEmptyRowsOpCPU(
     b.create<memref::StoreOp>(loc, zero_int, output_indices, store_index);
     // output_values
     auto default_value = b.create<memref::LoadOp>(
-        loc, sparse_fill_empty_rows_op.default_value());
+        loc, sparse_fill_empty_rows_op.getDefaultValue());
     b.create<memref::StoreOp>(loc, default_value, output_values, row_start_idx);
     b.create<scf::YieldOp>(loc, one);
 
diff --git a/tao_compiler/mlir/disc/transforms/placement_utils.cc b/tao_compiler/mlir/disc/transforms/placement_utils.cc
index 15a564e3e4b..b9fb52b72ea 100644
--- a/tao_compiler/mlir/disc/transforms/placement_utils.cc
+++ b/tao_compiler/mlir/disc/transforms/placement_utils.cc
@@ -109,7 +109,7 @@ ShapeOperandList getShapeCalcOperandList(Operation* op) {
       target_name = custom_call.call_target_name().str();
     } else {
       auto custom_call = cast<lmhlo_disc::CustomCallOp>(op);
-      target_name = custom_call.call_target_name().str();
+      target_name = custom_call.getCallTargetName().str();
     }
     auto iter = kCustomCallShapeOperandListMap.find(target_name);
     if (iter != kCustomCallShapeOperandListMap.end()) {


Test patch:
diff --git a/tao_compiler/mlir/disc/transforms/tests/disc-lower-to-library-call.mlir b/tao_compiler/mlir/disc/transforms/tests/disc-lower-to-library-call.mlir
index 1cdf79e6887..219bd0071fd 100644
--- a/tao_compiler/mlir/disc/transforms/tests/disc-lower-to-library-call.mlir
+++ b/tao_compiler/mlir/disc/transforms/tests/disc-lower-to-library-call.mlir
@@ -444,3 +444,45 @@ func.func @dynamic_conv(%arg0: !disc_ral.context) {
   "disc_ral.send_output"(%arg0, %c0_2, %69) : (!disc_ral.context, index, memref<?x?x?x?xf32, "gpu">) -> ()
   return
 }
+
+// -----
+
+// CHECK-LABEL: quantized_dynamic_conv
+func.func @quantized_dynamic_conv(%arg0: !disc_ral.context) {
+  %c0 = arith.constant 0 : index
+  %0 = "disc_ral.recv_input"(%arg0, %c0) : (!disc_ral.context, index) -> memref<?x?x?x?xi8, "gpu">
+  %c1 = arith.constant 1 : index
+  %1 = "disc_ral.recv_input"(%arg0, %c1) : (!disc_ral.context, index) -> memref<?x?x?x?xi8, "gpu">
+  %c2 = arith.constant 2 : index
+  %2 = "disc_ral.recv_input"(%arg0, %c2) : (!disc_ral.context, index) -> memref<4xi32, "cpu">
+  %c3 = arith.constant 3 : index
+  %3 = "disc_ral.recv_input"(%arg0, %c3) : (!disc_ral.context, index) -> memref<f32, "gpu">
+  %c4 = arith.constant 4 : index
+  %4 = "disc_ral.recv_input"(%arg0, %c4) : (!disc_ral.context, index) -> memref<i32, "gpu">
+  %c5 = arith.constant 5 : index
+  %5 = "disc_ral.recv_input"(%arg0, %c5) : (!disc_ral.context, index) -> memref<f32, "gpu">
+  %c6 = arith.constant 6 : index
+  %6 = "disc_ral.recv_input"(%arg0, %c6) : (!disc_ral.context, index) -> memref<i32, "gpu">
+  %c7 = arith.constant 7 : index
+  %7 = "disc_ral.recv_input"(%arg0, %c7) : (!disc_ral.context, index) -> memref<f32, "gpu">
+  %c8 = arith.constant 8 : index
+  %8 = "disc_ral.recv_input"(%arg0, %c8) : (!disc_ral.context, index) -> memref<i32, "gpu">
+  %c9 = arith.constant 9 : index
+  %9 = "disc_ral.recv_input"(%arg0, %c9) : (!disc_ral.context, index) -> memref<?x?x?x?xi8, "gpu">
+  // CHECK: memref.store {{.*}}, %[[PADDING:.*]][%c16] : memref<17xi32, "cpu">
+  // CHECK: disc_ral.dispatch
+  // CHECK-SAME: backend_config = "gpu"
+  // CHECK-SAME: call_target_name = "ral_qconv"
+  // CHECK-NOT: lmhlo_disc.quantized_dynamic_conv
+  "lmhlo_disc.quantized_dynamic_conv"(%0, %1, %2, %3, %4, %5, %6, %7, %8, %9) {
+    axis = dense<> : tensor<0xi64>,
+    batch_group_count = 1 : i64,
+    dimension_numbers = #mhlo.conv<[b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1]>,
+    feature_group_count = 1 : i64,
+    rhs_dilation = dense<1> : tensor<2xi64>,
+    use_dynamic = false,
+    use_symmetric = true,
+    window_strides = dense<3> : tensor<2xi64>
+  } : (memref<?x?x?x?xi8, "gpu">, memref<?x?x?x?xi8, "gpu">, memref<4xi32, "cpu">, memref<f32, "gpu">, memref<i32, "gpu">, memref<f32, "gpu">, memref<i32, "gpu">, memref<f32, "gpu">, memref<i32, "gpu">, memref<?x?x?x?xi8, "gpu">) -> ()
+  return
+}