Skip to content

Commit

Permalink
Integrates/llvm 20240621 (iree-org#17723)
Browse files Browse the repository at this point in the history
Bump stablehlo to:
openxla/stablehlo@57d16b1
Bump llvm-project to:
llvm/llvm-project@c83d9e9
  • Loading branch information
nirvedhmeshram authored Jun 21, 2024
1 parent ac418d1 commit 7b58c71
Show file tree
Hide file tree
Showing 6 changed files with 22 additions and 22 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -377,8 +377,8 @@ func.func @scatter_ui32(%arg0: tensor<1xui32>, %arg1: tensor<1x1xi32>, %arg2: te
// CHECK: %[[ARG0:[a-zA-Z0-9]+]]
// CHECK: %[[ARG1:[a-zA-Z0-9]+]]
// CHECK: %[[ARG2:[a-zA-Z0-9]+]]
// CHECK: %[[BITCAST0:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] : tensor<1xui32> to tensor<1xi32>
// CHECK: %[[BITCAST2:.+]] = builtin.unrealized_conversion_cast %[[ARG2]] : tensor<1xui32> to tensor<1xi32>
// CHECK-DAG: %[[BITCAST0:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] : tensor<1xui32> to tensor<1xi32>
// CHECK-DAG: %[[BITCAST2:.+]] = builtin.unrealized_conversion_cast %[[ARG2]] : tensor<1xui32> to tensor<1xi32>
// CHECK: %[[SCATTER:.+]] = iree_linalg_ext.scatter
// CHECK-SAME: unique_indices(true)
// CHECK-SAME: ins(%[[BITCAST2]], %[[ARG1]] : tensor<1xi32>, tensor<1x1xi32>)
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -28,8 +28,8 @@ func.func @addf_bf16(%arg0 : bf16, %arg1 : bf16) -> bf16 {
// CHECK-LABEL: @addf_bf16
// CHECK-SAME: %[[ARG0:.+]]: bf16,
// CHECK-SAME: %[[ARG1:.+]]: bf16
// CHECK: %[[EXT0:.+]] = arith.extf %[[ARG0]] : bf16 to f32
// CHECK: %[[EXT1:.+]] = arith.extf %[[ARG1]] : bf16 to f32
// CHECK-DAG: %[[EXT0:.+]] = arith.extf %[[ARG0]] : bf16 to f32
// CHECK-DAG: %[[EXT1:.+]] = arith.extf %[[ARG1]] : bf16 to f32
// CHECK: %[[ADD:.+]] = arith.addf %[[EXT0]], %[[EXT1]] : f32
// CHECK: %[[TRUNC:.+]] = arith.truncf %[[ADD]] : f32 to bf16

Expand All @@ -43,8 +43,8 @@ func.func @addf_vector_bf16(%arg0 : vector<4xbf16>, %arg1 : vector<4xbf16>) -> v
// CHECK-LABEL: @addf_vector_bf16
// CHECK-SAME: %[[ARG0:.+]]: vector<4xbf16>,
// CHECK-SAME: %[[ARG1:.+]]: vector<4xbf16>
// CHECK: %[[EXT0:.+]] = arith.extf %[[ARG0]] : vector<4xbf16> to vector<4xf32>
// CHECK: %[[EXT1:.+]] = arith.extf %[[ARG1]] : vector<4xbf16> to vector<4xf32>
// CHECK-DAG: %[[EXT0:.+]] = arith.extf %[[ARG0]] : vector<4xbf16> to vector<4xf32>
// CHECK-DAG: %[[EXT1:.+]] = arith.extf %[[ARG1]] : vector<4xbf16> to vector<4xf32>
// CHECK: %[[ADD:.+]] = arith.addf %[[EXT0]], %[[EXT1]] : vector<4xf32>
// CHECK: %[[TRUNC:.+]] = arith.truncf %[[ADD]] : vector<4xf32> to vector<4xbf16>

Expand Down Expand Up @@ -107,17 +107,17 @@ func.func @store_reduction_bf16(%arg0 : vector<3xbf16>, %arg1 : vector<3xbf16>,
}

// CHECK-LABEL: @store_reduction_bf16
// CHECK: %[[CST:.+]] = arith.constant dense<1.000000e+00> : vector<bf16>
// CHECK: %[[VAL0:.+]] = arith.extf %arg0 : vector<3xbf16> to vector<3xf32>
// CHECK: %[[VAL1:.+]] = arith.extf %arg1 : vector<3xbf16> to vector<3xf32>
// CHECK: %[[VAL2:.+]] = vector.extractelement %[[CST]][] : vector<bf16>
// CHECK: %[[VAL3:.+]] = arith.extf %[[VAL2]] : bf16 to f32
// CHECK: %[[VAL4:.+]] = arith.mulf %[[VAL0]], %[[VAL1]] : vector<3xf32>
// CHECK: %[[VAL5:.+]] = vector.reduction <add>, %[[VAL4]], %[[VAL3]] : vector<3xf32> into f32
// CHECK: %[[VAL6:.+]] = arith.truncf %[[VAL5]] : f32 to bf16
// CHECK: %[[VAL7:.+]] = vector.broadcast %[[VAL6]] : bf16 to vector<bf16>
// CHECK: %[[VAL8:.+]] = vector.extractelement %[[VAL7]][] : vector<bf16>
// CHECK: memref.store %[[VAL8]], %arg2[] : memref<bf16>
// CHECK: %[[CST:.+]] = arith.constant dense<1.000000e+00> : vector<bf16>
// CHECK-DAG: %[[VAL0:.+]] = arith.extf %arg0 : vector<3xbf16> to vector<3xf32>
// CHECK-DAG: %[[VAL1:.+]] = arith.extf %arg1 : vector<3xbf16> to vector<3xf32>
// CHECK: %[[VAL2:.+]] = vector.extractelement %[[CST]][] : vector<bf16>
// CHECK: %[[VAL3:.+]] = arith.extf %[[VAL2]] : bf16 to f32
// CHECK: %[[VAL4:.+]] = arith.mulf %[[VAL0]], %[[VAL1]] : vector<3xf32>
// CHECK: %[[VAL5:.+]] = vector.reduction <add>, %[[VAL4]], %[[VAL3]] : vector<3xf32> into f32
// CHECK: %[[VAL6:.+]] = arith.truncf %[[VAL5]] : f32 to bf16
// CHECK: %[[VAL7:.+]] = vector.broadcast %[[VAL6]] : bf16 to vector<bf16>
// CHECK: %[[VAL8:.+]] = vector.extractelement %[[VAL7]][] : vector<bf16>
// CHECK: memref.store %[[VAL8]], %arg2[] : memref<bf16>

// -----

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -797,7 +797,7 @@ static void rewriteForallToLanes(RewriterBase &rewriter, scf::ForallOp forallOp,
Location loc = forallOp->getLoc();
assert(isLaneMappableForall(forallOp) && "mapping non-lane forall op");

Value laneId = rewriter.create<gpu::LaneIdOp>(loc);
Value laneId = rewriter.create<gpu::LaneIdOp>(loc, /*upperBound=*/nullptr);
rewriter.eraseOp(forallOp.getTerminator());
rewriter.setInsertionPoint(forallOp);
rewriter.inlineBlockBefore(forallOp.getBody(), forallOp, {laneId});
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -53,9 +53,9 @@ static Value convertOpTypeFromI32(IREE::HAL::InterfaceConstantLoadOp loadOp,
// AnyFloat
Value resizedValue = extractElementOp.getResult();
if (sourceBitWidth > destBitWidth) {
return builder.create<arith::TruncFOp>(loc, opType, extractElementOp);
return builder.create<arith::TruncFOp>(loc, opType, resizedValue);
} else if (sourceBitWidth < destBitWidth) {
return builder.create<arith::ExtFOp>(loc, opType, extractElementOp);
return builder.create<arith::ExtFOp>(loc, opType, resizedValue);
}
return builder.create<arith::BitcastOp>(loc, opType, resizedValue);
}
Expand Down
2 changes: 1 addition & 1 deletion third_party/llvm-project
Submodule llvm-project updated 2750 files
2 changes: 1 addition & 1 deletion third_party/stablehlo
Submodule stablehlo updated 72 files
+2 −0 BUILD.bazel
+2 −2 WORKSPACE.bazel
+1 −1 build_tools/llvm_version.txt
+67 −0 build_tools/math/README.md
+109 −0 build_tools/math/generate_ChloDecompositionPatternsMath.py
+448 −0 build_tools/math/generate_tests.py
+25 −0 docs/generated/stablehlo_passes.md
+14 −0 stablehlo/api/PortableApi.cpp
+6 −1 stablehlo/api/PortableApi.h
+1 −1 stablehlo/dialect/Version.h
+13 −0 stablehlo/integrations/python/PortableApi.cpp
+7 −0 stablehlo/integrations/python/tests/stablehlo.py
+64 −2 stablehlo/reference/Api.cpp
+13 −0 stablehlo/testdata/quantized/abs_shape_float32_20_20_qi8.mlir
+229 −0 stablehlo/testdata/quantized/add_qi8.mlir
+13 −0 stablehlo/testdata/quantized/cbrt_shape_float32_20_20_qi8.mlir
+46 −0 stablehlo/testdata/quantized/ceil_qi8.mlir
+13 −0 stablehlo/testdata/quantized/ceil_shape_float32_20_20_qi8.mlir
+13 −0 stablehlo/testdata/quantized/cos_shape_float32_20_20_qi8.mlir
+68 −0 stablehlo/testdata/quantized/cosine_qi8.mlir
+15 −0 stablehlo/testdata/quantized/div_dtypes_lhs_float32_2__rhs_float32_2_qi8.mlir
+13 −0 stablehlo/testdata/quantized/div_singularity_0_by_0_lhs_float32_2__rhs_float32_2_qi8.mlir
+13 −0 stablehlo/testdata/quantized/div_singularity_inf_by_inf_lhs_float32_1__rhs_float32_1_qi8.mlir
+15 −0 stablehlo/testdata/quantized/div_singularity_negative_by_0_lhs_float32_2__rhs_float32_2_qi8.mlir
+15 −0 stablehlo/testdata/quantized/div_singularity_positive_by_0_lhs_float32_2__rhs_float32_2_qi8.mlir
+13 −0 stablehlo/testdata/quantized/exp_shape_float32_20_20_qi8.mlir
+13 −0 stablehlo/testdata/quantized/expm1_shape_float32_20_20_qi8.mlir
+46 −0 stablehlo/testdata/quantized/floor_qi8.mlir
+13 −0 stablehlo/testdata/quantized/floor_shape_float32_20_20_qi8.mlir
+13 −0 stablehlo/testdata/quantized/integer_pow_dtypes_shape_float32_20_30__y_2_qi8.mlir
+13 −0 stablehlo/testdata/quantized/log1p_shape_float32_20_20_qi8.mlir
+13 −0 stablehlo/testdata/quantized/log_shape_float32_20_20_qi8.mlir
+15 −0 stablehlo/testdata/quantized/mul_dtypes_lhs_float32_20_20__rhs_float32_20_20_qi8.mlir
+181 −0 stablehlo/testdata/quantized/multiply_qi8.mlir
+13 −0 stablehlo/testdata/quantized/neg_shape_float32_20_20_qi8.mlir
+178 −0 stablehlo/testdata/quantized/negate_qi8.mlir
+13 −0 stablehlo/testdata/quantized/sign_shape_float32_20_20_qi8.mlir
+13 −0 stablehlo/testdata/quantized/sign_special_0_dtype_float32_qi8.mlir
+13 −0 stablehlo/testdata/quantized/sin_shape_float32_20_20_qi8.mlir
+68 −0 stablehlo/testdata/quantized/sine_qi8.mlir
+13 −0 stablehlo/testdata/quantized/sqrt_shape_float32_20_20_qi8.mlir
+15 −0 stablehlo/testdata/quantized/sub_dtypes_lhs_float32_20_20__rhs_float32_20_20_qi8.mlir
+193 −0 stablehlo/testdata/quantized/subtract_qi8.mlir
+68 −0 stablehlo/testdata/quantized/tanh_qi8.mlir
+13 −0 stablehlo/testdata/quantized/tanh_shape_float32_20_20_qi8.mlir
+1 −0 stablehlo/tests/BUILD.bazel
+83 −0 stablehlo/tests/CheckOps.cpp
+3 −0 stablehlo/tests/CheckOps.h
+42 −0 stablehlo/tests/CheckOps.td
+1 −1 stablehlo/tests/TestUtils.cpp
+309 −39 stablehlo/tests/chlo/chlo_legalize_to_stablehlo.mlir
+54 −0 stablehlo/tests/interpret/quantized_ops.mlir
+19 −0 stablehlo/tests/math/asin_complex128.mlir
+19 −0 stablehlo/tests/math/asin_complex64.mlir
+19 −0 stablehlo/tests/math/asin_float32.mlir
+19 −0 stablehlo/tests/math/asin_float64.mlir
+72 −0 stablehlo/tests/math/ulp_difference_float32.mlir
+72 −0 stablehlo/tests/math/ulp_difference_float64.mlir
+285 −9 stablehlo/tests/ops_stablehlo_quantized.mlir
+26 −8 stablehlo/tests/stablehlo_legalize_quant_to_int.mlir
+11 −0 stablehlo/tools/StablehloTranslateMain.cpp
+1 −0 stablehlo/transforms/CMakeLists.txt
+8 −19 stablehlo/transforms/ChloDecompositionPatterns.td
+386 −0 stablehlo/transforms/ChloDecompositionPatternsMath.td
+8 −0 stablehlo/transforms/ChloLegalizeToStablehlo.cpp
+16 −0 stablehlo/transforms/PassPipelines.cpp
+10 −0 stablehlo/transforms/Passes.h
+31 −0 stablehlo/transforms/Passes.td
+1 −1 stablehlo/transforms/StablehloCanonicalizeDynamism.cpp
+13 −0 stablehlo/transforms/StablehloLegalizeQuantToInt.cpp
+136 −0 stablehlo/transforms/StablehloLegalizeQuantizedOpToQDQ.cpp
+1 −1 stablehlo/transforms/StablehloRefineShapes.cpp

0 comments on commit 7b58c71

Please sign in to comment.