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