Integrate LLVM to llvm/llvm-project@c53000a7282b (#24935) Integrate LLVM to llvm/llvm-project@c53000a7282b from branch `sm-iree-integrates/llvm-20260916` Carries adaptations for the following upstream commits: * llvm/llvm-project@40d4395bb229, llvm/llvm-project@65b48e45ea94, llvm/llvm-project@b570e2aa0aa4: Arith header no longer transitive. * llvm/llvm-project@3b99e79991be, llvm/llvm-project@70f5696eaf02, llvm/llvm-project@80164074c01f, llvm/llvm-project@f7dd24bbe0e9: Linalg unary and binary named ops removed; tests use linalg.elementwise <kind>. * llvm/llvm-project@2eb384fccbf8: generalizeNamedOp returns FailureOr<LinalgOp>. * llvm/llvm-project@173a02c755bf: LLVM dialect strict property assembly format in CHECK lines. * llvm/llvm-project@fc696f791e06: aggregate builders with mixed attributes are deprecated and fatal under Bazel's -Werror; call sites use typed builders, copy the source op's properties, or build through OperationState. * llvm/llvm-project@90cebef14116: MemCpyOpt forms llvm.memset in no-builtins dispatches; the device bitcode library now supplies the libcall. * llvm/llvm-project@492410de73b0: TargetID::getCanonicalFeatureString is now getCanonicalTargetIDString. * llvm/llvm-project@9d4a7d05d2ef: TargetOptions::AllowFPOpFusion is gone. FMA formation has followed the contract flag since llvm/llvm-project@bb625414d6ff. Assisted-by: claude code --------- Signed-off-by: Jelle Schühmacher <schuehmacher@roofline.ai> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com>
diff --git a/compiler/plugins/input/StableHLO/Conversion/MapStableHLOToScalarOp.h b/compiler/plugins/input/StableHLO/Conversion/MapStableHLOToScalarOp.h index 9754694..7da10fa 100644 --- a/compiler/plugins/input/StableHLO/Conversion/MapStableHLOToScalarOp.h +++ b/compiler/plugins/input/StableHLO/Conversion/MapStableHLOToScalarOp.h
@@ -197,7 +197,7 @@ Value operator()(Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> /*argTypes*/, ValueRange args, OpBuilder *b) { return StdScalarOp::create(*b, loc, resultTypes, args, - ArrayRef<NamedAttribute>{}); + typename StdScalarOp::Properties{}); } }; @@ -208,7 +208,7 @@ Type elementType = getElementTypeOrSelf(argTypes.front()); if (SupportedType{}(elementType)) { return StdScalarOp::create(*b, loc, resultTypes, args, - ArrayRef<NamedAttribute>{}); + typename StdScalarOp::Properties{}); } return MapStableHloOpToScalarOpImpl<Args...>{}(loc, resultTypes, argTypes, args, b); @@ -689,7 +689,7 @@ mlir::arith::UIToFPOp::areCastCompatible(convertedSourceType, targetType)) { return mlir::arith::UIToFPOp::create(*b, loc, resultTypes, args, - ArrayRef<NamedAttribute>{}); + mlir::arith::UIToFPOp::Properties{}); } if (mlir::arith::SIToFPOp::areCastCompatible(sourceType, targetType)) { return mlir::arith::SIToFPOp::create(*b, loc, resultTypes, args, @@ -700,18 +700,20 @@ auto res = cast<FloatType>(targetType); if (src.getWidth() > res.getWidth()) { return mlir::arith::TruncFOp::create(*b, loc, resultTypes, args, - ArrayRef<NamedAttribute>{}); + mlir::arith::TruncFOp::Properties{}); } if (src.getWidth() < res.getWidth()) { return mlir::arith::ExtFOp::create(*b, loc, resultTypes, args, - ArrayRef<NamedAttribute>{}); + mlir::arith::ExtFOp::Properties{}); } // There's no direct conversion between different 16 bit floating point // types, so go through 32 bit float. if (sourceType != targetType) { assert(sourceType.isBF16() || targetType.isBF16()); - Value ext = arith::ExtFOp::create(*b, loc, b->getF32Type(), args); - return arith::TruncFOp::create(*b, loc, resultTypes, ext); + Value ext = arith::ExtFOp::create(*b, loc, TypeRange{b->getF32Type()}, + args, arith::ExtFOp::Properties{}); + return arith::TruncFOp::create(*b, loc, resultTypes, ValueRange{ext}, + arith::TruncFOp::Properties{}); } // No conversion is needed for identical float types. return args.front(); @@ -737,13 +739,13 @@ auto res = cast<IntegerType>(targetType); if (src.getWidth() > res.getWidth()) { return mlir::arith::TruncIOp::create(*b, loc, resultTypes, args, - ArrayRef<NamedAttribute>{}); + mlir::arith::TruncIOp::Properties{}); } if (src.getWidth() < res.getWidth()) { // Special case boolean values, so they get casted to `1` instead of `-1`. if (IsUnsignedIntegerType{}(src)) { return mlir::arith::ExtUIOp::create(*b, loc, resultTypes, args, - ArrayRef<NamedAttribute>{}); + mlir::arith::ExtUIOp::Properties{}); } return mlir::arith::ExtSIOp::create(*b, loc, resultTypes, args, ArrayRef<NamedAttribute>{});
diff --git a/compiler/plugins/input/StableHLO/Conversion/Preprocessing/StableHLOToStableHLO.cpp b/compiler/plugins/input/StableHLO/Conversion/Preprocessing/StableHLOToStableHLO.cpp index e152d9c..94b13af 100644 --- a/compiler/plugins/input/StableHLO/Conversion/Preprocessing/StableHLOToStableHLO.cpp +++ b/compiler/plugins/input/StableHLO/Conversion/Preprocessing/StableHLOToStableHLO.cpp
@@ -109,7 +109,8 @@ SmallVector<Value, 2> operands = {transposed, op.getRhs()}; auto newConv = mlir::stablehlo::ConvolutionOp::create( - rewriter, op.getLoc(), op.getType(), operands, op->getAttrs()); + rewriter, op.getLoc(), TypeRange{op.getType()}, operands, + op.getProperties(), op->getDiscardableAttrDictionary().getValue()); newConv.setDimensionNumbersAttr(newDimensionNumbers); rewriter.replaceOp(op, newConv.getResult()); @@ -177,7 +178,8 @@ SmallVector<Value, 2> operands = {op.getLhs(), transposeKernel}; mlir::stablehlo::ConvolutionOp newConv = mlir::stablehlo::ConvolutionOp::create( - rewriter, op.getLoc(), op.getType(), operands, op->getAttrs()); + rewriter, op.getLoc(), TypeRange{op.getType()}, operands, + op.getProperties(), op->getDiscardableAttrDictionary().getValue()); newConv.setDimensionNumbersAttr(newDimensionNumbers); rewriter.replaceOp(op, {newConv.getResult()}); @@ -242,8 +244,10 @@ SmallVector<Value, 2> operands = {op.getLhs(), op.getRhs()}; auto newConv = mlir::stablehlo::ConvolutionOp::create( rewriter, op.getLoc(), - RankedTensorType::get(convShape, resultType.getElementType()), operands, - op->getAttrs()); + TypeRange{ + RankedTensorType::get(convShape, resultType.getElementType())}, + operands, op.getProperties(), + op->getDiscardableAttrDictionary().getValue()); newConv.setDimensionNumbersAttr(newDimensionNumbers); auto transposed = mlir::stablehlo::TransposeOp::create( @@ -1717,8 +1721,9 @@ auto resultShape = bcastOperandType.getShape(); auto resultType = RankedTensorType::get(resultShape, elementType); - Value result = ElementwiseOpT::create(rewriter, op.getLoc(), resultType, - bcastOperands); + Value result = ElementwiseOpT::create( + rewriter, op.getLoc(), TypeRange{resultType}, bcastOperands, + op.getProperties(), op->getDiscardableAttrDictionary().getValue()); rewriter.replaceOpWithNewOp<mlir::stablehlo::BroadcastInDimOp>( op, op.getType(), result, bcastOps[0].getBroadcastDimensionsAttr()); @@ -1803,8 +1808,9 @@ op, "fused operands do not satisfy op type constraints"); } - rewriter.replaceOpWithNewOp<Op>(op, op->getResultTypes(), operands, - op->getAttrs()); + rewriter.replaceOpWithNewOp<Op>( + op, op->getResultTypes(), operands, op.getProperties(), + op->getDiscardableAttrDictionary().getValue()); return success(); } };
diff --git a/compiler/plugins/target/ROCM/ROCMTarget.cpp b/compiler/plugins/target/ROCM/ROCMTarget.cpp index 1b83dcd..254f3a2 100644 --- a/compiler/plugins/target/ROCM/ROCMTarget.cpp +++ b/compiler/plugins/target/ROCM/ROCMTarget.cpp
@@ -448,7 +448,7 @@ if (failed(targetID)) { return nullptr; } - format = targetID->getCanonicalFeatureString(); + format = targetID->getCanonicalTargetIDString(); } else { format = "rocm-hsaco-fb"; // legacy HIP } @@ -836,7 +836,6 @@ return variantOp.emitError() << "cannot initialize target triple"; } llvm::TargetOptions opt; - opt.AllowFPOpFusion = llvm::FPOpFusion::Fast; // Be extra cautious while this is less tested, and prevent unknown // fallbacks from global isel. // @@ -1178,7 +1177,7 @@ break; } case ContainerType::AMDGPU: { - std::string canonicalTargetID = targetID->getCanonicalFeatureString(); + std::string canonicalTargetID = targetID->getCanonicalTargetIDString(); executableBinaryFormat = executableBuilder.getStringAttr(canonicalTargetID); binaryContainer = serializeAMDGPUBinaryContainer(
diff --git a/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ArithToF32.cpp b/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ArithToF32.cpp index 7641108..390abff 100644 --- a/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ArithToF32.cpp +++ b/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ArithToF32.cpp
@@ -55,7 +55,8 @@ } if (inputETy.getIntOrFloatBitWidth() < eTy.getIntOrFloatBitWidth()) { - return arith::ExtFOp::create(builder, loc, type, inputs[0]); + return arith::ExtFOp::create(builder, loc, type, inputs[0], + arith::FastMathFlagsAttr{}); } return nullptr; @@ -189,7 +190,9 @@ resultEType.getWidth())) { return rewriter.notifyMatchFailure(op, "invalid width combination"); } - rewriter.replaceOpWithNewOp<OpTy>(op, resultType, op.getOperand()); + rewriter.replaceOpWithNewOp<OpTy>( + op, TypeRange{resultType}, ValueRange{op.getOperand()}, + op.getProperties(), op->getDiscardableAttrDictionary().getValue()); return success(); } };
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUGeneralizeNamedOps.cpp b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUGeneralizeNamedOps.cpp index 295250e..6b1f0e5 100644 --- a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUGeneralizeNamedOps.cpp +++ b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUGeneralizeNamedOps.cpp
@@ -35,7 +35,7 @@ getCompilationInfo(linalgOp); rewriter.setInsertionPoint(linalgOp); - FailureOr<linalg::GenericOp> generalizedOp = + FailureOr<linalg::LinalgOp> generalizedOp = linalg::generalizeNamedOp(rewriter, linalgOp); if (failed(generalizedOp)) { linalgOp->emitOpError("failed to generalize operation");
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_apply_tiling_level.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_apply_tiling_level.mlir index c88983d..04067ef 100644 --- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_apply_tiling_level.mlir +++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_apply_tiling_level.mlir
@@ -491,7 +491,7 @@ func.func @swap_expand_shape_with_extract_slice(%0: tensor<60xf32>) -> tensor<2x3x10xf32> { %expand = tensor.expand_shape %0 [[0, 1, 2]] output_shape [2, 3, 10] : tensor<60xf32> into tensor<2x3x10xf32> %empty = tensor.empty() : tensor<2x3x10xf32> - %exp = linalg.exp {lowering_config = #config} ins(%expand : tensor<2x3x10xf32>) outs(%empty : tensor<2x3x10xf32>) -> tensor<2x3x10xf32> + %exp = linalg.elementwise <exp> {lowering_config = #config} ins(%expand : tensor<2x3x10xf32>) outs(%empty : tensor<2x3x10xf32>) -> tensor<2x3x10xf32> return %exp : tensor<2x3x10xf32> } } @@ -501,7 +501,7 @@ // THREAD: %[[LINEAR_IDX:.+]] = affine.linearize_index disjoint [%[[X]], %[[Y]], %[[Z]]] by (2, 3, 10) // THREAD: %[[SLICE:.+]] = tensor.extract_slice %{{.*}}[%[[LINEAR_IDX]]] [5] [1] : tensor<60xf32> to tensor<5xf32> // THREAD: %[[EXPAND:.+]] = tensor.expand_shape %[[SLICE]] {{\[\[}}0, 1, 2]] output_shape [1, 1, 5] -// THREAD: linalg.exp {{.*}} ins(%[[EXPAND]] +// THREAD: linalg.elementwise <exp> {{.*}} ins(%[[EXPAND]] // ----- @@ -510,7 +510,7 @@ func.func @swap_expand_shape_with_extract_slice_full_inner_dim(%0: tensor<120xf32>) -> tensor<3x4x10xf32> { %expand = tensor.expand_shape %0 [[0, 1, 2]] output_shape [3, 4, 10] : tensor<120xf32> into tensor<3x4x10xf32> %empty = tensor.empty() : tensor<3x4x10xf32> - %exp = linalg.exp {lowering_config = #config} ins(%expand : tensor<3x4x10xf32>) outs(%empty : tensor<3x4x10xf32>) -> tensor<3x4x10xf32> + %exp = linalg.elementwise <exp> {lowering_config = #config} ins(%expand : tensor<3x4x10xf32>) outs(%empty : tensor<3x4x10xf32>) -> tensor<3x4x10xf32> return %exp : tensor<3x4x10xf32> } } @@ -521,7 +521,7 @@ // THREAD: %[[LINEAR_IDX:.+]] = affine.linearize_index disjoint [%[[X]], %[[Y]], %[[C0]]] by (3, 4, 10) // THREAD: %[[SLICE:.+]] = tensor.extract_slice %{{.*}}[%[[LINEAR_IDX]]] [20] [1] : tensor<120xf32> to tensor<20xf32> // THREAD: %[[EXPAND:.+]] = tensor.expand_shape %[[SLICE]] {{\[\[}}0, 1, 2]] output_shape [1, 2, 10] -// THREAD: linalg.exp {{.*}} ins(%[[EXPAND]] +// THREAD: linalg.elementwise <exp> {{.*}} ins(%[[EXPAND]] // ----- @@ -530,7 +530,7 @@ func.func @no_swap_expand_shape_with_extract_slice_non_contiguous(%0: tensor<120xf32>) -> tensor<3x4x10xf32> { %expand = tensor.expand_shape %0 [[0, 1, 2]] output_shape [3, 4, 10] : tensor<120xf32> into tensor<3x4x10xf32> %empty = tensor.empty() : tensor<3x4x10xf32> - %exp = linalg.exp {lowering_config = #config} ins(%expand : tensor<3x4x10xf32>) outs(%empty : tensor<3x4x10xf32>) -> tensor<3x4x10xf32> + %exp = linalg.elementwise <exp> {lowering_config = #config} ins(%expand : tensor<3x4x10xf32>) outs(%empty : tensor<3x4x10xf32>) -> tensor<3x4x10xf32> return %exp : tensor<3x4x10xf32> } } @@ -538,7 +538,7 @@ // THREAD-LABEL: func.func @no_swap_expand_shape_with_extract_slice_non_contiguous // THREAD: tensor.expand_shape // THREAD: scf.forall -// THREAD: linalg.exp +// THREAD: linalg.elementwise <exp> // ----- @@ -547,7 +547,7 @@ func.func @swap_expand_shape_with_extract_slice_multiple_expanded_dims(%0: tensor<120x56xf32>) -> tensor<3x4x10x7x8xf32> { %expand = tensor.expand_shape %0 [[0, 1, 2], [3, 4]] output_shape [3, 4, 10, 7, 8] : tensor<120x56xf32> into tensor<3x4x10x7x8xf32> %empty = tensor.empty() : tensor<3x4x10x7x8xf32> - %exp = linalg.exp {lowering_config = #config} + %exp = linalg.elementwise <exp> {lowering_config = #config} ins(%expand : tensor<3x4x10x7x8xf32>) outs(%empty : tensor<3x4x10x7x8xf32>) -> tensor<3x4x10x7x8xf32> return %exp : tensor<3x4x10x7x8xf32> } @@ -560,7 +560,7 @@ // THREAD: %[[LINEAR_IDX1:.+]] = affine.linearize_index disjoint [%[[ID2]], %[[ID3]]] by (7, 8) // THREAD: %[[SLICE:.+]] = tensor.extract_slice %{{.*}}[%[[LINEAR_IDX0]], %[[LINEAR_IDX1]]] [20, 4] [1, 1] // THREAD: %[[EXPAND:.+]] = tensor.expand_shape %[[SLICE]] {{\[\[}}0, 1, 2], [3, 4]] output_shape [1, 2, 10, 1, 4] -// THREAD: linalg.exp {{.*}} ins(%[[EXPAND]] +// THREAD: linalg.elementwise <exp> {{.*}} ins(%[[EXPAND]] // -----
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_fuse_and_hoist_forall.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_fuse_and_hoist_forall.mlir index ff1f36c..2170cc3 100644 --- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_fuse_and_hoist_forall.mlir +++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_fuse_and_hoist_forall.mlir
@@ -210,7 +210,7 @@ scf.yield %9 : tensor<128x128xf16> } %transpose = linalg.transpose ins(%8: tensor<128x128xf16>) outs(%empty: tensor<128x128xf16>) permutation = [1, 0] - %ceil = linalg.ceil ins(%transpose: tensor<128x128xf16>) outs(%empty: tensor<128x128xf16>) -> tensor<128x128xf16> + %ceil = linalg.elementwise <ceil> ins(%transpose: tensor<128x128xf16>) outs(%empty: tensor<128x128xf16>) -> tensor<128x128xf16> return %ceil : tensor<128x128xf16> } @@ -220,7 +220,7 @@ // CHECK: %[[LOOP:.+]] = scf.for {{.*}} -> (tensor<2x4xf16>) // CHECK: linalg.copy // CHECK: %[[T:.+]] = linalg.transpose ins(%[[LOOP]] : tensor<2x4xf16>) -// CHECK: linalg.ceil ins(%[[T]] : tensor<4x2xf16>) {{.*}} -> tensor<4x2xf16> +// CHECK: linalg.elementwise <ceil> ins(%[[T]] : tensor<4x2xf16>) {{.*}} -> tensor<4x2xf16> // CHECK: scf.forall.in_parallel // CHECK: scf.forall.in_parallel // CHECK: return @@ -250,7 +250,7 @@ scf.yield %9 : tensor<128x128xf16> } %transpose_input = linalg.transpose ins(%3: tensor<128x128xf16>) outs(%empty: tensor<128x128xf16>) permutation = [1, 0] - %add = linalg.add + %add = linalg.elementwise <add> ins(%8, %transpose_input : tensor<128x128xf16>, tensor<128x128xf16>) outs(%empty: tensor<128x128xf16>) -> tensor<128x128xf16> return %add : tensor<128x128xf16> @@ -264,7 +264,7 @@ // CHECK: %[[LOOP:.+]] = scf.for {{.*}} -> (tensor<2x4xf16>) // CHECK: linalg.copy // CHECK: %[[T:.+]] = linalg.transpose ins(%{{.*}} : tensor<4x2xf16>) -// CHECK: linalg.add ins(%[[LOOP]], %[[T]] : tensor<2x4xf16>, tensor<2x4xf16>) {{.*}} -> tensor<2x4xf16> +// CHECK: linalg.elementwise <add> ins(%[[LOOP]], %[[T]] : tensor<2x4xf16>, tensor<2x4xf16>) {{.*}} -> tensor<2x4xf16> // CHECK: scf.forall.in_parallel // CHECK: scf.forall.in_parallel // CHECK: return @@ -501,7 +501,7 @@ tensor.parallel_insert_slice %16 into %arg8[%arg5, %arg6] [2, 2] [1, 1] : tensor<2x2xf16> into tensor<128x128xf16> } } {mapping = [#gpu.thread<linear_dim_1>, #gpu.thread<linear_dim_0>]} - %add = linalg.add + %add = linalg.elementwise <add> ins(%10#0, %10#1 : tensor<128x128xf16>, tensor<128x128xf16>) outs(%empty: tensor<128x128xf16>) -> tensor<128x128xf16> return %add : tensor<128x128xf16> @@ -510,7 +510,7 @@ // CHECK-LABEL: func @fuse_multi_use( // CHECK: scf.forall // CHECK: linalg.copy -// CHECK: linalg.add +// CHECK: linalg.elementwise <add> // CHECK: scf.forall.in_parallel // CHECK: return @@ -537,7 +537,7 @@ tensor.parallel_insert_slice %17 into %arg8[%arg6, %arg5] [2, 2] [1, 1] : tensor<2x2xf16> into tensor<128x128xf16> } } {mapping = [#gpu.thread<linear_dim_1>, #gpu.thread<linear_dim_0>]} - %add = linalg.add + %add = linalg.elementwise <add> ins(%10#0, %10#1 : tensor<128x128xf16>, tensor<128x128xf16>) outs(%empty: tensor<128x128xf16>) -> tensor<128x128xf16> return %add : tensor<128x128xf16> @@ -770,7 +770,7 @@ %extracted_slice_3 = tensor.extract_slice %arg7[0, 0, %arg6] [1, 1, 1] [1, 1, 1] : tensor<1x1x64xf32> to tensor<1x1x1xf32> %extracted_slice_4 = tensor.extract_slice %arg8[0, 0, %arg6] [1, 1, 1] [1, 1, 1] : tensor<1x1x64xf32> to tensor<1x1x1xf32> %16 = linalg.copy ins(%extracted_slice_2 : tensor<1x1x1xf32>) outs(%extracted_slice_3 : tensor<1x1x1xf32>) -> tensor<1x1x1xf32> - %17 = linalg.add ins(%extracted_slice_2, %extracted_slice_2 : tensor<1x1x1xf32>, tensor<1x1x1xf32>) outs(%extracted_slice_4 : tensor<1x1x1xf32>) -> tensor<1x1x1xf32> + %17 = linalg.elementwise <add> ins(%extracted_slice_2, %extracted_slice_2 : tensor<1x1x1xf32>, tensor<1x1x1xf32>) outs(%extracted_slice_4 : tensor<1x1x1xf32>) -> tensor<1x1x1xf32> scf.forall.in_parallel { tensor.parallel_insert_slice %16 into %arg7[0, 0, %arg6] [1, 1, 1] [1, 1, 1] : tensor<1x1x1xf32> into tensor<1x1x64xf32> tensor.parallel_insert_slice %17 into %arg8[0, 0, %arg6] [1, 1, 1] [1, 1, 1] : tensor<1x1x1xf32> into tensor<1x1x64xf32> @@ -793,7 +793,7 @@ // CHECK-DAG: %[[OUT_SLICE0:.+]] = tensor.extract_slice %[[INIT_ARG0]][%[[TID0]], %[[TID1]], %[[TID2]]] [1, 1, 1] // CHECK-DAG: %[[OUT_SLICE1:.+]] = tensor.extract_slice %[[INIT_ARG1]][%[[TID0]], %[[TID1]], %[[TID2]]] [1, 1, 1] // CHECK: %[[COPY:.+]] = linalg.copy ins(%[[IN_SLICE]]{{.*}} outs(%[[OUT_SLICE0]] -// CHECK: %[[ADD:.+]] = linalg.add ins(%[[IN_SLICE]], %[[IN_SLICE]]{{.*}} outs(%[[OUT_SLICE1]] +// CHECK: %[[ADD:.+]] = linalg.elementwise <add> ins(%[[IN_SLICE]], %[[IN_SLICE]]{{.*}} outs(%[[OUT_SLICE1]] // CHECK: scf.forall.in_parallel { // CHECK-DAG: tensor.parallel_insert_slice %[[COPY]] into %[[INIT_ARG0]][%[[TID0]], %[[TID1]], %[[TID2]]] [1, 1, 1] // CHECK-DAG: tensor.parallel_insert_slice %[[ADD]] into %[[INIT_ARG1]][%[[TID0]], %[[TID1]], %[[TID2]]] [1, 1, 1] @@ -809,7 +809,7 @@ %out0_slice = tensor.extract_slice %out0[%arg3] [1] [1] : tensor<64xf32> to tensor<1xf32> %out1_slice = tensor.extract_slice %out1[%arg3] [1] [1] : tensor<64xf32> to tensor<1xf32> %copied = linalg.copy ins(%arg0 : tensor<1xf32>) outs(%out0_slice : tensor<1xf32>) -> tensor<1xf32> - %add = linalg.add ins(%copied, %arg1 : tensor<1xf32>, tensor<1xf32>) outs(%out1_slice : tensor<1xf32>) -> tensor<1xf32> + %add = linalg.elementwise <add> ins(%copied, %arg1 : tensor<1xf32>, tensor<1xf32>) outs(%out1_slice : tensor<1xf32>) -> tensor<1xf32> scf.forall.in_parallel { tensor.parallel_insert_slice %copied into %out0[%arg3] [1] [1] : tensor<1xf32> into tensor<64xf32> tensor.parallel_insert_slice %add into %out1[%arg3] [1] [1] : tensor<1xf32> into tensor<64xf32> @@ -833,7 +833,7 @@ // CHECK: %[[FORALL:.+]]:2 = scf.forall (%[[TID:.+]]) in (64) // CHECK-SAME: shared_outs(%[[OUT0:.+]] = %[[EMPTY]], %[[OUT1:.+]] = %[[EMPTY]]) // CHECK: %[[COPY:.+]] = linalg.copy ins(%[[ARG0]] -// CHECK: %[[ADD:.+]] = linalg.add ins(%[[COPY]], %[[ARG1]] +// CHECK: %[[ADD:.+]] = linalg.elementwise <add> ins(%[[COPY]], %[[ARG1]] // CHECK: %[[GENERIC:.+]] = linalg.generic{{.*}} ins(%[[ADD]] // CHECK-DAG: tensor.parallel_insert_slice %[[COPY]] into %[[OUT0]] // CHECK-DAG: tensor.parallel_insert_slice %[[GENERIC]] into %[[OUT1]] @@ -996,7 +996,7 @@ %slice0 = tensor.extract_slice %0[%i, %j] [4, 4] [1, 1] : tensor<128x128xf16> to tensor<4x4xf16> %slice1 = tensor.extract_slice %1[%i, %j] [4, 4] [1, 1] : tensor<128x128xf16> to tensor<4x4xf16> %slice_out = tensor.extract_slice %out[%i, %j] [4, 4] [1, 1] : tensor<128x128xf16> to tensor<4x4xf16> - %add = linalg.add ins(%slice0, %slice1 : tensor<4x4xf16>, tensor<4x4xf16>) outs(%slice_out : tensor<4x4xf16>) -> tensor<4x4xf16> + %add = linalg.elementwise <add> ins(%slice0, %slice1 : tensor<4x4xf16>, tensor<4x4xf16>) outs(%slice_out : tensor<4x4xf16>) -> tensor<4x4xf16> scf.forall.in_parallel { tensor.parallel_insert_slice %add into %out[%i, %j] [4, 4] [1, 1] : tensor<4x4xf16> into tensor<128x128xf16> } @@ -1017,5 +1017,5 @@ // CHECK: %[[BARRIER2:.+]] = iree_gpu.barrier_region ins(%[[SWIZZLE2]] : tensor<128x128xf16>) { // CHECK: scf.for // CHECK: linalg.copy -// CHECK: linalg.add +// CHECK: linalg.elementwise <add> // CHECK: return
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_greedily_distribute_to_threads.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_greedily_distribute_to_threads.mlir index 25c052e..323ec65 100644 --- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_greedily_distribute_to_threads.mlir +++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_greedily_distribute_to_threads.mlir
@@ -59,7 +59,7 @@ %extracted_slice = tensor.extract_slice %3[0, %arg0] [64, 8] [1, 1] : tensor<64x64xf32> to tensor<64x8xf32> %extracted_slice_0 = tensor.extract_slice %4[0, %arg0] [64, 8] [1, 1] : tensor<64x64xf32> to tensor<64x8xf32> %extracted_slice_1 = tensor.extract_slice %arg1[0, %arg0] [64, 8] [1, 1] : tensor<64x64xf32> to tensor<64x8xf32> - %7 = linalg.add + %7 = linalg.elementwise <add> ins(%extracted_slice, %extracted_slice_0 : tensor<64x8xf32>, tensor<64x8xf32>) outs(%extracted_slice_1 : tensor<64x8xf32>) -> tensor<64x8xf32> %insert = tensor.insert_slice %7 into %arg1[0, %arg0] [64, 8] [1, 1] : tensor<64x8xf32> into tensor<64x64xf32> @@ -71,7 +71,7 @@ // CHECK-LABEL: func.func @in_nested_region // CHECK: scf.for {{.*}} iter_args // CHECK: scf.forall -// CHECK: linalg.add {{.*}} -> tensor<1x4xf32> +// CHECK: linalg.elementwise <add> {{.*}} -> tensor<1x4xf32> // ----- @@ -90,18 +90,18 @@ %extracted_slice_2 = tensor.extract_slice %extracted_slice[%arg5, %arg6] [1, 4] [1, 1] : tensor<64x8xf32> to tensor<1x4xf32> %extracted_slice_3 = tensor.extract_slice %extracted_slice_0[%arg5, %arg6] [1, 4] [1, 1] : tensor<64x8xf32> to tensor<1x4xf32> %extracted_slice_4 = tensor.extract_slice %arg7[%arg5, %arg6] [1, 4] [1, 1] : tensor<64x8xf32> to tensor<1x4xf32> - %2 = linalg.add ins(%extracted_slice_2, %extracted_slice_3 : tensor<1x4xf32>, tensor<1x4xf32>) outs(%extracted_slice_4 : tensor<1x4xf32>) -> tensor<1x4xf32> + %2 = linalg.elementwise <add> ins(%extracted_slice_2, %extracted_slice_3 : tensor<1x4xf32>, tensor<1x4xf32>) outs(%extracted_slice_4 : tensor<1x4xf32>) -> tensor<1x4xf32> scf.forall.in_parallel { tensor.parallel_insert_slice %2 into %arg7[%arg5, %arg6] [1, 4] [1, 1] : tensor<1x4xf32> into tensor<64x8xf32> } } {mapping = [#gpu.thread<linear_dim_1>, #gpu.thread<linear_dim_0>]} - %7 = linalg.add + %7 = linalg.elementwise <add> ins(%1, %1 : tensor<64x8xf32>, tensor<64x8xf32>) outs(%1 : tensor<64x8xf32>) -> tensor<64x8xf32> %insert = tensor.insert_slice %7 into %arg1[0, %arg0] [64, 8] [1, 1] : tensor<64x8xf32> into tensor<64x64xf32> scf.yield %insert : tensor<64x64xf32> } - %8 = linalg.add + %8 = linalg.elementwise <add> ins(%6, %6 : tensor<64x64xf32>, tensor<64x64xf32>) outs(%6 : tensor<64x64xf32>) -> tensor<64x64xf32> return %8 : tensor<64x64xf32> @@ -115,16 +115,16 @@ // Verify that the existing forall stays the same. // CHECK: scf.forall {{.*}} shared_outs(%[[ITER:.+]] = %[[INIT]]) // CHECK: %[[DEST:.+]] = tensor.extract_slice %[[ITER]] -// CHECK: linalg.add {{.*}} outs(%[[DEST]] : tensor<1x4xf32> +// CHECK: linalg.elementwise <add> {{.*}} outs(%[[DEST]] : tensor<1x4xf32> // CHECK: scf.forall.in_parallel // CHECK: %[[DIST_ADD:.+]] = scf.forall -// CHECK: linalg.add +// CHECK: linalg.elementwise <add> // CHECK: scf.forall.in_parallel // CHECK: %[[INSERT:.+]] = tensor.insert_slice %[[DIST_ADD]] // CHECK: scf.yield %[[INSERT]] // CHECK: %[[RES_ADD:.+]] = scf.forall -// CHECK: linalg.add +// CHECK: linalg.elementwise <add> // CHECK: scf.forall.in_parallel // CHECK: return %[[RES_ADD]] @@ -135,7 +135,7 @@ translation_info = #iree_codegen.translation_info<pipeline = #iree_gpu.pipeline<TileAndFuse> workgroup_size = [64, 1, 1] subgroup_size = 64> } { %add_empty = tensor.empty() : tensor<64x256xf32> - %6 = linalg.add + %6 = linalg.elementwise <add> ins(%3, %4 : tensor<64x256xf32>, tensor<64x256xf32>) outs(%add_empty : tensor<64x256xf32>) -> tensor<64x256xf32> %transpose_empty = tensor.empty() : tensor<256x64xf32> @@ -147,12 +147,12 @@ // CHECK-LABEL: func.func @multiple_use_tilable_op // CHECK: %[[ADD_DIST:.+]] = scf.forall -// CHECK: %[[ADD:.+]] = linalg.add {{.*}} -> tensor<1x4xf32> +// CHECK: %[[ADD:.+]] = linalg.elementwise <add> {{.*}} -> tensor<1x4xf32> // CHECK: scf.forall.in_parallel // CHECK: tensor.parallel_insert_slice %[[ADD]] // CHECK: mapping = [#gpu.thread<linear_dim_1>, #gpu.thread<linear_dim_0>] // CHECK: %[[T_DIST:.+]] = scf.forall -// CHECK: %[[FUSED_ADD:.+]] = linalg.add {{.*}} -> tensor<4x1xf32> +// CHECK: %[[FUSED_ADD:.+]] = linalg.elementwise <add> {{.*}} -> tensor<4x1xf32> // CHECK: %[[T:.+]] = linalg.transpose ins(%[[FUSED_ADD]] // CHECK: scf.forall.in_parallel // CHECK: tensor.parallel_insert_slice %[[T]]
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/add_fmfs.mlir b/compiler/src/iree/compiler/Codegen/Common/test/add_fmfs.mlir index 6ef697b..9cfd607 100644 --- a/compiler/src/iree/compiler/Codegen/Common/test/add_fmfs.mlir +++ b/compiler/src/iree/compiler/Codegen/Common/test/add_fmfs.mlir
@@ -9,5 +9,5 @@ llvm.return %add : f32 } -// CHECK: llvm.fmul %{{.*}}, %{{.*}} {fastmathFlags = #llvm.fastmath<contract>} : f32 -// CHECK: llvm.fadd %{{.*}}, %{{.*}} {fastmathFlags = #llvm.fastmath<contract>} : f32 +// CHECK: llvm.fmul %{{.*}}, %{{.*}} fastmath<contract> : f32 +// CHECK: llvm.fadd %{{.*}}, %{{.*}} fastmath<contract> : f32
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/optimize_tensor_insert_extract_slices.mlir b/compiler/src/iree/compiler/Codegen/Common/test/optimize_tensor_insert_extract_slices.mlir index fb3345f..4392e21 100644 --- a/compiler/src/iree/compiler/Codegen/Common/test/optimize_tensor_insert_extract_slices.mlir +++ b/compiler/src/iree/compiler/Codegen/Common/test/optimize_tensor_insert_extract_slices.mlir
@@ -527,7 +527,7 @@ %loop = scf.for %i = %c0 to %c8 step %c1 iter_args(%arg1 = %init) -> tensor<64x64xf32> { %slice1 = tensor.extract_slice %arg1[1, 0][8, 1][1, 1] : tensor<64x64xf32> to tensor<8x1xf32> %slice2 = tensor.extract_slice %init[0, %i][8, 1][1, 1] : tensor<64x64xf32> to tensor<8x1xf32> - %out = linalg.add ins(%slice1, %slice2 : tensor<8x1xf32>, tensor<8x1xf32>) outs(%empty: tensor<8x1xf32>) -> tensor<8x1xf32> + %out = linalg.elementwise <add> ins(%slice1, %slice2 : tensor<8x1xf32>, tensor<8x1xf32>) outs(%empty: tensor<8x1xf32>) -> tensor<8x1xf32> %td = tensor.insert_slice %out into %t[1, 0][8, 1][1, 1] : tensor<8x1xf32> into tensor<64x64xf32> scf.yield %td : tensor<64x64xf32> } @@ -539,7 +539,7 @@ // CHECK: scf.for {{.*}} iter_args(%[[IV:.+]] = {{.*}}) // CHECK: %[[SLICE:.+]] = tensor.extract_slice // CHECK-NOT: tensor.extract_slice -// CHECK: linalg.add ins(%[[IV]], %[[SLICE]] : {{.*}}) +// CHECK: linalg.elementwise <add> ins(%[[IV]], %[[SLICE]] : {{.*}}) // CHECK: scf.yield // CHECK: tensor.insert_slice @@ -553,7 +553,7 @@ %loop = scf.for %i = %c0 to %c8 step %c1 iter_args(%arg1 = %init) -> tensor<64x64xf32> { %slice1 = tensor.extract_slice %arg1[1, 0][8, 1][1, 1] : tensor<64x64xf32> to tensor<8x1xf32> %slice2 = tensor.extract_slice %init[0, %i][8, 1][1, 1] : tensor<64x64xf32> to tensor<8x1xf32> - %out = linalg.add ins(%slice1, %slice2 : tensor<8x1xf32>, tensor<8x1xf32>) outs(%empty: tensor<8x1xf32>) -> tensor<8x1xf32> + %out = linalg.elementwise <add> ins(%slice1, %slice2 : tensor<8x1xf32>, tensor<8x1xf32>) outs(%empty: tensor<8x1xf32>) -> tensor<8x1xf32> %td = tensor.insert_slice %out into %t[0, 0][8, 1][1, 1] : tensor<8x1xf32> into tensor<64x64xf32> scf.yield %td : tensor<64x64xf32> }
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/tile_and_distribute_workgroups_using_forall.mlir b/compiler/src/iree/compiler/Codegen/Common/test/tile_and_distribute_workgroups_using_forall.mlir index 65a8d8d..cfdc43a 100644 --- a/compiler/src/iree/compiler/Codegen/Common/test/tile_and_distribute_workgroups_using_forall.mlir +++ b/compiler/src/iree/compiler/Codegen/Common/test/tile_and_distribute_workgroups_using_forall.mlir
@@ -644,7 +644,7 @@ %arg2: tensor<3x1xi32>) -> tensor<3x2048x2048xf32> { %cst = arith.constant 0.000000e+00 : f32 %0 = tensor.empty() : tensor<3x2048x2048xf32> - %1 = linalg.add {lowering_config = #iree_gpu.lowering_config<{workgroup = [1, 1, 256]}>} + %1 = linalg.elementwise <add> {lowering_config = #iree_gpu.lowering_config<{workgroup = [1, 1, 256]}>} ins(%arg0, %arg1 : tensor<3x2048x2048xf32>, tensor<3x2048x2048xf32>) outs(%0 : tensor<3x2048x2048xf32>) -> tensor<3x2048x2048xf32> %2 = iree_linalg_ext.scatter dimension_map = [0] unique_indices(true) ins(%1, %arg2 : tensor<3x2048x2048xf32>, tensor<3x1xi32>) outs(%0 : tensor<3x2048x2048xf32>) { @@ -659,7 +659,7 @@ // CHECK-SAME: %[[RHS:[A-Za-z0-9]+]]: tensor<3x2048x2048xf32> // CHECK-SAME: %[[IND:[A-Za-z0-9]+]]: tensor<3x1xi32> // CHECK: %[[RESULT:.+]] = scf.forall (%[[ID0:.+]], %[[ID1:.+]], %[[ID2:[A-Za-z0-9]+]]) {{.*}} shared_outs(%[[DEST:.+]] = %{{.*}}) -// CHECK-DAG: %[[SRC:.+]] = linalg.add +// CHECK-DAG: %[[SRC:.+]] = linalg.elementwise <add> // CHECK-DAG: %[[IND_SLICE:.+]] = tensor.extract_slice %[[IND]][%[[ID0]], 0] {{.*}} : tensor<3x1xi32> to tensor<1x1xi32> // CHECK-DAG: %[[DEST_SLICE:.+]] = tensor.extract_slice %[[DEST]][0, %[[ID1]], %[[ID2]]] {{.*}} to tensor<3x1x256xf32> // CHECK: %[[SCATTER:.+]] = iree_linalg_ext.scatter dimension_map = [0] unique_indices(true)
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/tile_large_tensors.mlir b/compiler/src/iree/compiler/Codegen/Common/test/tile_large_tensors.mlir index cc079cd..bc7da9c 100644 --- a/compiler/src/iree/compiler/Codegen/Common/test/tile_large_tensors.mlir +++ b/compiler/src/iree/compiler/Codegen/Common/test/tile_large_tensors.mlir
@@ -72,7 +72,7 @@ %extracted_slice = tensor.extract_slice %3[0, %arg0] [64, 8] [1, 1] : tensor<64x64xf32> to tensor<64x8xf32> %extracted_slice_0 = tensor.extract_slice %4[0, %arg0] [64, 8] [1, 1] : tensor<64x64xf32> to tensor<64x8xf32> %extracted_slice_1 = tensor.extract_slice %arg1[0, %arg0] [64, 8] [1, 1] : tensor<64x64xf32> to tensor<64x8xf32> - %7 = linalg.add + %7 = linalg.elementwise <add> ins(%extracted_slice, %extracted_slice_0 : tensor<64x8xf32>, tensor<64x8xf32>) outs(%extracted_slice_1 : tensor<64x8xf32>) -> tensor<64x8xf32> %insert = tensor.insert_slice %7 into %arg1[0, %arg0] [64, 8] [1, 1] : tensor<64x8xf32> into tensor<64x64xf32> @@ -84,16 +84,16 @@ // CHECK-LABEL: func.func @in_nested_region // CHECK: scf.for // CHECK: scf.for %{{.*}} = %c0 to %c64 step %c1 -// CHECK: linalg.add {{.*}} -> tensor<1x8xf32> +// CHECK: linalg.elementwise <add> {{.*}} -> tensor<1x8xf32> // ----- func.func @multiple_use_tilable_op(%3: tensor<64x256xf32>, %4: tensor<64x256xf32>) -> (tensor<64x256xf32>, tensor<64x256xf32>) { %empty = tensor.empty() : tensor<64x256xf32> - %6 = linalg.add + %6 = linalg.elementwise <add> ins(%3, %4 : tensor<64x256xf32>, tensor<64x256xf32>) outs(%empty : tensor<64x256xf32>) -> tensor<64x256xf32> - %7 = linalg.exp + %7 = linalg.elementwise <exp> ins(%6 : tensor<64x256xf32>) outs(%empty : tensor<64x256xf32>) -> tensor<64x256xf32> return %6, %7 : tensor<64x256xf32>, tensor<64x256xf32> @@ -101,10 +101,10 @@ // CHECK-LABEL: func.func @multiple_use_tilable_op // CHECK: %[[ADD_TILING:.+]] = scf.for -// CHECK: linalg.add {{.*}} -> tensor<1x64xf32> +// CHECK: linalg.elementwise <add> {{.*}} -> tensor<1x64xf32> // CHECK: %[[EXP_TILING:.+]] = scf.for -// CHECK: %[[FUSED_ADD:.+]] = linalg.add {{.*}} -> tensor<1x64xf32> -// CHECK: linalg.exp ins(%[[FUSED_ADD]] +// CHECK: %[[FUSED_ADD:.+]] = linalg.elementwise <add> {{.*}} -> tensor<1x64xf32> +// CHECK: linalg.elementwise <exp> ins(%[[FUSED_ADD]] // CHECK: return %[[ADD_TILING]], %[[EXP_TILING]] // -----
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/CPU/IR/IREECPUAttrs.cpp b/compiler/src/iree/compiler/Codegen/Dialect/CPU/IR/IREECPUAttrs.cpp index 357920d..058946b 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/CPU/IR/IREECPUAttrs.cpp +++ b/compiler/src/iree/compiler/Codegen/Dialect/CPU/IR/IREECPUAttrs.cpp
@@ -980,7 +980,8 @@ auto vt = cast<VectorType>(v.getType()); auto wideTy = VectorType::get(vt.getShape(), wider); if (isa<FloatType>(vt.getElementType())) { - return arith::ExtFOp::create(builder, loc, wideTy, v); + return arith::ExtFOp::create(builder, loc, wideTy, v, + arith::FastMathFlagsAttr{}); } return arith::ExtSIOp::create(builder, loc, wideTy, v); }; @@ -1209,7 +1210,8 @@ } auto wideTy = VectorType::get(vt.getShape(), accElem); if (isa<FloatType>(accElem)) { - return arith::ExtFOp::create(builder, loc, wideTy, v); + return arith::ExtFOp::create(builder, loc, wideTy, v, + arith::FastMathFlagsAttr{}); } return unsignedSrc ? Value(arith::ExtUIOp::create(builder, loc, wideTy, v)) : Value(arith::ExtSIOp::create(builder, loc, wideTy, v));
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/UKernelOps.cpp b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/UKernelOps.cpp index d94a724..5617050 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/UKernelOps.cpp +++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/UKernelOps.cpp
@@ -336,8 +336,10 @@ nonTensorResultValues.push_back(result); } - auto bufferOp = OpTy::create(rewriter, op->getLoc(), nonTensorResultTypes, - bufferOpOperands, op->getAttrs()); + auto bufferOp = + OpTy::create(rewriter, op->getLoc(), nonTensorResultTypes, + bufferOpOperands, cast<OpTy>(op).getProperties(), + op->getDiscardableAttrDictionary().getValue()); SmallVector<Value> bufferizedResults = cast<DestinationStyleOpInterface>(bufferOp.getOperation()) .getDpsInits();
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/PCF/Transforms/test/fuse_consumers.mlir b/compiler/src/iree/compiler/Codegen/Dialect/PCF/Transforms/test/fuse_consumers.mlir index 2eddca1..0561754 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/PCF/Transforms/test/fuse_consumers.mlir +++ b/compiler/src/iree/compiler/Codegen/Dialect/PCF/Transforms/test/fuse_consumers.mlir
@@ -140,7 +140,7 @@ pcf.write_slice %cst_7 into %ref1[%id0, %id1] [4, 5] [1, 1] : tensor<4x5xi32> into !pcf.sref<?x?xi32, sync(#pcf.test_scope)> pcf.return } - %1 = linalg.add ins(%0#0, %0#1 : tensor<?x?xi32>, tensor<?x?xi32>) outs(%dest: tensor<?x?xi32>) -> tensor<?x?xi32> + %1 = linalg.elementwise <add> ins(%0#0, %0#1 : tensor<?x?xi32>, tensor<?x?xi32>) outs(%dest: tensor<?x?xi32>) -> tensor<?x?xi32> return %1 : tensor<?x?xi32> } @@ -159,7 +159,7 @@ // CHECK: cf.assert // CHECK: %[[SLICE:.+]] = tensor.extract_slice %[[DEST]][%[[ID0]], %[[ID1]]] [4, 5] -// CHECK: %[[COPY:.+]] = linalg.add ins(%[[CST5]], %[[CST7]]{{.*}} outs(%[[SLICE]] +// CHECK: %[[COPY:.+]] = linalg.elementwise <add> ins(%[[CST5]], %[[CST7]]{{.*}} outs(%[[SLICE]] // CHECK: pcf.write_slice %[[COPY]] into %[[REF]][%[[ID0]], %[[ID1]]] // CHECK: pcf.return // CHECK: return %[[GENERIC]] @@ -267,7 +267,7 @@ pcf.write_slice %cst_7 into %ref[%id] [3] [1] : tensor<3xi32> into !pcf.sref<?xi32, sync(#pcf.test_scope)> pcf.return } - %1 = linalg.add ins(%0, %0 : tensor<?xi32>, tensor<?xi32>) outs(%dest: tensor<?xi32>) -> tensor<?xi32> + %1 = linalg.elementwise <add> ins(%0, %0 : tensor<?xi32>, tensor<?xi32>) outs(%dest: tensor<?xi32>) -> tensor<?xi32> return %1 : tensor<?xi32> } @@ -282,13 +282,13 @@ // CHECK: ^bb1 // CHECK: %[[SLICE5:.+]] = tensor.extract_slice %[[DEST]][%[[ID]]] [3] -// CHECK: %[[ADD5:.+]] = linalg.add ins(%[[CST5]], %[[CST5]]{{.*}} outs(%[[SLICE5]] +// CHECK: %[[ADD5:.+]] = linalg.elementwise <add> ins(%[[CST5]], %[[CST5]]{{.*}} outs(%[[SLICE5]] // CHECK: pcf.write_slice %[[ADD5]] into %[[REF]][%[[ID]]] // CHECK: pcf.return // CHECK: ^bb2 // CHECK: %[[SLICE7:.+]] = tensor.extract_slice %[[DEST]][%[[ID]]] [3] -// CHECK: %[[ADD7:.+]] = linalg.add ins(%[[CST7]], %[[CST7]]{{.*}} outs(%[[SLICE7]] +// CHECK: %[[ADD7:.+]] = linalg.elementwise <add> ins(%[[CST7]], %[[CST7]]{{.*}} outs(%[[SLICE7]] // CHECK: pcf.write_slice %[[ADD7]] into %[[REF]][%[[ID]]] // CHECK: pcf.return // CHECK: return %[[GENERIC]] @@ -304,9 +304,9 @@ pcf.write_slice %cst into %ref[%id0, %id1] [4, 5] [1, 1] : tensor<4x5xf32> into !pcf.sref<?x?xf32, sync(#pcf.test_scope)> pcf.return } - %1 = linalg.exp ins(%0: tensor<?x?xf32>) outs(%dest: tensor<?x?xf32>) -> tensor<?x?xf32> - %2 = linalg.sqrt ins(%0: tensor<?x?xf32>) outs(%dest: tensor<?x?xf32>) -> tensor<?x?xf32> - %3 = linalg.add ins(%1, %2 : tensor<?x?xf32>, tensor<?x?xf32>) outs(%dest: tensor<?x?xf32>) -> tensor<?x?xf32> + %1 = linalg.elementwise <exp> ins(%0: tensor<?x?xf32>) outs(%dest: tensor<?x?xf32>) -> tensor<?x?xf32> + %2 = linalg.elementwise <sqrt> ins(%0: tensor<?x?xf32>) outs(%dest: tensor<?x?xf32>) -> tensor<?x?xf32> + %3 = linalg.elementwise <add> ins(%1, %2 : tensor<?x?xf32>, tensor<?x?xf32>) outs(%dest: tensor<?x?xf32>) -> tensor<?x?xf32> return %3 : tensor<?x?xf32> } @@ -319,11 +319,11 @@ // CHECK-NEXT: execute(%[[REF:.+]] = %[[DEST]])[%[[ID0:[A-Za-z0-9_]+]]: index, %[[ID1:[A-Za-z0-9_]+]]: index // CHECK: %[[SLICE1:.+]] = tensor.extract_slice %[[DEST]][%[[ID0]], %[[ID1]]] [4, 5] -// CHECK: %[[SQRT:.+]] = linalg.sqrt ins(%[[CST]]{{.*}} outs(%[[SLICE1]] +// CHECK: %[[SQRT:.+]] = linalg.elementwise <sqrt> ins(%[[CST]]{{.*}} outs(%[[SLICE1]] // CHECK: %[[SLICE0:.+]] = tensor.extract_slice %[[DEST]][%[[ID0]], %[[ID1]]] [4, 5] -// CHECK: %[[EXP:.+]] = linalg.exp ins(%[[CST]]{{.*}} outs(%[[SLICE0]] +// CHECK: %[[EXP:.+]] = linalg.elementwise <exp> ins(%[[CST]]{{.*}} outs(%[[SLICE0]] // CHECK: %[[SLICE2:.+]] = tensor.extract_slice %[[DEST]][%[[ID0]], %[[ID1]]] [4, 5] -// CHECK: %[[ADD:.+]] = linalg.add ins(%[[EXP]], %[[SQRT]]{{.*}} outs(%[[SLICE2]] +// CHECK: %[[ADD:.+]] = linalg.elementwise <add> ins(%[[EXP]], %[[SQRT]]{{.*}} outs(%[[SLICE2]] // CHECK: pcf.write_slice %[[ADD]] into %[[REF]][%[[ID0]], %[[ID1]]] // CHECK: pcf.return // CHECK: return %[[GENERIC]] @@ -455,14 +455,14 @@ pcf.write_slice %cst_7 into %ref1[%id1, %id0] [5, 4] [1, 1] : tensor<5x4xi32> into !pcf.sref<?x?xi32, sync(#pcf.test_scope)> pcf.return } - %1 = linalg.add ins(%0#0, %0#1 : tensor<?x?xi32>, tensor<?x?xi32>) outs(%dest: tensor<?x?xi32>) -> tensor<?x?xi32> + %1 = linalg.elementwise <add> ins(%0#0, %0#1 : tensor<?x?xi32>, tensor<?x?xi32>) outs(%dest: tensor<?x?xi32>) -> tensor<?x?xi32> return %1 : tensor<?x?xi32> } // CHECK-LABEL: @no_fuse_iteration_space_mismatch // CHECK: %[[GENERIC:.+]]:2 = pcf.generic scope(#pcf.test_scope) -// CHECK: %[[ADD:.+]] = linalg.add ins(%[[GENERIC]]#0, %[[GENERIC]]#1 +// CHECK: %[[ADD:.+]] = linalg.elementwise <add> ins(%[[GENERIC]]#0, %[[GENERIC]]#1 // CHECK: return %[[ADD]] // ----- @@ -481,14 +481,14 @@ } pcf.return } - %1 = linalg.add ins(%0#0, %0#1 : tensor<?x?xi32>, tensor<?x?xi32>) outs(%dest: tensor<?x?xi32>) -> tensor<?x?xi32> + %1 = linalg.elementwise <add> ins(%0#0, %0#1 : tensor<?x?xi32>, tensor<?x?xi32>) outs(%dest: tensor<?x?xi32>) -> tensor<?x?xi32> return %1 : tensor<?x?xi32> } // CHECK-LABEL: @no_fuse_no_insertion_point // CHECK: %[[GENERIC:.+]]:2 = pcf.generic scope(#pcf.test_scope) -// CHECK: %[[ADD:.+]] = linalg.add ins(%[[GENERIC]]#0, %[[GENERIC]]#1 +// CHECK: %[[ADD:.+]] = linalg.elementwise <add> ins(%[[GENERIC]]#0, %[[GENERIC]]#1 // CHECK: return %[[ADD]] // -----
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/PCF/Transforms/test/fuse_producers.mlir b/compiler/src/iree/compiler/Codegen/Dialect/PCF/Transforms/test/fuse_producers.mlir index 42c1a68..6c94654 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/PCF/Transforms/test/fuse_producers.mlir +++ b/compiler/src/iree/compiler/Codegen/Dialect/PCF/Transforms/test/fuse_producers.mlir
@@ -19,7 +19,7 @@ : (!pcf.sref<8x16xf32, sync(#pcf.test_scope)>) -> (tensor<8x16xf32>) { %slice = pcf.read_slice %ref[%id0, %id1] [4, 8] [1, 1] : !pcf.sref<8x16xf32, sync(#pcf.test_scope)> to tensor<4x8xf32> - %result = linalg.exp ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> + %result = linalg.elementwise <exp> ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> pcf.write_slice %result into %ref[%id0, %id1] [4, 8] [1, 1] : tensor<4x8xf32> into !pcf.sref<8x16xf32, sync(#pcf.test_scope)> pcf.return } @@ -36,7 +36,7 @@ // CHECK-NEXT: execute(%[[REF:.+]] = %[[DEST]])[%[[ID0:[A-Za-z0-9_]+]]: index, %[[ID1:[A-Za-z0-9_]+]]: index // CHECK: %[[EXTRACT:.+]] = tensor.extract_slice %[[DEST]][%[[ID0]], %[[ID1]]] [4, 8] [1, 1] // CHECK: %[[TILED_FILL:.+]] = linalg.fill ins(%[[CST]]{{.*}} outs(%[[EXTRACT]] : tensor<4x8xf32>) -// CHECK: %[[EXP:.+]] = linalg.exp ins(%[[TILED_FILL]]{{.*}} outs(%[[TILED_FILL]] +// CHECK: %[[EXP:.+]] = linalg.elementwise <exp> ins(%[[TILED_FILL]]{{.*}} outs(%[[TILED_FILL]] // CHECK: pcf.write_slice %[[EXP]] into %[[REF]][%[[ID0]], %[[ID1]]] [4, 8] [1, 1] // CHECK: pcf.return // CHECK: return %[[GENERIC]] @@ -52,7 +52,7 @@ : (!pcf.sref<8x16xf32, sync(#pcf.test_scope)>) -> (tensor<8x16xf32>) { %slice = pcf.read_slice %ref[%id0, %id1] [4, 8] [1, 1] : !pcf.sref<8x16xf32, sync(#pcf.test_scope)> to tensor<4x8xf32> - %result = linalg.exp ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> + %result = linalg.elementwise <exp> ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> pcf.write_slice %result into %ref[%id0, %id1] [4, 8] [1, 1] : tensor<4x8xf32> into !pcf.sref<8x16xf32, sync(#pcf.test_scope)> pcf.return } @@ -68,7 +68,7 @@ // CHECK-NEXT: execute(%[[REF:.+]] = %[[DEST]])[%[[ID0:[A-Za-z0-9_]+]]: index, %[[ID1:[A-Za-z0-9_]+]]: index // CHECK: %[[EXTRACT:.+]] = tensor.extract_slice %[[DEST]][%[[ID0]], %[[ID1]]] [4, 8] [1, 1] // CHECK: %[[TILED_FILL:.+]] = linalg.fill ins(%[[CST]]{{.*}} outs(%[[EXTRACT]] : tensor<4x8xf32>) -// CHECK: %[[EXP:.+]] = linalg.exp ins(%[[TILED_FILL]]{{.*}} outs(%[[TILED_FILL]] +// CHECK: %[[EXP:.+]] = linalg.elementwise <exp> ins(%[[TILED_FILL]]{{.*}} outs(%[[TILED_FILL]] // CHECK: pcf.write_slice %[[EXP]] into %[[REF]][%[[ID0]], %[[ID1]]] [4, 8] [1, 1] // CHECK: pcf.return // CHECK: return %[[LOOP]] @@ -85,7 +85,7 @@ : (!pcf.sref<8x16xf32, sync(#pcf.test_scope)>) -> (tensor<8x16xf32>) { %slice = pcf.read_slice %ref[%id0, %id1] [4, 8] [1, 1] : !pcf.sref<8x16xf32, sync(#pcf.test_scope)> to tensor<4x8xf32> - %result = linalg.exp ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> + %result = linalg.elementwise <exp> ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> pcf.write_slice %result into %ref[%id0, %id1] [4, 8] [1, 1] : tensor<4x8xf32> into !pcf.sref<8x16xf32, sync(#pcf.test_scope)> pcf.return } @@ -106,7 +106,7 @@ // CHECK: %[[INPUT_SLICE:.+]] = tensor.extract_slice %[[INPUT]][%[[ID1]], %[[ID0]]] [8, 4] [1, 1] : tensor<16x8xf32> to tensor<8x4xf32> // CHECK: %[[OUT_SLICE:.+]] = tensor.extract_slice %[[EMPTY]][%[[ID0]], %[[ID1]]] [4, 8] [1, 1] : tensor<8x16xf32> to tensor<4x8xf32> // CHECK: %[[TRANSPOSED:.+]] = linalg.transpose ins(%[[INPUT_SLICE]] : tensor<8x4xf32>) outs(%[[OUT_SLICE]] : tensor<4x8xf32>) permutation = [1, 0] -// CHECK: %[[EXP:.+]] = linalg.exp ins(%[[TRANSPOSED]]{{.*}} outs(%[[TRANSPOSED]] +// CHECK: %[[EXP:.+]] = linalg.elementwise <exp> ins(%[[TRANSPOSED]]{{.*}} outs(%[[TRANSPOSED]] // CHECK: pcf.write_slice %[[EXP]] into %[[REF]][%[[ID0]], %[[ID1]]] [4, 8] [1, 1] // CHECK: pcf.return // CHECK: return %[[GENERIC]] @@ -122,10 +122,10 @@ : (!pcf.sref<8x16xf32, sync(#pcf.test_scope)>) -> (tensor<8x16xf32>) { %slice0 = pcf.read_slice %ref[%id0, 0] [4, 8] [1, 1] : !pcf.sref<8x16xf32, sync(#pcf.test_scope)> to tensor<4x8xf32> - %result0 = linalg.exp ins(%slice0 : tensor<4x8xf32>) outs(%slice0 : tensor<4x8xf32>) -> tensor<4x8xf32> + %result0 = linalg.elementwise <exp> ins(%slice0 : tensor<4x8xf32>) outs(%slice0 : tensor<4x8xf32>) -> tensor<4x8xf32> pcf.write_slice %result0 into %ref[%id0, 0] [4, 8] [1, 1] : tensor<4x8xf32> into !pcf.sref<8x16xf32, sync(#pcf.test_scope)> %slice1 = pcf.read_slice %ref[%id0, 8] [4, 8] [1, 1] : !pcf.sref<8x16xf32, sync(#pcf.test_scope)> to tensor<4x8xf32> - %result1 = linalg.exp ins(%slice1 : tensor<4x8xf32>) outs(%slice1 : tensor<4x8xf32>) -> tensor<4x8xf32> + %result1 = linalg.elementwise <exp> ins(%slice1 : tensor<4x8xf32>) outs(%slice1 : tensor<4x8xf32>) -> tensor<4x8xf32> pcf.write_slice %result1 into %ref[%id0, 8] [4, 8] [1, 1] : tensor<4x8xf32> into !pcf.sref<8x16xf32, sync(#pcf.test_scope)> pcf.return } @@ -141,11 +141,11 @@ // CHECK-NEXT: execute(%[[REF:.+]] = %[[DEST]])[%[[ID0:[A-Za-z0-9_]+]]: index // CHECK: %[[EXT0:.+]] = tensor.extract_slice %[[DEST]][%[[ID0]], 0] [4, 8] [1, 1] // CHECK: %[[FILL0:.+]] = linalg.fill ins(%[[CST]]{{.*}} outs(%[[EXT0]] -// CHECK: %[[EXP0:.+]] = linalg.exp ins(%[[FILL0]] +// CHECK: %[[EXP0:.+]] = linalg.elementwise <exp> ins(%[[FILL0]] // CHECK: pcf.write_slice %[[EXP0]] into %[[REF]][%[[ID0]], 0] [4, 8] [1, 1] // CHECK: %[[EXT1:.+]] = tensor.extract_slice %[[DEST]][%[[ID0]], 8] [4, 8] [1, 1] // CHECK: %[[FILL1:.+]] = linalg.fill ins(%[[CST]]{{.*}} outs(%[[EXT1]] -// CHECK: %[[EXP1:.+]] = linalg.exp ins(%[[FILL1]] +// CHECK: %[[EXP1:.+]] = linalg.elementwise <exp> ins(%[[FILL1]] // CHECK: pcf.write_slice %[[EXP1]] into %[[REF]][%[[ID0]], 8] [4, 8] [1, 1] // CHECK: pcf.return // CHECK: return %[[GENERIC]] @@ -161,7 +161,7 @@ : (!pcf.sref<8x16xf32, sync(#pcf.test_scope)>) -> (tensor<8x16xf32>) { %slice = pcf.read_slice %ref[%id0, %id1] [4, 8] [1, 1] : !pcf.sref<8x16xf32, sync(#pcf.test_scope)> to tensor<4x8xf32> - %result = linalg.exp ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> + %result = linalg.elementwise <exp> ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> pcf.write_slice %result into %ref[%id0, %id1] [4, 8] [1, 1] : tensor<4x8xf32> into !pcf.sref<8x16xf32, sync(#pcf.test_scope)> pcf.return } @@ -193,7 +193,7 @@ -> (tensor<8x16xf32>, tensor<4x8xf32>) { %slice0 = pcf.read_slice %ref0[%id0, %id1] [4, 8] [1, 1] : !pcf.sref<8x16xf32, sync(#pcf.test_scope)> to tensor<4x8xf32> %slice1 = pcf.read_slice %ref1[0, 0] [4, 8] [1, 1] : !pcf.sref<4x8xf32, sync(#pcf.test_scope)> to tensor<4x8xf32> - %add = linalg.add ins(%slice0, %slice1 : tensor<4x8xf32>, tensor<4x8xf32>) outs(%slice1 : tensor<4x8xf32>) -> tensor<4x8xf32> + %add = linalg.elementwise <add> ins(%slice0, %slice1 : tensor<4x8xf32>, tensor<4x8xf32>) outs(%slice1 : tensor<4x8xf32>) -> tensor<4x8xf32> pcf.write_slice %add into %ref0[%id0, %id1] [4, 8] [1, 1] : tensor<4x8xf32> into !pcf.sref<8x16xf32, sync(#pcf.test_scope)> pcf.return } @@ -210,7 +210,7 @@ // CHECK-NEXT: execute(%[[REF0:.+]] = %[[ARG0]], %[[REF1:.+]] = %[[DEST]]) // CHECK: pcf.read_slice %[[REF0]][%{{.+}}, %{{.+}}] [4, 8] [1, 1] // CHECK: %[[TILED_FILL:.+]] = linalg.fill ins(%[[CST]]{{.*}} outs(%[[DEST]] : tensor<4x8xf32>) -// CHECK: linalg.add +// CHECK: linalg.elementwise <add> // CHECK: pcf.write_slice // CHECK: pcf.return @@ -264,7 +264,7 @@ : (!pcf.sref<8x16xf32, #pcf.test_scope>) -> (tensor<8x16xf32>) { %slice = pcf.read_slice %ref[%id0, %id1] [4, 8] [1, 1] : !pcf.sref<8x16xf32, #pcf.test_scope> to tensor<4x8xf32> - %result = linalg.exp ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> + %result = linalg.elementwise <exp> ins(%slice : tensor<4x8xf32>) outs(%slice : tensor<4x8xf32>) -> tensor<4x8xf32> pcf.write_slice %result into %ref[%id0, %id1] [4, 8] [1, 1] : tensor<4x8xf32> into !pcf.sref<8x16xf32, #pcf.test_scope> pcf.return }
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/BUILD.bazel b/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/BUILD.bazel index 00d0053..a1f6917 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/BUILD.bazel +++ b/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/BUILD.bazel
@@ -82,6 +82,7 @@ "//compiler/src/iree/compiler/Utils", "@llvm-project//llvm:Support", "@llvm-project//mlir:AffineDialect", + "@llvm-project//mlir:ArithDialect", "@llvm-project//mlir:ControlFlowInterfaces", "@llvm-project//mlir:DialectUtils", "@llvm-project//mlir:IR",
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/CMakeLists.txt b/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/CMakeLists.txt index cffa3b3..48399d0 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/CMakeLists.txt +++ b/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/CMakeLists.txt
@@ -43,6 +43,7 @@ ::IREEVectorExtOpsGen LLVMSupport MLIRAffineDialect + MLIRArithDialect MLIRControlFlowInterfaces MLIRIR MLIRSideEffectInterfaces
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/VectorExtOps.cpp b/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/VectorExtOps.cpp index 9485176..5d4b8db 100644 --- a/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/VectorExtOps.cpp +++ b/compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/VectorExtOps.cpp
@@ -7,6 +7,7 @@ #include "iree/compiler/Codegen/Dialect/VectorExt/IR/VectorExtOps.h" #include "llvm/Support/InterleavedRange.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" +#include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/Dialect/Utils/IndexingUtils.h" #include "mlir/Dialect/Vector/IR/VectorOps.h" #include "mlir/IR/Matchers.h"
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.cpp index a145397..4936d6d 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.cpp +++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/DispatchABI.cpp
@@ -1160,11 +1160,14 @@ /*context=*/importFunc.second, /*reserved=*/nullPtrValue, }; - auto callOp = - LLVM::CallOp::create(builder, loc, TypeRange{builder.getI32Type()}, args); - callOp.getProperties().operandSegmentSizes = { - static_cast<int32_t>(args.size()), 0}; - callOp.getProperties().op_bundle_sizes = builder.getDenseI32ArrayAttr({}); + LLVM::CallOp::Properties properties; + LLVM::CallOp::populateDefaultProperties( + OperationName(LLVM::CallOp::getOperationName(), builder.getContext()), + properties); + properties.operandSegmentSizes = {static_cast<int32_t>(args.size()), 0}; + properties.op_bundle_sizes = builder.getDenseI32ArrayAttr({}); + auto callOp = LLVM::CallOp::create( + builder, loc, TypeRange{builder.getI32Type()}, args, properties); return callOp.getResult(); } @@ -1331,9 +1334,10 @@ Location loc = forOp->getLoc(); if (cConv == IREE::HAL::CallingConvention::Default) { - auto callOp = - LLVM::CallOp::create(rewriter, loc, abiFunctionType->getReturnTypes(), - allArgsList, forOp->getAttrs()); + auto callOp = LLVM::CallOp::create( + rewriter, loc, abiFunctionType->getReturnTypes(), allArgsList, + cast<LLVM::CallOp>(forOp).getProperties(), + forOp->getDiscardableAttrDictionary().getValue()); callOp.getProperties().operandSegmentSizes = { static_cast<int32_t>(allArgsList.size()), 0}; callOp.getProperties().op_bundle_sizes = rewriter.getDenseI32ArrayAttr({});
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/ExpandF16OpToF32.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/ExpandF16OpToF32.cpp index 65321dc..c7e59c4 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMCPU/ExpandF16OpToF32.cpp +++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/ExpandF16OpToF32.cpp
@@ -43,10 +43,13 @@ operands.push_back(operand); continue; } - Value ext = arith::ExtFOp::create(rewriter, loc, f32Type, operand); + Value ext = arith::ExtFOp::create(rewriter, loc, f32Type, operand, + arith::FastMathFlagsAttr{}); operands.push_back(ext); } - Value newOp = Op::create(rewriter, loc, f32Type, operands); + Value newOp = Op::create(rewriter, loc, TypeRange{f32Type}, operands, + op.getProperties(), + op->getDiscardableAttrDictionary().getValue()); rewriter.replaceOpWithNewOp<arith::TruncFOp>(op, resultType, newOp); return success();
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/apply_scale_lowering.mlir b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/apply_scale_lowering.mlir index 82345a8..379eead 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/apply_scale_lowering.mlir +++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/apply_scale_lowering.mlir
@@ -172,7 +172,7 @@ // CHECK-LABEL: llvm.func @apply_scale_zve32x // CHECK-DAG: %[[RHS:.+]] = llvm.mlir.constant(dense<19689> : vector<2xi32>) : vector<2xi32> // CHECK-DAG: %[[RHSEXT:.+]] = llvm.mlir.constant(dense<19689> : vector<2xi64>) : vector<2xi64> -// CHECK: %[[LHS:.+]] = llvm.load %{{.+}} {alignment = 4 : i64} : !llvm.ptr -> vector<2xi32> +// CHECK: %[[LHS:.+]] = llvm.load %{{.+}} <alignment = 4> : !llvm.ptr -> vector<2xi32> // CHECK: %[[LHSEXT:.+]] = llvm.sext %[[LHS]] : vector<2xi32> to vector<2xi64> // CHECK: %[[MULEXT:.*]] = llvm.mul %[[LHSEXT]], %[[RHSEXT]] : vector<2xi64> // CHECK: %[[MULLOW:.*]] = llvm.mul %[[LHS]], %[[RHS]] : vector<2xi32> @@ -222,7 +222,7 @@ // CHECK-LABEL: llvm.func @apply_scale_zve32f // CHECK-DAG: %[[RHS:.+]] = llvm.mlir.constant(dense<19689> : vector<2xi32>) : vector<2xi32> // CHECK-DAG: %[[RHSEXT:.+]] = llvm.mlir.constant(dense<19689> : vector<2xi64>) : vector<2xi64> -// CHECK: %[[LHS:.+]] = llvm.load %{{.+}} {alignment = 4 : i64} : !llvm.ptr -> vector<2xi32> +// CHECK: %[[LHS:.+]] = llvm.load %{{.+}} <alignment = 4> : !llvm.ptr -> vector<2xi32> // CHECK: %[[LHSEXT:.+]] = llvm.sext %[[LHS]] : vector<2xi32> to vector<2xi64> // CHECK: %[[MULEXT:.*]] = llvm.mul %[[LHSEXT]], %[[RHSEXT]] : vector<2xi64> // CHECK: %[[MULLOW:.*]] = llvm.mul %[[LHS]], %[[RHS]] : vector<2xi32>
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/convert_to_llvm.mlir b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/convert_to_llvm.mlir index f6b2b48..d996359 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/convert_to_llvm.mlir +++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/convert_to_llvm.mlir
@@ -182,7 +182,7 @@ } // CHECK-LABEL: llvm.func @negative_no_gather_lowering( -// CHECK: llvm.intr.masked.gather {{.*}} : (vector<64x!llvm.ptr>, vector<64xi1>, vector<64xf32>) -> vector<64xf32> +// CHECK: llvm.intr.masked.gather({{.*}}), alignment({{[0-9]+}}) : (vector<64x!llvm.ptr>, vector<64xi1>, vector<64xf32>) -> vector<64xf32> // -----
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/pipeline_bf16_native_converts_tests.mlir b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/pipeline_bf16_native_converts_tests.mlir index 25aeb5a..711985b 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/pipeline_bf16_native_converts_tests.mlir +++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/pipeline_bf16_native_converts_tests.mlir
@@ -26,7 +26,7 @@ %lhs = iree_tensor_ext.dispatch.tensor.load %0, offsets = [0], sizes = [1024], strides = [1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<1024xbf16>> -> tensor<1024xbf16> %rhs = iree_tensor_ext.dispatch.tensor.load %1, offsets = [0], sizes = [1024], strides = [1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<1024xbf16>> -> tensor<1024xbf16> %init = tensor.empty() : tensor<1024xbf16> - %add = linalg.add ins(%lhs, %rhs : tensor<1024xbf16>, tensor<1024xbf16>) outs(%init : tensor<1024xbf16>) -> tensor<1024xbf16> + %add = linalg.elementwise <add> ins(%lhs, %rhs : tensor<1024xbf16>, tensor<1024xbf16>) outs(%init : tensor<1024xbf16>) -> tensor<1024xbf16> iree_tensor_ext.dispatch.tensor.store %add, %2, offsets = [0], sizes = [1024], strides = [1] : tensor<1024xbf16> -> !iree_tensor_ext.dispatch.tensor<writeonly:tensor<1024xbf16>> return }
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorLowering.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorLowering.cpp index ef38d14..42223f8 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorLowering.cpp +++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorLowering.cpp
@@ -83,7 +83,8 @@ } if (isa<FloatType>(dstElementType)) { - return arith::ExtFOp::create(rewriter, loc, promotedType, v); + return arith::ExtFOp::create(rewriter, loc, promotedType, v, + arith::FastMathFlagsAttr{}); } // For integer types, vector.contract only supports signless integer types // and promotion happens via sign extension. @@ -214,13 +215,15 @@ if (lhsVecType.getElementType() != elemType) { Type promotedType = lhsVecType.clone(elemType); - lhs = arith::ExtFOp::create(rewriter, loc, promotedType, lhs); + lhs = arith::ExtFOp::create(rewriter, loc, promotedType, lhs, + arith::FastMathFlagsAttr{}); lhsVecType = cast<VectorType>(lhs.getType()); } if (rhsVecType.getElementType() != elemType) { Type promotedType = rhsVecType.clone(elemType); - rhs = arith::ExtFOp::create(rewriter, loc, promotedType, rhs); + rhs = arith::ExtFOp::create(rewriter, loc, promotedType, rhs, + arith::FastMathFlagsAttr{}); rhsVecType = cast<VectorType>(rhs.getType()); }
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse.mlir index ecb26f6..7aa1bc3 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse.mlir +++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse.mlir
@@ -247,14 +247,14 @@ %c0 = arith.constant 0 : index %8 = tensor.dim %11, %c0 : tensor<?x256xf16> %13 = tensor.empty(%8) : tensor<?x256xf16> - %15 = linalg.add ins(%11, %12 : tensor<?x256xf16>, tensor<?x256xf16>) outs(%13 : tensor<?x256xf16>) -> tensor<?x256xf16> + %15 = linalg.elementwise <add> ins(%11, %12 : tensor<?x256xf16>, tensor<?x256xf16>) outs(%13 : tensor<?x256xf16>) -> tensor<?x256xf16> return %15 : tensor<?x256xf16> } } // CHECK-LABEL: func.func @elementwise_dynamic_dim( // CHECK-SAME: #iree_codegen.translation_info<pipeline = #iree_gpu.pipeline<TileAndFuse> workgroup_size = [64, 1, 1] subgroup_size = 64> -// CHECK: linalg.add {{.*}}lowering_config = #iree_gpu.lowering_config +// CHECK: linalg.elementwise <add> {{.*}}lowering_config = #iree_gpu.lowering_config // CHECK-SAME: thread = [1, 4] // CHECK-SAME: workgroup = [1, 256] @@ -262,7 +262,7 @@ func.func @elementwise_unaligned(%11: tensor<180x180xf16>, %12: tensor<180x180xf16>) -> tensor<180x180xf16> { %13 = tensor.empty() : tensor<180x180xf16> - %15 = linalg.add ins(%11, %12 : tensor<180x180xf16>, tensor<180x180xf16>) outs(%13 : tensor<180x180xf16>) -> tensor<180x180xf16> + %15 = linalg.elementwise <add> ins(%11, %12 : tensor<180x180xf16>, tensor<180x180xf16>) outs(%13 : tensor<180x180xf16>) -> tensor<180x180xf16> return %15 : tensor<180x180xf16> } @@ -273,7 +273,7 @@ func.func @elementwise_large_rank(%11: tensor<3x5x7x11x13x17x19x23xf16>, %12: tensor<3x5x7x11x13x17x19x23xf16>) -> tensor<3x5x7x11x13x17x19x23xf16> { %13 = tensor.empty() : tensor<3x5x7x11x13x17x19x23xf16> - %15 = linalg.add ins(%11, %12 : tensor<3x5x7x11x13x17x19x23xf16>, tensor<3x5x7x11x13x17x19x23xf16>) outs(%13 : tensor<3x5x7x11x13x17x19x23xf16>) -> tensor<3x5x7x11x13x17x19x23xf16> + %15 = linalg.elementwise <add> ins(%11, %12 : tensor<3x5x7x11x13x17x19x23xf16>, tensor<3x5x7x11x13x17x19x23xf16>) outs(%13 : tensor<3x5x7x11x13x17x19x23xf16>) -> tensor<3x5x7x11x13x17x19x23xf16> return %15 : tensor<3x5x7x11x13x17x19x23xf16> } @@ -562,7 +562,7 @@ %arg2: tensor<3x1xi32>) -> tensor<3x2048x2048xf32> { %cst = arith.constant 0.000000e+00 : f32 %0 = tensor.empty() : tensor<3x2048x2048xf32> - %1 = linalg.add ins(%arg0, %arg1 : tensor<3x2048x2048xf32>, tensor<3x2048x2048xf32>) + %1 = linalg.elementwise <add> ins(%arg0, %arg1 : tensor<3x2048x2048xf32>, tensor<3x2048x2048xf32>) outs(%0 : tensor<3x2048x2048xf32>) -> tensor<3x2048x2048xf32> %2 = iree_linalg_ext.scatter dimension_map = [0] unique_indices(true) ins(%1, %arg2 : tensor<3x2048x2048xf32>, tensor<3x1xi32>) outs(%0 : tensor<3x2048x2048xf32>) { @@ -575,7 +575,7 @@ // CHECK-LABEL: func.func @elementwise_scatter( // CHECK-SAME: #iree_codegen.translation_info<pipeline = #iree_gpu.pipeline<TileAndFuse> workgroup_size = [64, 1, 1] subgroup_size = 64 -// CHECK: linalg.add {{.*}}lowering_config = #iree_gpu.lowering_config +// CHECK: linalg.elementwise <add> {{.*}}lowering_config = #iree_gpu.lowering_config // CHECK-SAME: thread = [1, 1, 4] // CHECK-SAME: workgroup = [1, 1, 256]
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_tile_and_fuse.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_tile_and_fuse.mlir index c5c20f7..3449a2c 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_tile_and_fuse.mlir +++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/pipeline_tile_and_fuse.mlir
@@ -608,7 +608,7 @@ %3 = iree_tensor_ext.dispatch.tensor.load %0, offsets = [0, 0], sizes = [1, 3], strides = [1, 1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<1x3xf32>> -> tensor<1x3xf32> %4 = iree_tensor_ext.dispatch.tensor.load %1, offsets = [0, 0], sizes = [1, 3], strides = [1, 1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<1x3xf32>> -> tensor<1x3xf32> %5 = tensor.empty() : tensor<1x3xf32> - %6 = linalg.add {lowering_config = #lowering_config} ins(%3, %4 : tensor<1x3xf32>, tensor<1x3xf32>) outs(%5 : tensor<1x3xf32>) -> tensor<1x3xf32> + %6 = linalg.elementwise <add> {lowering_config = #lowering_config} ins(%3, %4 : tensor<1x3xf32>, tensor<1x3xf32>) outs(%5 : tensor<1x3xf32>) -> tensor<1x3xf32> iree_tensor_ext.dispatch.tensor.store %6, %2, offsets = [0, 0], sizes = [1, 3], strides = [1, 1] : tensor<1x3xf32> -> !iree_tensor_ext.dispatch.tensor<writeonly:tensor<1x3xf32>> return }
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir index 77fcac8..8f5337b 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir +++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
@@ -422,8 +422,8 @@ } // CHECK-LABEL: llvm.func @masked_load_store // CHECK: %[[MASK_BIT:.+]] = llvm.icmp "sgt" {{.*}} : vector<1xi64> -// CHECK: llvm.intr.masked.load %{{.*}}, %[[MASK_BIT]] -// CHECK: llvm.intr.masked.store %{{.*}}, %[[MASK_BIT]] +// CHECK: llvm.intr.masked.load(%{{.*}}, %[[MASK_BIT]] +// CHECK: llvm.intr.masked.store(%{{.*}}, %{{.*}}, %[[MASK_BIT]] // ----- // Test workgroup size lowering
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir index a15ca43..beab53c 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir +++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
@@ -95,8 +95,8 @@ } // CHECK-LABEL: llvm.func @masked_load_store // CHECK: %[[MASK_BIT:.+]] = llvm.icmp "sgt" {{.*}} : vector<1xi64> -// CHECK: llvm.intr.masked.load %{{.*}}, %[[MASK_BIT]] -// CHECK: llvm.intr.masked.store %{{.*}}, %[[MASK_BIT]] +// CHECK: llvm.intr.masked.load(%{{.*}}, %[[MASK_BIT]] +// CHECK: llvm.intr.masked.store(%{{.*}}, %{{.*}}, %[[MASK_BIT]] // ----- // Test workgroup size lowering
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir index aa1bfb2..d47baba 100644 --- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir +++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
@@ -198,7 +198,7 @@ } // CHECK-LABEL: llvm.func @reduction -// CHECK: "llvm.intr.vector.reduce.fadd"({{.*}}) {{.*}} : (f32, vector<4xf32>) -> f32 +// CHECK: llvm.intr.vector.reduce.fadd({{.*}}){{.*}} : (f32, vector<4xf32>) -> f32 // ----- @@ -264,7 +264,7 @@ } // CHECK-LABEL: llvm.func @vector_reduction_dispatch -// CHECK-COUNT-4: "llvm.intr.vector.reduce.fadd"({{.*}}) {{.*}} : (f32, vector<4xf32>) -> f32 +// CHECK-COUNT-4: llvm.intr.vector.reduce.fadd({{.*}}){{.*}} : (f32, vector<4xf32>) -> f32 // CHECK: llvm.store %{{.*}} : vector<4xf32>, !llvm.ptr<1> // ----- @@ -376,7 +376,7 @@ // CHECK: llvm.load {{.*}} : !llvm.ptr<3> -> vector<1xf32> // CHECK-COUNT-2: nvvm.shfl.sync bfly // CHECK: llvm.fdiv %{{.*}}, %{{.*}} -// CHECK: llvm.store %{{.*}}, %{{.*}} {alignment = 4 : i64} : vector<4xf32>, !llvm.ptr<1> +// CHECK: llvm.store %{{.*}}, %{{.*}} <alignment = 4> : vector<4xf32>, !llvm.ptr<1> // ----- @@ -434,8 +434,8 @@ // Check that bufferization is emitting correct code for the temp shared // memory alloc. // SM80-LABEL: llvm.func @shared_mem_transpose -// SM80: llvm.load %{{.*}} {alignment = 4 : i64} : !llvm.ptr<1> -> vector<4xf32> -// SM80: llvm.store %{{.*}}, %{{.*}} {alignment = 4 : i64} : vector<4xf32>, !llvm.ptr<3> +// SM80: llvm.load %{{.*}} <alignment = 4> : !llvm.ptr<1> -> vector<4xf32> +// SM80: llvm.store %{{.*}}, %{{.*}} <alignment = 4> : vector<4xf32>, !llvm.ptr<3> // -----
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Conversion/TypeConverter.cpp b/compiler/src/iree/compiler/Dialect/HAL/Conversion/TypeConverter.cpp index 9127ef5..9d71572 100644 --- a/compiler/src/iree/compiler/Dialect/HAL/Conversion/TypeConverter.cpp +++ b/compiler/src/iree/compiler/Dialect/HAL/Conversion/TypeConverter.cpp
@@ -41,7 +41,9 @@ ValueRange inputs, Location loc) -> Value { assert(inputs.size() == 1); if (isa<TensorType>(inputs[0].getType())) { - return IREE::HAL::TensorExportOp::create(builder, loc, type, inputs[0]); + return IREE::HAL::TensorExportOp::create( + builder, loc, TypeRange{type}, ValueRange{inputs[0]}, + IREE::HAL::TensorExportOp::Properties{}); } else if (isa<IREE::HAL::BufferViewType>(inputs[0].getType())) { return IREE::HAL::BufferViewBufferOp::create(builder, loc, type, inputs[0]); @@ -58,7 +60,9 @@ auto inputValue = inputs[0]; auto inputType = inputValue.getType(); if (isa<TensorType>(inputType)) { - return IREE::HAL::TensorExportOp::create(builder, loc, type, inputValue); + return IREE::HAL::TensorExportOp::create( + builder, loc, TypeRange{type}, ValueRange{inputValue}, + IREE::HAL::TensorExportOp::Properties{}); } else if (isa<IREE::HAL::BufferType>(inputType)) { // Look for the buffer view this buffer came from, if any. // If we don't have the origin buffer view then we can't know the shape
diff --git a/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/DecomposeMapStore.cpp b/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/DecomposeMapStore.cpp index 9b105cf..0be8fc2 100644 --- a/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/DecomposeMapStore.cpp +++ b/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/DecomposeMapStore.cpp
@@ -357,7 +357,8 @@ SmallVector<Value> operands = {flatOutputBuffer, offsets[0], indexVector, maskVector, inputVector}; rewriter.replaceOpWithNewOp<vector::ScatterOp>( - mapStoreOp, /*resultTypes=*/TypeRange{}, operands); + mapStoreOp, /*resultTypes=*/TypeRange{}, operands, + vector::ScatterOp::Properties{}); return success(); }
diff --git a/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/SplitReduction.cpp b/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/SplitReduction.cpp index 5065a4e..402009f 100644 --- a/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/SplitReduction.cpp +++ b/compiler/src/iree/compiler/Dialect/LinalgExt/Transforms/SplitReduction.cpp
@@ -683,7 +683,8 @@ unsigned outBitWidth = outType.getIntOrFloatBitWidth(); if (outBitWidth > inBitWidth) { - inCast = arith::ExtFOp::create(b, loc, outVal.getType(), in); + inCast = arith::ExtFOp::create(b, loc, outVal.getType(), in, + arith::FastMathFlagsAttr{}); } else if (outBitWidth < inBitWidth) { inCast = arith::TruncFOp::create(b, loc, outVal.getType(), in); }
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/Patterns.cpp b/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/Patterns.cpp index f4ff1dc..cda7c6b 100644 --- a/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/Patterns.cpp +++ b/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/Patterns.cpp
@@ -65,7 +65,8 @@ // Transfer to unknown lifetime. auto unknownType = rewriter.getType<IREE::Stream::ResourceType>(); auto constantSize = rewriter.createOrFold<IREE::Stream::ResourceSizeOp>( - constantOp.getLoc(), rewriter.getIndexType(), newOp.getResult()); + constantOp.getLoc(), rewriter.getIndexType(), newOp.getResult(), + /*affinity=*/IREE::Stream::AffinityAttr{}); auto transferOp = IREE::Stream::AsyncTransferOp::create( rewriter, constantOp.getLoc(), unknownType, newOp.getResult(), constantSize, constantSize, @@ -116,7 +117,8 @@ // Transfer to unknown lifetime. auto unknownType = rewriter.getType<IREE::Stream::ResourceType>(); auto constantSize = rewriter.createOrFold<IREE::Stream::ResourceSizeOp>( - constantOp.getLoc(), rewriter.getIndexType(), newOp.getResult()); + constantOp.getLoc(), rewriter.getIndexType(), newOp.getResult(), + /*affinity=*/IREE::Stream::AffinityAttr{}); auto transferOp = IREE::Stream::AsyncTransferOp::create( rewriter, constantOp.getLoc(), unknownType, newOp.getResult(), constantSize, constantSize,
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Conversion/StandardToStream/Patterns.cpp b/compiler/src/iree/compiler/Dialect/Stream/Conversion/StandardToStream/Patterns.cpp index 690913f..5c3c198 100644 --- a/compiler/src/iree/compiler/Dialect/Stream/Conversion/StandardToStream/Patterns.cpp +++ b/compiler/src/iree/compiler/Dialect/Stream/Conversion/StandardToStream/Patterns.cpp
@@ -62,7 +62,8 @@ auto unknownType = rewriter.getType<IREE::Stream::ResourceType>(); auto constantSize = rewriter.createOrFold<IREE::Stream::ResourceSizeOp>( - constantOp.getLoc(), rewriter.getIndexType(), newOp.getResult()); + constantOp.getLoc(), rewriter.getIndexType(), newOp.getResult(), + /*affinity=*/IREE::Stream::AffinityAttr{}); auto transferOp = IREE::Stream::AsyncTransferOp::create( rewriter, constantOp.getLoc(), unknownType, newOp.getResult(), constantSize, constantSize,
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Conversion/UtilToStream/Patterns.cpp b/compiler/src/iree/compiler/Dialect/Stream/Conversion/UtilToStream/Patterns.cpp index 00d8e28..b5c5956 100644 --- a/compiler/src/iree/compiler/Dialect/Stream/Conversion/UtilToStream/Patterns.cpp +++ b/compiler/src/iree/compiler/Dialect/Stream/Conversion/UtilToStream/Patterns.cpp
@@ -273,7 +273,8 @@ TypeAttr::get(globalOp.getType()), /*result_encoding_dims=*/ValueRange{}, affinityAttr); initialValueSize = IREE::Stream::ResourceSizeOp::create( - rewriter, globalOp.getLoc(), indexType, initialValue); + rewriter, globalOp.getLoc(), indexType, initialValue, + /*affinity=*/IREE::Stream::AffinityAttr{}); } resourceOp.createStoreOp(globalOp.getLoc(), initialValue, rewriter); resourceSizeOp.createStoreOp(globalOp.getLoc(), initialValueSize,
diff --git a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamTypes.cpp b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamTypes.cpp index 7bbb612..4fa44d6 100644 --- a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamTypes.cpp +++ b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamTypes.cpp
@@ -560,7 +560,8 @@ Value ResourceType::inferSizeFromValue(Location loc, Value value, OpBuilder &builder) const { return builder.createOrFold<IREE::Stream::ResourceSizeOp>( - loc, builder.getIndexType(), value); + loc, builder.getIndexType(), value, + /*affinity=*/IREE::Stream::AffinityAttr{}); } Value ResourceType::createSubrangeOp(Location loc, Value resource,
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Transforms/PropagateTimepoints.cpp b/compiler/src/iree/compiler/Dialect/Stream/Transforms/PropagateTimepoints.cpp index 2d15adf..3b35447 100644 --- a/compiler/src/iree/compiler/Dialect/Stream/Transforms/PropagateTimepoints.cpp +++ b/compiler/src/iree/compiler/Dialect/Stream/Transforms/PropagateTimepoints.cpp
@@ -240,7 +240,8 @@ // If we couldn't find anything we could use we'll insert the size query. The // hope is that more program analysis could take care of this for us. - return IREE::Stream::ResourceSizeOp::create(builder, loc, resourceValue); + return IREE::Stream::ResourceSizeOp::create( + builder, loc, resourceValue, /*affinity=*/IREE::Stream::AffinityAttr{}); } // Recursively expands resources into (timepoint, resource) pairs within the @@ -355,7 +356,8 @@ replacementExceptions.insert(resultSize.getDefiningOp()); } else { auto sizeOp = IREE::Stream::ResourceSizeOp::create( - builder, op.getLoc(), op.getLoadedGlobalValue()); + builder, op.getLoc(), op.getLoadedGlobalValue(), + /*affinity=*/IREE::Stream::AffinityAttr{}); replacementExceptions.insert(sizeOp); resultSize = sizeOp.getResult(); } @@ -483,9 +485,10 @@ auto newResult = newOp.getResult(newIdx++); auto newTimepoint = newOp.getResult(newIdx++); resourceTimepointMap.map(newResult, newTimepoint); - auto newResultSize = - IREE::Stream::ResourceSizeOp::create(builder, op.getLoc(), newResult) - .getResult(); + auto newResultSize = IREE::Stream::ResourceSizeOp::create( + builder, op.getLoc(), newResult, + /*affinity=*/IREE::Stream::AffinityAttr{}) + .getResult(); auto awaitOp = IREE::Stream::TimepointAwaitOp::create( builder, op.getLoc(), newResult, newResultSize, newTimepoint); oldResult.replaceAllUsesWith(awaitOp.getResults().front());
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp b/compiler/src/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp index 2ea9618..e7975c4 100644 --- a/compiler/src/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp +++ b/compiler/src/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp
@@ -170,7 +170,7 @@ // placeholder (as we may later decide it's ok to transition on a // particular device). auto resultSize = rewriter.createOrFold<IREE::Stream::ResourceSizeOp>( - op->getLoc(), result); + op->getLoc(), result, /*affinity=*/IREE::Stream::AffinityAttr{}); auto affinityAttr = getOpAffinity(op); auto transferOp = IREE::Stream::AsyncTransferOp::create( rewriter, op->getLoc(), newType, result, resultSize, resultSize,
diff --git a/compiler/src/iree/compiler/Dialect/Util/Conversion/ConversionPatterns.h b/compiler/src/iree/compiler/Dialect/Util/Conversion/ConversionPatterns.h index 1c01336..ca403a3 100644 --- a/compiler/src/iree/compiler/Dialect/Util/Conversion/ConversionPatterns.h +++ b/compiler/src/iree/compiler/Dialect/Util/Conversion/ConversionPatterns.h
@@ -33,21 +33,34 @@ newResultTypes.push_back(newTypes.front()); } - SmallVector<NamedAttribute> newAttrs; - if (failed(convertTypeAttributes(op->getAttrs(), newAttrs))) { + ArrayRef<NamedAttribute> oldAttrs = + op->getDiscardableAttrDictionary().getValue(); + // Null for ops with empty properties; such ops reject a properties attr. + auto propsDict = + dyn_cast_if_present<DictionaryAttr>(op->getPropertiesAsAttribute()); + ArrayRef<NamedAttribute> oldProps; + if (propsDict) { + oldProps = propsDict.getValue(); + } + SmallVector<NamedAttribute> newAttrs, newProps; + if (failed(convertTypeAttributes(oldAttrs, newAttrs)) || + failed(convertTypeAttributes(oldProps, newProps))) { return rewriter.notifyMatchFailure(op, "failed converting type attributes"); } if (newResultTypes == op->getResultTypes() && - op->getOperands() == adaptor.getOperands() && - newAttrs == op->getAttrs()) { + op->getOperands() == adaptor.getOperands() && newAttrs == oldAttrs && + newProps == oldProps) { return rewriter.notifyMatchFailure(op, "op does not need transformation"); } - auto newOp = T::create(rewriter, op.getLoc(), newResultTypes, - adaptor.getOperands(), newAttrs); - rewriter.replaceOp(op, newOp->getResults()); + OperationState state(op.getLoc(), T::getOperationName(), + adaptor.getOperands(), newResultTypes, newAttrs); + if (propsDict) { + state.propertiesAttr = rewriter.getDictionaryAttr(newProps); + } + rewriter.replaceOp(op, rewriter.create(state)); return success(); }
diff --git a/compiler/src/iree/compiler/Dialect/Util/Transforms/OptimizeIntArithmetic.cpp b/compiler/src/iree/compiler/Dialect/Util/Transforms/OptimizeIntArithmetic.cpp index a6635f9..558d0b6 100644 --- a/compiler/src/iree/compiler/Dialect/Util/Transforms/OptimizeIntArithmetic.cpp +++ b/compiler/src/iree/compiler/Dialect/Util/Transforms/OptimizeIntArithmetic.cpp
@@ -105,8 +105,11 @@ if (failed(staticallyLegalToConvertToUnsignedOp(solver, op, indexIsI64))) { return failure(); } - rewriter.replaceOpWithNewOp<Unsigned>(op, op->getResultTypes(), - op->getOperands(), op->getAttrs()); + OperationState state(op.getLoc(), Unsigned::getOperationName(), + op->getOperands(), op->getResultTypes(), + op->getDiscardableAttrDictionary().getValue()); + state.propertiesAttr = op->getPropertiesAsAttribute(); + rewriter.replaceOp(op, rewriter.create(state)); return success(); }
diff --git a/compiler/src/iree/compiler/DispatchCreation/test/fuse_encoding_ops_into_dispatch_regions.mlir b/compiler/src/iree/compiler/DispatchCreation/test/fuse_encoding_ops_into_dispatch_regions.mlir index 8e9f8ba..ffc81ea 100644 --- a/compiler/src/iree/compiler/DispatchCreation/test/fuse_encoding_ops_into_dispatch_regions.mlir +++ b/compiler/src/iree/compiler/DispatchCreation/test/fuse_encoding_ops_into_dispatch_regions.mlir
@@ -164,7 +164,7 @@ %cst = arith.constant 0.000000e+00 : f32 %0 = tensor.empty() : tensor<32x32xf32> %1 = flow.dispatch.region -> (tensor<32x32xf32>) { - %3 = linalg.add ins(%arg0, %arg0 : tensor<32x32xf32>, tensor<32x32xf32>) + %3 = linalg.elementwise <add> ins(%arg0, %arg0 : tensor<32x32xf32>, tensor<32x32xf32>) outs(%0 : tensor<32x32xf32>) -> tensor<32x32xf32> flow.return %3 : tensor<32x32xf32> } @@ -176,7 +176,7 @@ // CHECK: #[[$ENCODING:.+]] = #iree_encoding.testing<> // CHECK-LABEL: @reshape_fusion // CHECK: %[[DISPATCH0:.+]] = flow.dispatch.region -> (tensor<16x64xf32, #[[$ENCODING]]>) -// CHECK: linalg.add +// CHECK: linalg.elementwise <add> // CHECK: tensor.collapse_shape // CHECK: tensor.expand_shape // CHECK: %[[SET_ENCODING:.+]] = iree_encoding.set_encoding
diff --git a/compiler/src/iree/compiler/DispatchCreation/test/hoist_encoding_ops.mlir b/compiler/src/iree/compiler/DispatchCreation/test/hoist_encoding_ops.mlir index e205928..7751424 100644 --- a/compiler/src/iree/compiler/DispatchCreation/test/hoist_encoding_ops.mlir +++ b/compiler/src/iree/compiler/DispatchCreation/test/hoist_encoding_ops.mlir
@@ -201,7 +201,7 @@ %init = tensor.empty() : tensor<640x320xf32> %0 = flow.dispatch.region -> (tensor<640x320xf32>) { %input = util.global.load @weight : tensor<640x320xf32> - %1 = linalg.elementwise kind=#linalg.elementwise_kind<exp> + %1 = linalg.elementwise <exp> ins(%input: tensor<640x320xf32>) outs(%init: tensor<640x320xf32>) -> tensor<640x320xf32>
diff --git a/compiler/src/iree/compiler/ExternalInterfaces/HALExternalModels.cpp b/compiler/src/iree/compiler/ExternalInterfaces/HALExternalModels.cpp index a65d533..3272b1e 100644 --- a/compiler/src/iree/compiler/ExternalInterfaces/HALExternalModels.cpp +++ b/compiler/src/iree/compiler/ExternalInterfaces/HALExternalModels.cpp
@@ -170,8 +170,8 @@ // Create new import before the op. auto importOp = IREE::Stream::TimepointImportOp::create( builder, callOp.getLoc(), - builder.getType<IREE::Stream::TimepointType>(), - ValueRange{operand}); + builder.getType<IREE::Stream::TimepointType>(), ValueRange{operand}, + /*affinity=*/IREE::Stream::AffinityAttr{}); timepoints.push_back(importOp.getResultTimepoint()); } } @@ -216,8 +216,8 @@ // Create new import after the call. auto importOp = IREE::Stream::TimepointImportOp::create( builder, callOp.getLoc(), - builder.getType<IREE::Stream::TimepointType>(), - ValueRange{signalFence}); + builder.getType<IREE::Stream::TimepointType>(), ValueRange{signalFence}, + /*affinity=*/IREE::Stream::AffinityAttr{}); return importOp.getResultTimepoint(); } };
diff --git a/compiler/src/iree/compiler/GlobalOptimization/GeneralizeLinalgNamedOps.cpp b/compiler/src/iree/compiler/GlobalOptimization/GeneralizeLinalgNamedOps.cpp index cdaa420..3e36566 100644 --- a/compiler/src/iree/compiler/GlobalOptimization/GeneralizeLinalgNamedOps.cpp +++ b/compiler/src/iree/compiler/GlobalOptimization/GeneralizeLinalgNamedOps.cpp
@@ -44,13 +44,9 @@ namedOpCandidates.push_back(linalgOp); return; } - if (isa_and_nonnull<linalg::AbsOp, linalg::AddOp, linalg::BroadcastOp, - linalg::CeilOp, linalg::CopyOp, linalg::DivOp, - linalg::DivUnsignedOp, linalg::ExpOp, linalg::FloorOp, - linalg::LogOp, linalg::MapOp, linalg::MaxOp, - linalg::MulOp, linalg::NegFOp, linalg::ReduceOp, - linalg::SubOp, linalg::TransposeOp>( - linalgOp.getOperation()) || + if (isa_and_nonnull<linalg::BroadcastOp, linalg::CopyOp, + linalg::ElementwiseOp, linalg::MapOp, linalg::ReduceOp, + linalg::TransposeOp>(linalgOp.getOperation()) || linalg::isaConvolutionOpInterface(linalgOp)) { namedOpCandidates.push_back(linalgOp); } @@ -59,7 +55,7 @@ IRRewriter rewriter(&getContext()); for (auto linalgOp : namedOpCandidates) { rewriter.setInsertionPoint(linalgOp); - FailureOr<linalg::GenericOp> generalizedOp = + FailureOr<linalg::LinalgOp> generalizedOp = linalg::generalizeNamedOp(rewriter, linalgOp); if (failed(generalizedOp)) { linalgOp->emitOpError("failed to generalize operation");
diff --git a/compiler/src/iree/compiler/GlobalOptimization/PropagateLinalgTranspose.cpp b/compiler/src/iree/compiler/GlobalOptimization/PropagateLinalgTranspose.cpp index 5b65e42..c92fa0d 100644 --- a/compiler/src/iree/compiler/GlobalOptimization/PropagateLinalgTranspose.cpp +++ b/compiler/src/iree/compiler/GlobalOptimization/PropagateLinalgTranspose.cpp
@@ -157,7 +157,12 @@ } OpBuilder::InsertionGuard guard(rewriter); rewriter.setInsertionPoint(linalgOp); - return linalg::generalizeNamedOp(rewriter, linalgOp); + FailureOr<linalg::LinalgOp> generalizedOp = + linalg::generalizeNamedOp(rewriter, linalgOp); + if (failed(generalizedOp)) { + return failure(); + } + return cast<linalg::GenericOp>(generalizedOp->getOperation()); } //===----------------------------------------------------------------------===//
diff --git a/compiler/src/iree/compiler/GlobalOptimization/test/generalize_named_ops.mlir b/compiler/src/iree/compiler/GlobalOptimization/test/generalize_named_ops.mlir index 6ec0de9..2f8158c 100644 --- a/compiler/src/iree/compiler/GlobalOptimization/test/generalize_named_ops.mlir +++ b/compiler/src/iree/compiler/GlobalOptimization/test/generalize_named_ops.mlir
@@ -6,7 +6,7 @@ %d0 = tensor.dim %arg0, %c0 : tensor<?x?xf32> %d1 = tensor.dim %arg0, %c1 : tensor<?x?xf32> %empty = tensor.empty(%d0, %d1): tensor<?x?xf32> - %add = linalg.add ins(%arg0, %arg1 : tensor<?x?xf32>, tensor<?x?xf32>) + %add = linalg.elementwise <add> ins(%arg0, %arg1 : tensor<?x?xf32>, tensor<?x?xf32>) outs(%empty : tensor<?x?xf32>) -> tensor<?x?xf32> util.return %add : tensor<?x?xf32> } @@ -23,7 +23,7 @@ %d1 = tensor.dim %arg0, %c1 : tensor<?x?xf32> %dispatch = flow.dispatch.region[] -> (tensor<?x?xf32>{%d0, %d1}) { %empty = tensor.empty(%d0, %d1): tensor<?x?xf32> - %add = linalg.add ins(%arg0, %arg1 : tensor<?x?xf32>, tensor<?x?xf32>) + %add = linalg.elementwise <add> ins(%arg0, %arg1 : tensor<?x?xf32>, tensor<?x?xf32>) outs(%empty : tensor<?x?xf32>) -> tensor<?x?xf32> flow.return %add : tensor<?x?xf32> } @@ -31,7 +31,7 @@ } // CHECK-LABEL: util.func public @no_generalize_op_within_dispatch // CHECK: %[[DISPATCH:.+]] = flow.dispatch.region -// CHECK: %[[ADD:.+]] = linalg.add +// CHECK: %[[ADD:.+]] = linalg.elementwise <add> // CHECK: flow.return %[[ADD]] // CHECK: util.return %[[DISPATCH]]
diff --git a/compiler/src/iree/compiler/InputConversion/Common/ConvertPrimitiveType.cpp b/compiler/src/iree/compiler/InputConversion/Common/ConvertPrimitiveType.cpp index fe40994..7f0817f 100644 --- a/compiler/src/iree/compiler/InputConversion/Common/ConvertPrimitiveType.cpp +++ b/compiler/src/iree/compiler/InputConversion/Common/ConvertPrimitiveType.cpp
@@ -52,7 +52,8 @@ } if (inputETy.getIntOrFloatBitWidth() < eTy.getIntOrFloatBitWidth()) { - return arith::ExtFOp::create(builder, loc, type, inputs[0]); + return arith::ExtFOp::create(builder, loc, type, inputs[0], + arith::FastMathFlagsAttr{}); } return nullptr; @@ -238,7 +239,9 @@ resultEType.getWidth())) { return rewriter.notifyMatchFailure(op, "invalid width combination"); } - rewriter.replaceOpWithNewOp<OpTy>(op, resultType, op.getOperand()); + rewriter.replaceOpWithNewOp<OpTy>( + op, TypeRange{resultType}, ValueRange{op.getOperand()}, + op.getProperties(), op->getDiscardableAttrDictionary().getValue()); return success(); } };
diff --git a/compiler/src/iree/compiler/Preprocessing/Common/Convert1X1FilterConv2DToMatmul.cpp b/compiler/src/iree/compiler/Preprocessing/Common/Convert1X1FilterConv2DToMatmul.cpp index e72a9b3..f5a3f5a 100644 --- a/compiler/src/iree/compiler/Preprocessing/Common/Convert1X1FilterConv2DToMatmul.cpp +++ b/compiler/src/iree/compiler/Preprocessing/Common/Convert1X1FilterConv2DToMatmul.cpp
@@ -72,7 +72,8 @@ newMaps[0] = AffineMap::get(inputMap.getNumDims(), inputMap.getNumSymbols(), newExprs, rewriter.getContext()); - auto genericOp = linalg::generalizeNamedOp(rewriter, convOp).value(); + auto genericOp = cast<linalg::GenericOp>( + linalg::generalizeNamedOp(rewriter, convOp)->getOperation()); genericOp.setIndexingMapsAttr(rewriter.getAffineMapArrayAttr(newMaps)); return success(); }
diff --git a/compiler/src/iree/compiler/Preprocessing/Common/GeneralizeLinalgMatMul.cpp b/compiler/src/iree/compiler/Preprocessing/Common/GeneralizeLinalgMatMul.cpp index c58e2c0..f2f42e8 100644 --- a/compiler/src/iree/compiler/Preprocessing/Common/GeneralizeLinalgMatMul.cpp +++ b/compiler/src/iree/compiler/Preprocessing/Common/GeneralizeLinalgMatMul.cpp
@@ -39,7 +39,7 @@ for (auto linalgOp : namedOpCandidates) { rewriter.setInsertionPoint(linalgOp); - FailureOr<linalg::GenericOp> generalizedOp = + FailureOr<linalg::LinalgOp> generalizedOp = linalg::generalizeNamedOp(rewriter, linalgOp); if (failed(generalizedOp)) { linalgOp->emitOpError("failed to generalize operation");
diff --git a/compiler/src/iree/compiler/Preprocessing/Common/PromoteLinalgOutputs.cpp b/compiler/src/iree/compiler/Preprocessing/Common/PromoteLinalgOutputs.cpp index c9af88b..295d720 100644 --- a/compiler/src/iree/compiler/Preprocessing/Common/PromoteLinalgOutputs.cpp +++ b/compiler/src/iree/compiler/Preprocessing/Common/PromoteLinalgOutputs.cpp
@@ -45,7 +45,8 @@ rewriter, loc, TypeRange{promoteOutputType}, ValueRange{output}, ValueRange{promoteEmpty}, maps, iteratorTypes, [&](OpBuilder &b, Location loc, ValueRange args) { - Value result = arith::ExtFOp::create(b, loc, destType, args[0]); + Value result = arith::ExtFOp::create(b, loc, destType, args[0], + arith::FastMathFlagsAttr{}); linalg::YieldOp::create(b, loc, result); }) ->getResult(0);
diff --git a/compiler/src/iree/compiler/Preprocessing/Common/test/make_single_dispatch_for_function.mlir b/compiler/src/iree/compiler/Preprocessing/Common/test/make_single_dispatch_for_function.mlir index 1a55271..5362afe 100644 --- a/compiler/src/iree/compiler/Preprocessing/Common/test/make_single_dispatch_for_function.mlir +++ b/compiler/src/iree/compiler/Preprocessing/Common/test/make_single_dispatch_for_function.mlir
@@ -67,10 +67,10 @@ %0 = hal.tensor.import wait(%arg1) => %arg0 : !hal.buffer_view -> tensor<10xf32> %1 = hal.tensor.import wait(%arg3) => %arg2 : !hal.buffer_view -> tensor<10xf32> %2 = tensor.empty() : tensor<10xf32> - %3 = linalg.add ins(%0, %1 : tensor<10xf32>, tensor<10xf32>) + %3 = linalg.elementwise <add> ins(%0, %1 : tensor<10xf32>, tensor<10xf32>) outs(%2 : tensor<10xf32>) -> tensor<10xf32> %4 = hal.tensor.import wait(%arg5) => %arg4 : !hal.buffer_view -> tensor<10xf32> - %5 = linalg.add ins(%3, %4 : tensor<10xf32>, tensor<10xf32>) + %5 = linalg.elementwise <add> ins(%3, %4 : tensor<10xf32>, tensor<10xf32>) outs(%2 : tensor<10xf32>) -> tensor<10xf32> util.return %3, %5 : tensor<10xf32>, tensor<10xf32> } @@ -86,8 +86,8 @@ // CHECK-DAG: %[[IMPORT2:.+]] = hal.tensor.import wait(%[[ARG5]]) => %[[ARG4]] // CHECK: %[[DISPATCH:.+]]:2 = flow.dispatch.region // CHECK: %[[EMPTY:.+]] = tensor.empty() -// CHECK: %[[ADD1:.+]] = linalg.add ins(%[[IMPORT0]], %[[IMPORT1]] : -// CHECK: %[[ADD2:.+]] = linalg.add ins(%[[ADD1]], %[[IMPORT2]] : +// CHECK: %[[ADD1:.+]] = linalg.elementwise <add> ins(%[[IMPORT0]], %[[IMPORT1]] : +// CHECK: %[[ADD2:.+]] = linalg.elementwise <add> ins(%[[ADD1]], %[[IMPORT2]] : // CHECK: flow.return %[[ADD1]], %[[ADD2]] // CHECK: return %[[DISPATCH]]#0, %[[DISPATCH]]#1 @@ -98,11 +98,11 @@ %arg2 : tensor<10xf32>, %arg3 : !hal.fence) -> (!hal.buffer_view, !hal.buffer_view) { %0 = tensor.empty() : tensor<10xf32> - %1 = linalg.add ins(%arg0, %arg1 : tensor<10xf32>, tensor<10xf32>) + %1 = linalg.elementwise <add> ins(%arg0, %arg1 : tensor<10xf32>, tensor<10xf32>) outs(%0 : tensor<10xf32>) -> tensor<10xf32> %2 = hal.tensor.barrier join(%1 : tensor<10xf32>) => %arg3 : !hal.fence %3 = hal.tensor.export %2 : tensor<10xf32> -> !hal.buffer_view - %4 = linalg.add ins(%arg2, %1 : tensor<10xf32>, tensor<10xf32>) + %4 = linalg.elementwise <add> ins(%arg2, %1 : tensor<10xf32>, tensor<10xf32>) outs(%0 : tensor<10xf32>) -> tensor<10xf32> %5 = hal.tensor.barrier join(%4 : tensor<10xf32>) => %arg3 : !hal.fence %6 = hal.tensor.export %5 : tensor<10xf32> -> !hal.buffer_view @@ -115,8 +115,8 @@ // CHECK-SAME: %[[ARG3:[a-zA-Z0-9]+]]: !hal.fence // CHECK: %[[DISPATCH:.+]]:2 = flow.dispatch.region // CHECK: %[[EMPTY:.+]] = tensor.empty() -// CHECK: %[[ADD1:.+]] = linalg.add ins(%[[ARG0]], %[[ARG1]] : -// CHECK: %[[ADD2:.+]] = linalg.add ins(%[[ARG2]], %[[ADD1]] : +// CHECK: %[[ADD1:.+]] = linalg.elementwise <add> ins(%[[ARG0]], %[[ARG1]] : +// CHECK: %[[ADD2:.+]] = linalg.elementwise <add> ins(%[[ARG2]], %[[ADD1]] : // CHECK: flow.return %[[ADD1]], %[[ADD2]] // CHECK-DAG: %[[BARRIER0:.+]] = hal.tensor.barrier join(%[[DISPATCH]]#0 : // CHECK-DAG: %[[EXPORT0:.+]] = hal.tensor.export %[[BARRIER0]] @@ -132,7 +132,7 @@ %arg2 : tensor<10xf32>, %arg3 : !hal.fence) -> (!hal.buffer_view, !hal.buffer_view) { %0 = tensor.empty() : tensor<10xf32> - %1 = linalg.add ins(%arg0, %arg1 : tensor<10xf32>, tensor<10xf32>) + %1 = linalg.elementwise <add> ins(%arg0, %arg1 : tensor<10xf32>, tensor<10xf32>) outs(%0 : tensor<10xf32>) -> tensor<10xf32> %2 = hal.tensor.barrier join(%1 : tensor<10xf32>) => %arg3 : !hal.fence %3 = hal.tensor.export %2 : tensor<10xf32> -> !hal.buffer_view
diff --git a/runtime/src/iree/builtins/device/device_generic.c b/runtime/src/iree/builtins/device/device_generic.c index 73ee421..1427460 100644 --- a/runtime/src/iree/builtins/device/device_generic.c +++ b/runtime/src/iree/builtins/device/device_generic.c
@@ -173,4 +173,39 @@ return x * y + z; } +// LLVM codegen lowers oversized llvm.memset/memcpy/memmove intrinsics to +// libcalls, and embedded executables link with -nostdlib. +// +// Pointer-width unsigned integer for the string routines; bitcode is only ever +// retargeted between architectures of the same pointer width. +typedef __SIZE_TYPE__ iree_device_size_t; + +IREE_DEVICE_EXPORT void* memset(void* dest, int value, + iree_device_size_t count) { + uint8_t* d = (uint8_t*)dest; + for (iree_device_size_t i = 0; i < count; ++i) d[i] = (uint8_t)value; + return dest; +} + +IREE_DEVICE_EXPORT void* memcpy(void* IREE_DEVICE_RESTRICT dest, + const void* IREE_DEVICE_RESTRICT src, + iree_device_size_t count) { + uint8_t* d = (uint8_t*)dest; + const uint8_t* s = (const uint8_t*)src; + for (iree_device_size_t i = 0; i < count; ++i) d[i] = s[i]; + return dest; +} + +IREE_DEVICE_EXPORT void* memmove(void* dest, const void* src, + iree_device_size_t count) { + uint8_t* d = (uint8_t*)dest; + const uint8_t* s = (const uint8_t*)src; + if (d < s) { + for (iree_device_size_t i = 0; i < count; ++i) d[i] = s[i]; + } else { + for (iree_device_size_t i = count; i > 0; --i) d[i - 1] = s[i - 1]; + } + return dest; +} + #endif // IREE_DEVICE_STANDALONE
diff --git a/samples/transform_dialect/example_module.mlir b/samples/transform_dialect/example_module.mlir index a79ec1e..13d20e1 100644 --- a/samples/transform_dialect/example_module.mlir +++ b/samples/transform_dialect/example_module.mlir
@@ -8,7 +8,7 @@ // // module { // func.func @example_module(%A : !A_size, %B : !B_size, %C : !C_size) -> !O_size { -// %0 = linalg.add ins(%A, %A : !A_size, !A_size) +// %0 = linalg.elementwise <add> ins(%A, %A : !A_size, !A_size) // outs(%A : !A_size) -> !A_size // %1 = linalg.matmul ins(%0, %B : !A_size, !B_size) // outs(%C : !C_size) -> !C_size @@ -54,7 +54,7 @@ } { -// The linalg.add (expressed as a linalg.generic). +// The linalg.elementwise <add> (expressed as a linalg.generic). hal.executable private @example_module_dispatch_0 { hal.executable.variant public @vulkan_spirv_fb target(<"vulkan-spirv", "vulkan-spirv-fb", {iree_codegen.target_info = #target}>) { hal.executable.export public @example_module_dispatch_0_generic_80_f32 ordinal(0) layout(#pipeline_layout_0) count(%arg0: !hal.device) -> (index, index, index) {
diff --git a/tests/e2e/linalg/small_float_arith.mlir b/tests/e2e/linalg/small_float_arith.mlir index d81cb15..483c3bb 100644 --- a/tests/e2e/linalg/small_float_arith.mlir +++ b/tests/e2e/linalg/small_float_arith.mlir
@@ -1,7 +1,7 @@ func.func @add_f8E5M2FNUZ() { %input = util.unfoldable_constant dense<[0.0, 1.0, 2.0, 4.0]> : tensor<4xf8E5M2FNUZ> %init = tensor.empty() : tensor<4xf8E5M2FNUZ> - %add = linalg.add + %add = linalg.elementwise <add> ins(%input, %input : tensor<4xf8E5M2FNUZ>, tensor<4xf8E5M2FNUZ>) outs(%init : tensor<4xf8E5M2FNUZ>) -> tensor<4xf8E5M2FNUZ> check.expect_almost_eq_const(%add, dense<[0.0, 2.0, 4.0, 8.0]> : tensor<4xf8E5M2FNUZ>) : tensor<4xf8E5M2FNUZ> @@ -12,7 +12,7 @@ %lhs = util.unfoldable_constant dense<[4.0, 8.0, 16.0, 32.0]> : tensor<4xf8E5M2FNUZ> %rhs = util.unfoldable_constant dense<[1.0, 2.0, 4.0, 8.0]> : tensor<4xf8E5M2FNUZ> %init = tensor.empty() : tensor<4xf8E5M2FNUZ> - %sub = linalg.sub + %sub = linalg.elementwise <sub> ins(%lhs, %rhs : tensor<4xf8E5M2FNUZ>, tensor<4xf8E5M2FNUZ>) outs(%init : tensor<4xf8E5M2FNUZ>) -> tensor<4xf8E5M2FNUZ> check.expect_almost_eq_const(%sub, dense<[3.0, 6.0, 12.0, 24.0]> : tensor<4xf8E5M2FNUZ>) : tensor<4xf8E5M2FNUZ> @@ -23,7 +23,7 @@ %lhs = util.unfoldable_constant dense<[1.0, 2.0, 4.0, 8.0]> : tensor<4xf8E5M2FNUZ> %rhs = util.unfoldable_constant dense<[2.0, 2.0, 2.0, 2.0]> : tensor<4xf8E5M2FNUZ> %init = tensor.empty() : tensor<4xf8E5M2FNUZ> - %mul = linalg.mul + %mul = linalg.elementwise <mul> ins(%lhs, %rhs : tensor<4xf8E5M2FNUZ>, tensor<4xf8E5M2FNUZ>) outs(%init : tensor<4xf8E5M2FNUZ>) -> tensor<4xf8E5M2FNUZ> check.expect_almost_eq_const(%mul, dense<[2.0, 4.0, 8.0, 16.0]> : tensor<4xf8E5M2FNUZ>) : tensor<4xf8E5M2FNUZ> @@ -33,7 +33,7 @@ func.func @negf_f8E5M2FNUZ() { %input = util.unfoldable_constant dense<[1.0, -2.0, 4.0, -8.0]> : tensor<4xf8E5M2FNUZ> %init = tensor.empty() : tensor<4xf8E5M2FNUZ> - %neg = linalg.negf + %neg = linalg.elementwise <negf> ins(%input : tensor<4xf8E5M2FNUZ>) outs(%init : tensor<4xf8E5M2FNUZ>) -> tensor<4xf8E5M2FNUZ> check.expect_almost_eq_const(%neg, dense<[-1.0, 2.0, -4.0, 8.0]> : tensor<4xf8E5M2FNUZ>) : tensor<4xf8E5M2FNUZ> @@ -43,7 +43,7 @@ func.func @add_f8E4M3FNUZ() { %input = util.unfoldable_constant dense<[0.0, 1.0, 2.0, 4.0]> : tensor<4xf8E4M3FNUZ> %init = tensor.empty() : tensor<4xf8E4M3FNUZ> - %add = linalg.add + %add = linalg.elementwise <add> ins(%input, %input : tensor<4xf8E4M3FNUZ>, tensor<4xf8E4M3FNUZ>) outs(%init : tensor<4xf8E4M3FNUZ>) -> tensor<4xf8E4M3FNUZ> check.expect_almost_eq_const(%add, dense<[0.0, 2.0, 4.0, 8.0]> : tensor<4xf8E4M3FNUZ>) : tensor<4xf8E4M3FNUZ> @@ -54,7 +54,7 @@ %lhs = util.unfoldable_constant dense<[1.0, 2.0, 3.0, 4.0]> : tensor<4xf8E4M3FNUZ> %rhs = util.unfoldable_constant dense<[2.0, 2.0, 2.0, 2.0]> : tensor<4xf8E4M3FNUZ> %init = tensor.empty() : tensor<4xf8E4M3FNUZ> - %mul = linalg.mul + %mul = linalg.elementwise <mul> ins(%lhs, %rhs : tensor<4xf8E4M3FNUZ>, tensor<4xf8E4M3FNUZ>) outs(%init : tensor<4xf8E4M3FNUZ>) -> tensor<4xf8E4M3FNUZ> check.expect_almost_eq_const(%mul, dense<[2.0, 4.0, 6.0, 8.0]> : tensor<4xf8E4M3FNUZ>) : tensor<4xf8E4M3FNUZ> @@ -64,7 +64,7 @@ func.func @add_f8E5M2() { %input = util.unfoldable_constant dense<[0.0, 1.0, 2.0, 4.0]> : tensor<4xf8E5M2> %init = tensor.empty() : tensor<4xf8E5M2> - %add = linalg.add + %add = linalg.elementwise <add> ins(%input, %input : tensor<4xf8E5M2>, tensor<4xf8E5M2>) outs(%init : tensor<4xf8E5M2>) -> tensor<4xf8E5M2> check.expect_almost_eq_const(%add, dense<[0.0, 2.0, 4.0, 8.0]> : tensor<4xf8E5M2>) : tensor<4xf8E5M2> @@ -74,7 +74,7 @@ func.func @add_f8E4M3FN() { %input = util.unfoldable_constant dense<[0.0, 1.0, 2.0, 4.0]> : tensor<4xf8E4M3FN> %init = tensor.empty() : tensor<4xf8E4M3FN> - %add = linalg.add + %add = linalg.elementwise <add> ins(%input, %input : tensor<4xf8E4M3FN>, tensor<4xf8E4M3FN>) outs(%init : tensor<4xf8E4M3FN>) -> tensor<4xf8E4M3FN> check.expect_almost_eq_const(%add, dense<[0.0, 2.0, 4.0, 8.0]> : tensor<4xf8E4M3FN>) : tensor<4xf8E4M3FN> @@ -99,7 +99,7 @@ %lhs = util.unfoldable_constant dense<[0.0, 4.0, 32768.0, 49152.0, 0x80, -32768.0, 0.125, -4.0]> : tensor<8xf8E5M2FNUZ> %rhs = util.unfoldable_constant dense<[0.0, 4.0, 32768.0, 1.0, 1.0, -32768.0, 0.125, -4.0]> : tensor<8xf8E5M2FNUZ> %init = tensor.empty() : tensor<8xf8E5M2FNUZ> - %add = linalg.add + %add = linalg.elementwise <add> ins(%lhs, %rhs : tensor<8xf8E5M2FNUZ>, tensor<8xf8E5M2FNUZ>) outs(%init : tensor<8xf8E5M2FNUZ>) -> tensor<8xf8E5M2FNUZ> check.expect_almost_eq_const(%add, dense<[0.0, 8.0, 0x80, 49152.0, 0x80, 0x80, 0.25, -8.0]> : tensor<8xf8E5M2FNUZ>) : tensor<8xf8E5M2FNUZ> @@ -120,7 +120,7 @@ %lhs = util.unfoldable_constant dense<[0.0, 4.0, 192.0, 224.0, 0x80, -192.0, 0.125, -4.0]> : tensor<8xf8E4M3FNUZ> %rhs = util.unfoldable_constant dense<[0.0, 4.0, 192.0, 1.0, 1.0, -192.0, 0.125, -4.0]> : tensor<8xf8E4M3FNUZ> %init = tensor.empty() : tensor<8xf8E4M3FNUZ> - %add = linalg.add + %add = linalg.elementwise <add> ins(%lhs, %rhs : tensor<8xf8E4M3FNUZ>, tensor<8xf8E4M3FNUZ>) outs(%init : tensor<8xf8E4M3FNUZ>) -> tensor<8xf8E4M3FNUZ> check.expect_almost_eq_const(%add, dense<[0.0, 8.0, 0x80, 224.0, 0x80, 0x80, 0.25, -8.0]> : tensor<8xf8E4M3FNUZ>) : tensor<8xf8E4M3FNUZ> @@ -141,7 +141,7 @@ %lhs = util.unfoldable_constant dense<[0.0, -0.0, 32768.0, 0x7C, 0x7F, -32768.0, 0xFC, 0x7C]> : tensor<8xf8E5M2> %rhs = util.unfoldable_constant dense<[0.0, 0.0, 32768.0, 1.0, 1.0, -32768.0, -1.0, 0xFC]> : tensor<8xf8E5M2> %init = tensor.empty() : tensor<8xf8E5M2> - %add = linalg.add + %add = linalg.elementwise <add> ins(%lhs, %rhs : tensor<8xf8E5M2>, tensor<8xf8E5M2>) outs(%init : tensor<8xf8E5M2>) -> tensor<8xf8E5M2> check.expect_almost_eq_const(%add, dense<[0.0, 0.0, 0x7C, 0x7C, 0x7F, 0xFC, 0xFC, 0x7F]> : tensor<8xf8E5M2>) : tensor<8xf8E5M2> @@ -162,7 +162,7 @@ %lhs = util.unfoldable_constant dense<[0.0, -0.0, 384.0, 416.0, 0x7F, -384.0, 0.125, -4.0]> : tensor<8xf8E4M3FN> %rhs = util.unfoldable_constant dense<[0.0, 0.0, 384.0, 1.0, 1.0, -384.0, 0.125, -4.0]> : tensor<8xf8E4M3FN> %init = tensor.empty() : tensor<8xf8E4M3FN> - %add = linalg.add + %add = linalg.elementwise <add> ins(%lhs, %rhs : tensor<8xf8E4M3FN>, tensor<8xf8E4M3FN>) outs(%init : tensor<8xf8E4M3FN>) -> tensor<8xf8E4M3FN> check.expect_almost_eq_const(%add, dense<[0.0, 0.0, 0x7F, 416.0, 0x7F, 0x7F, 0.25, -8.0]> : tensor<8xf8E4M3FN>) : tensor<8xf8E4M3FN> @@ -185,7 +185,7 @@ func.func @negzero_negf_f8E5M2() { %input = util.unfoldable_constant dense<[0.0, 0.0, 0.0, 0.0]> : tensor<4xf8E5M2> %init = tensor.empty() : tensor<4xf8E5M2> - %neg = linalg.negf ins(%input : tensor<4xf8E5M2>) outs(%init : tensor<4xf8E5M2>) -> tensor<4xf8E5M2> + %neg = linalg.elementwise <negf> ins(%input : tensor<4xf8E5M2>) outs(%init : tensor<4xf8E5M2>) -> tensor<4xf8E5M2> %result_i8 = arith.bitcast %neg : tensor<4xf8E5M2> to tensor<4xi8> // -128 = 0x80 = -0.0 in f8E5M2 check.expect_eq_const(%result_i8, dense<[-128, -128, -128, -128]> : tensor<4xi8>) : tensor<4xi8> @@ -197,7 +197,7 @@ func.func @negzero_negf_f8E4M3FN() { %input = util.unfoldable_constant dense<[0.0, 0.0, 0.0, 0.0]> : tensor<4xf8E4M3FN> %init = tensor.empty() : tensor<4xf8E4M3FN> - %neg = linalg.negf ins(%input : tensor<4xf8E4M3FN>) outs(%init : tensor<4xf8E4M3FN>) -> tensor<4xf8E4M3FN> + %neg = linalg.elementwise <negf> ins(%input : tensor<4xf8E4M3FN>) outs(%init : tensor<4xf8E4M3FN>) -> tensor<4xf8E4M3FN> %result_i8 = arith.bitcast %neg : tensor<4xf8E4M3FN> to tensor<4xi8> // -128 = 0x80 = -0.0 in f8E4M3FN check.expect_eq_const(%result_i8, dense<[-128, -128, -128, -128]> : tensor<4xi8>) : tensor<4xi8> @@ -210,7 +210,7 @@ %input = util.unfoldable_constant dense<[-0.0, -0.0, -0.0, -0.0]> : tensor<4xf8E5M2> %one = util.unfoldable_constant dense<[1.0, 1.0, 1.0, 1.0]> : tensor<4xf8E5M2> %init = tensor.empty() : tensor<4xf8E5M2> - %mul = linalg.mul ins(%input, %one : tensor<4xf8E5M2>, tensor<4xf8E5M2>) outs(%init : tensor<4xf8E5M2>) -> tensor<4xf8E5M2> + %mul = linalg.elementwise <mul> ins(%input, %one : tensor<4xf8E5M2>, tensor<4xf8E5M2>) outs(%init : tensor<4xf8E5M2>) -> tensor<4xf8E5M2> %result_i8 = arith.bitcast %mul : tensor<4xf8E5M2> to tensor<4xi8> // -128 = 0x80 = -0.0 in f8E5M2 check.expect_eq_const(%result_i8, dense<[-128, -128, -128, -128]> : tensor<4xi8>) : tensor<4xi8> @@ -222,7 +222,7 @@ %input = util.unfoldable_constant dense<[-0.0, -0.0, -0.0, -0.0]> : tensor<4xf8E4M3FN> %one = util.unfoldable_constant dense<[1.0, 1.0, 1.0, 1.0]> : tensor<4xf8E4M3FN> %init = tensor.empty() : tensor<4xf8E4M3FN> - %mul = linalg.mul ins(%input, %one : tensor<4xf8E4M3FN>, tensor<4xf8E4M3FN>) outs(%init : tensor<4xf8E4M3FN>) -> tensor<4xf8E4M3FN> + %mul = linalg.elementwise <mul> ins(%input, %one : tensor<4xf8E4M3FN>, tensor<4xf8E4M3FN>) outs(%init : tensor<4xf8E4M3FN>) -> tensor<4xf8E4M3FN> %result_i8 = arith.bitcast %mul : tensor<4xf8E4M3FN> to tensor<4xi8> // -128 = 0x80 = -0.0 in f8E4M3FN check.expect_eq_const(%result_i8, dense<[-128, -128, -128, -128]> : tensor<4xi8>) : tensor<4xi8> @@ -369,7 +369,7 @@ %input_i8 = util.unfoldable_constant dense<[0x20, 0x64, 0xCA, 0x1E]> : tensor<4xi8> %input_f4 = flow.tensor.bitcast %input_i8 : tensor<4xi8> -> tensor<8xf4E2M1FN> %init = tensor.empty() : tensor<8xf4E2M1FN> - %add = linalg.add + %add = linalg.elementwise <add> ins(%input_f4, %input_f4 : tensor<8xf4E2M1FN>, tensor<8xf4E2M1FN>) outs(%init : tensor<8xf4E2M1FN>) -> tensor<8xf4E2M1FN> // Expected after doubling:
diff --git a/tests/e2e/parameters/encode_parameters.mlir b/tests/e2e/parameters/encode_parameters.mlir index 9613be1..417b155 100644 --- a/tests/e2e/parameters/encode_parameters.mlir +++ b/tests/e2e/parameters/encode_parameters.mlir
@@ -54,7 +54,7 @@ %c42 = arith.constant 42 : i32 %init = tensor.empty() : tensor<256xi32> %c42_tensor = linalg.fill ins(%c42 : i32) outs(%init : tensor<256xi32>) -> tensor<256xi32> - %added = linalg.add ins(%raw, %c42_tensor : tensor<256xi32>, tensor<256xi32>) outs(%init : tensor<256xi32>) -> tensor<256xi32> + %added = linalg.elementwise <add> ins(%raw, %c42_tensor : tensor<256xi32>, tensor<256xi32>) outs(%init : tensor<256xi32>) -> tensor<256xi32> util.global.store %added, @transformed : tensor<256xi32> util.return }
diff --git a/tests/e2e/parameters/export_parameters.mlir b/tests/e2e/parameters/export_parameters.mlir index ef8321e..ab770e3 100644 --- a/tests/e2e/parameters/export_parameters.mlir +++ b/tests/e2e/parameters/export_parameters.mlir
@@ -27,8 +27,8 @@ %empty = tensor.empty() : tensor<1x2xf32> %fill = linalg.fill ins(%cst : f32) outs(%empty : tensor<1x2xf32>) -> tensor<1x2xf32> %8 = linalg.matmul ins(%arg0, %6 : tensor<1x2xf32>, tensor<2x2xf32>) outs(%fill : tensor<1x2xf32>) -> tensor<1x2xf32> - %10 = linalg.add ins(%8, %5 : tensor<1x2xf32>, tensor<1x2xf32>) outs(%empty : tensor<1x2xf32>) -> tensor<1x2xf32> + %10 = linalg.elementwise <add> ins(%8, %5 : tensor<1x2xf32>, tensor<1x2xf32>) outs(%empty : tensor<1x2xf32>) -> tensor<1x2xf32> %12 = linalg.matmul ins(%10, %4 : tensor<1x2xf32>, tensor<2x2xf32>) outs(%fill : tensor<1x2xf32>) -> tensor<1x2xf32> - %14 = linalg.add ins(%12, %3 : tensor<1x2xf32>, tensor<1x2xf32>) outs(%empty : tensor<1x2xf32>) -> tensor<1x2xf32> + %14 = linalg.elementwise <add> ins(%12, %3 : tensor<1x2xf32>, tensor<1x2xf32>) outs(%empty : tensor<1x2xf32>) -> tensor<1x2xf32> return %14 : tensor<1x2xf32> }
diff --git a/tests/e2e/parameters/generate_splat_archive.mlir b/tests/e2e/parameters/generate_splat_archive.mlir index fb888c7..ba02399 100644 --- a/tests/e2e/parameters/generate_splat_archive.mlir +++ b/tests/e2e/parameters/generate_splat_archive.mlir
@@ -27,8 +27,8 @@ %empty = tensor.empty() : tensor<1x2xi32> %fill = linalg.fill ins(%cst : i32) outs(%empty : tensor<1x2xi32>) -> tensor<1x2xi32> %8 = linalg.matmul ins(%arg0, %6 : tensor<1x2xi32>, tensor<2x2xi32>) outs(%fill : tensor<1x2xi32>) -> tensor<1x2xi32> - %10 = linalg.add ins(%8, %5 : tensor<1x2xi32>, tensor<1x2xi32>) outs(%empty : tensor<1x2xi32>) -> tensor<1x2xi32> + %10 = linalg.elementwise <add> ins(%8, %5 : tensor<1x2xi32>, tensor<1x2xi32>) outs(%empty : tensor<1x2xi32>) -> tensor<1x2xi32> %12 = linalg.matmul ins(%10, %4 : tensor<1x2xi32>, tensor<2x2xi32>) outs(%fill : tensor<1x2xi32>) -> tensor<1x2xi32> - %14 = linalg.add ins(%12, %3 : tensor<1x2xi32>, tensor<1x2xi32>) outs(%empty : tensor<1x2xi32>) -> tensor<1x2xi32> + %14 = linalg.elementwise <add> ins(%12, %3 : tensor<1x2xi32>, tensor<1x2xi32>) outs(%empty : tensor<1x2xi32>) -> tensor<1x2xi32> return %14 : tensor<1x2xi32> }
diff --git a/third_party/llvm-project b/third_party/llvm-project index 6cce5bc..c53000a 160000 --- a/third_party/llvm-project +++ b/third_party/llvm-project
@@ -1 +1 @@ -Subproject commit 6cce5bca3b901ec32e250fb43e88eb0fe511af93 +Subproject commit c53000a7282bedb125ffaab8cb072df85596d095