Interpreter Error

57 views
Skip to first unread message

Axel Wong

unread,
Jun 2, 2024, 9:50:40 PMJun 2
to OpenXLA Discuss
Hi everyone,

I've encountered a runtime error with the following IR. There is no error during IR module verifying.

module {
  func.func @main(%arg0: tensor<1x3x32x32xf32>, %arg1: tensor<6x3x5x5xf32>, %arg2: tensor<6xf32>, %arg3: tensor<16x6x5x5xf32>, %arg4: tensor<16xf32>, %arg5: tensor<120x400xf32>, %arg6: tensor<120xf32>, %arg7: tensor<84x120xf32>, %arg8: tensor<84xf32>, %arg9: tensor<10x84xf32>, %arg10: tensor<10xf32>) -> tensor<1x10xf32> {
    %0 = stablehlo.uniform_quantize %arg0 : (tensor<1x3x32x32xf32>) -> tensor<1x3x32x32x!quant.uniform<i8:f32, 0.029222087934613228>>
    %1 = stablehlo.uniform_quantize %arg1 : (tensor<6x3x5x5xf32>) -> tensor<6x3x5x5x!quant.uniform<i8:f32, 0.0039700339548289776>>
    %2 = stablehlo.uniform_quantize %arg2 : (tensor<6xf32>) -> tensor<6x!quant.uniform<i8:f32, 0.0021052095107734203>>
    %3 = stablehlo.convolution(%0, %1) dim_numbers = [b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1], window = {stride = [1, 1], pad = [[0, 0], [0, 0]], lhs_dilate = [1, 1], rhs_dilate = [1, 1]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x3x32x32x!quant.uniform<i8:f32, 0.029222087934613228>>, tensor<6x3x5x5x!quant.uniform<i8:f32, 0.0039700339548289776>>) -> tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>
    %4 = stablehlo.broadcast_in_dim %2, dims = [1] : (tensor<6x!quant.uniform<i8:f32, 0.0021052095107734203>>) -> tensor<1x6x28x28x!quant.uniform<i8:f32, 0.0021052095107734203>>
    %5 = stablehlo.add %3, %4 : (tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>, tensor<1x6x28x28x!quant.uniform<i8:f32, 0.0021052095107734203>>) -> tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>
    %cst = stablehlo.constant dense<0.000000e+00> : tensor<1x6x28x28xf32>
    %6 = stablehlo.uniform_quantize %cst : (tensor<1x6x28x28xf32>) -> tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>
    %7 = stablehlo.maximum %5, %6 : tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>
    %cst_0 = stablehlo.constant dense<0xFF800000> : tensor<f32>
    %8 = stablehlo.uniform_quantize %cst_0 : (tensor<f32>) -> tensor<!quant.uniform<i8:f32, 0.053985215723514557>>
    %9 = "stablehlo.reduce_window"(%7, %8) <{base_dilations = array<i64: 1, 1, 1, 1>, padding = dense<0> : tensor<4x2xi64>, window_dilations = array<i64: 1, 1, 1, 1>, window_dimensions = array<i64: 1, 1, 2, 2>, window_strides = array<i64: 1, 1, 2, 2>}> ({
    ^bb0(%arg11: tensor<!quant.uniform<i8:f32, 0.053985215723514557>>, %arg12: tensor<!quant.uniform<i8:f32, 0.053985215723514557>>):
      %37 = stablehlo.maximum %arg11, %arg12 : tensor<!quant.uniform<i8:f32, 0.053985215723514557>>
      stablehlo.return %37 : tensor<!quant.uniform<i8:f32, 0.053985215723514557>>
    }) : (tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>, tensor<!quant.uniform<i8:f32, 0.053985215723514557>>) -> tensor<1x6x14x14x!quant.uniform<i8:f32, 0.04167218878865242>>
    %10 = stablehlo.uniform_quantize %arg3 : (tensor<16x6x5x5xf32>) -> tensor<16x6x5x5x!quant.uniform<i8:f32, 0.016297046095132828>>
    %11 = stablehlo.uniform_quantize %arg4 : (tensor<16xf32>) -> tensor<16x!quant.uniform<i8:f32, 0.004838157445192337>>
    %12 = stablehlo.convolution(%9, %10) dim_numbers = [b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1], window = {stride = [1, 1], pad = [[0, 0], [0, 0]], lhs_dilate = [1, 1], rhs_dilate = [1, 1]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x6x14x14x!quant.uniform<i8:f32, 0.04167218878865242>>, tensor<16x6x5x5x!quant.uniform<i8:f32, 0.016297046095132828>>) -> tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>
    %13 = stablehlo.broadcast_in_dim %11, dims = [1] : (tensor<16x!quant.uniform<i8:f32, 0.004838157445192337>>) -> tensor<1x16x10x10x!quant.uniform<i8:f32, 0.004838157445192337>>
    %14 = stablehlo.add %12, %13 : (tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>, tensor<1x16x10x10x!quant.uniform<i8:f32, 0.004838157445192337>>) -> tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>
    %cst_1 = stablehlo.constant dense<0.000000e+00> : tensor<1x16x10x10xf32>
    %15 = stablehlo.uniform_quantize %cst_1 : (tensor<1x16x10x10xf32>) -> tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>
    %16 = stablehlo.maximum %14, %15 : tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>
    %cst_2 = stablehlo.constant dense<0xFF800000> : tensor<f32>
    %17 = stablehlo.uniform_quantize %cst_2 : (tensor<f32>) -> tensor<!quant.uniform<i8:f32, 0.20967985689640045>>
    %18 = "stablehlo.reduce_window"(%16, %17) <{base_dilations = array<i64: 1, 1, 1, 1>, padding = dense<0> : tensor<4x2xi64>, window_dilations = array<i64: 1, 1, 1, 1>, window_dimensions = array<i64: 1, 1, 2, 2>, window_strides = array<i64: 1, 1, 2, 2>}> ({
    ^bb0(%arg11: tensor<!quant.uniform<i8:f32, 0.20967985689640045>>, %arg12: tensor<!quant.uniform<i8:f32, 0.20967985689640045>>):
      %37 = stablehlo.maximum %arg11, %arg12 : tensor<!quant.uniform<i8:f32, 0.20967985689640045>>
      stablehlo.return %37 : tensor<!quant.uniform<i8:f32, 0.20967985689640045>>
    }) : (tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>, tensor<!quant.uniform<i8:f32, 0.20967985689640045>>) -> tensor<1x16x5x5x!quant.uniform<i8:f32, 0.11828774958848953>>
    %19 = stablehlo.reshape %18 : (tensor<1x16x5x5x!quant.uniform<i8:f32, 0.11828774958848953>>) -> tensor<1x400x!quant.uniform<i8:f32, 0.11828774958848953>>
    %20 = stablehlo.uniform_quantize %arg5 : (tensor<120x400xf32>) -> tensor<120x400x!quant.uniform<i8:f32, 0.0099556939676403999>>
    %21 = stablehlo.uniform_quantize %arg6 : (tensor<120xf32>) -> tensor<120x!quant.uniform<i8:f32, 0.0091858189553022385>>
    %22 = stablehlo.transpose %20, dims = [1, 0] : (tensor<120x400x!quant.uniform<i8:f32, 0.0099556939676403999>>) -> tensor<400x120x!quant.uniform<i8:f32, 0.0099556939676403999>>
    %23 = stablehlo.dot_general %19, %22, contracting_dims = [1] x [0] : (tensor<1x400x!quant.uniform<i8:f32, 0.11828774958848953>>, tensor<400x120x!quant.uniform<i8:f32, 0.0099556939676403999>>) -> tensor<1x120x!quant.uniform<i8:f32, 0.46472719311714172>>
    %cst_3 = stablehlo.constant dense<0.000000e+00> : tensor<1x120xf32>
    %24 = stablehlo.uniform_quantize %cst_3 : (tensor<1x120xf32>) -> tensor<1x120x!quant.uniform<i8:f32, 0.46472719311714172>>
    %25 = stablehlo.maximum %23, %24 : tensor<1x120x!quant.uniform<i8:f32, 0.46472719311714172>>
    %26 = stablehlo.uniform_quantize %arg7 : (tensor<84x120xf32>) -> tensor<84x120x!quant.uniform<i8:f32, 0.011763181537389755>>
    %27 = stablehlo.uniform_quantize %arg8 : (tensor<84xf32>) -> tensor<84x!quant.uniform<i8:f32, 0.008067096583545208>>
    %28 = stablehlo.transpose %26, dims = [1, 0] : (tensor<84x120x!quant.uniform<i8:f32, 0.011763181537389755>>) -> tensor<120x84x!quant.uniform<i8:f32, 0.011763181537389755>>
    %29 = stablehlo.dot_general %25, %28, contracting_dims = [1] x [0] : (tensor<1x120x!quant.uniform<i8:f32, 0.46472719311714172>>, tensor<120x84x!quant.uniform<i8:f32, 0.011763181537389755>>) -> tensor<1x84x!quant.uniform<i8:f32, 0.40483295917510986>>
    %cst_4 = stablehlo.constant dense<0.000000e+00> : tensor<1x84xf32>
    %30 = stablehlo.uniform_quantize %cst_4 : (tensor<1x84xf32>) -> tensor<1x84x!quant.uniform<i8:f32, 0.40483295917510986>>
    %31 = stablehlo.maximum %29, %30 : tensor<1x84x!quant.uniform<i8:f32, 0.40483295917510986>>
    %32 = stablehlo.uniform_quantize %arg9 : (tensor<10x84xf32>) -> tensor<10x84x!quant.uniform<i8:f32, 0.013255398720502853>>
    %33 = stablehlo.uniform_quantize %arg10 : (tensor<10xf32>) -> tensor<10x!quant.uniform<i8:f32, 0.0026047099381685257>>
    %34 = stablehlo.transpose %32, dims = [1, 0] : (tensor<10x84x!quant.uniform<i8:f32, 0.013255398720502853>>) -> tensor<84x10x!quant.uniform<i8:f32, 0.013255398720502853>>
    %35 = stablehlo.dot_general %31, %34, contracting_dims = [1] x [0] : (tensor<1x84x!quant.uniform<i8:f32, 0.40483295917510986>>, tensor<84x10x!quant.uniform<i8:f32, 0.013255398720502853>>) -> tensor<1x10x!quant.uniform<i8:f32, 0.18991459906101227>>
    %36 = stablehlo.uniform_dequantize %35 : (tensor<1x10x!quant.uniform<i8:f32, 0.18991459906101227>>) -> tensor<1x10xf32>
    return %36 : tensor<1x10xf32>
  }
}

LLVM ERROR: Unsupported op: %0 = stablehlo.uniform_quantize %arg0 : (tensor<1x3x32x32xf32>) -> tensor<1x3x32x32x!quant.uniform<i8:f32, 0.029222087934613228>>

Aborted (core dumped)

Appreciate any help and ideas.

Thanks
Axel

Sandeep Dasgupta

unread,
Jun 3, 2024, 1:51:55 PMJun 3
to Axel Wong, OpenXLA Discuss
Hello Axel
I have just added an https://github.com/openxla/stablehlo/issues/2373 to address similar cases. The interpreter does not have support for the quantized types yet and that is something we would like to explore further. 

Also, figuring out the solution presented in the post  may not be applicable directly as the pass mhlo-quant-lelaglize-to-int met some issue which we are still exploring. We will get back asap. 

Regards,
Sandeep

--
You received this message because you are subscribed to the Google Groups "OpenXLA Discuss" group.
To unsubscribe from this group and stop receiving emails from it, send an email to openxla-discu...@openxla.org.
To view this discussion on the web visit https://groups.google.com/a/openxla.org/d/msgid/openxla-discuss/d84e84e3-422c-4765-aa38-07d999c5efdan%40openxla.org.
For more options, visit https://groups.google.com/a/openxla.org/d/optout.
Reply all
Reply to author
Forward
0 new messages