Forking dynamic behavior from flow.tensor.constant. (#17034) Needed this sooner than I expected due to requiring the ConstantLike trait on the constant op in order for constant materialization to work. Now the special behavior is moved to the flow.tensor.dynamic_constant op. Follow-up to #17024.
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowBase.td b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowBase.td index 1f2cc20..78b9036 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowBase.td +++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowBase.td
@@ -11,6 +11,7 @@ include "iree/compiler/Dialect/Util/IR/UtilBase.td" include "iree/compiler/Dialect/Util/IR/UtilTypes.td" include "mlir/IR/AttrTypeBase.td" +include "mlir/IR/BuiltinAttributeInterfaces.td" //===----------------------------------------------------------------------===// // IREE execution flow dialect @@ -233,4 +234,33 @@ let cppNamespace = "mlir::iree_compiler::IREE::Flow"; } +//===----------------------------------------------------------------------===// +// Parameter storage attributes +//===----------------------------------------------------------------------===// + +def FLOW_NamedParameterAttr : + AttrDef<Flow_Dialect, "NamedParameter", [ + TypedAttrInterface, + ]> { + let mnemonic = "parameter.named"; + let summary = [{named parameter referenced an optional scope and key}]; + let description = [{ + Species an externally-defined parameter that can be referenced by an + optional scope defining a set of parameters and a key uniquely identifying + the parameter within its scope. + }]; + let parameters = (ins + AttributeSelfTypeParameter<"">:$type, + OptionalParameter<"StringAttr">:$scope, + AttrParameter<"StringAttr", "">:$key, + OptionalParameter<"DictionaryAttr">:$config + ); + let assemblyFormat = [{ + `<` + custom<ParameterReference>($scope, $key) + (`,` $config^)? + `>` + }]; +} + #endif // IREE_DIALECT_FLOW_BASE
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp index 5b0d2e2..577ad48 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp +++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp
@@ -74,8 +74,12 @@ Operation *FlowDialect::materializeConstant(OpBuilder &builder, Attribute value, Type type, Location loc) { - if (arith::ConstantOp::isBuildableWith(value, type)) + if (arith::ConstantOp::isBuildableWith(value, type)) { return builder.create<arith::ConstantOp>(loc, type, cast<TypedAttr>(value)); + } else if (IREE::Flow::TensorConstantOp::isBuildableWith(value, type)) { + return builder.create<IREE::Flow::TensorConstantOp>(loc, type, + cast<TypedAttr>(value)); + } return nullptr; }
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp index d37b28c..a2079c3 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp +++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp
@@ -779,7 +779,13 @@ // flow.tensor.constant //===----------------------------------------------------------------------===// -OpFoldResult TensorConstantOp::fold(FoldAdaptor operands) { +OpFoldResult TensorConstantOp::fold(FoldAdaptor operands) { return getValue(); } + +//===----------------------------------------------------------------------===// +// flow.tensor.dynamic_constant +//===----------------------------------------------------------------------===// + +OpFoldResult TensorDynamicConstantOp::fold(FoldAdaptor operands) { auto dynamicType = getType(); if (dynamicType.getNumDynamicDims() == 0) { return getValue(); @@ -787,6 +793,43 @@ return {}; } +namespace { + +struct ExpandDynamicShapeConstant + : public OpRewritePattern<TensorDynamicConstantOp> { + using OpRewritePattern<TensorDynamicConstantOp>::OpRewritePattern; + LogicalResult matchAndRewrite(TensorDynamicConstantOp op, + PatternRewriter &rewriter) const override { + auto constantOp = rewriter.create<IREE::Flow::TensorConstantOp>( + op.getLoc(), op.getValue()); + auto dynamicType = op.getType(); + auto staticType = cast<ShapedType>(op.getValue().getType()); + SmallVector<Value> dynamicDims; + for (int64_t i = 0; i < dynamicType.getRank(); ++i) { + if (dynamicType.isDynamicDim(i)) { + auto dimValue = rewriter + .create<arith::ConstantIndexOp>( + op.getLoc(), staticType.getDimSize(i)) + .getResult(); + dynamicDims.push_back(rewriter + .create<IREE::Util::OptimizationBarrierOp>( + op.getLoc(), dimValue) + .getResult(0)); + } + } + rewriter.replaceOpWithNewOp<IREE::Flow::TensorReshapeOp>( + op, dynamicType, constantOp.getResult(), dynamicDims); + return success(); + } +}; + +} // namespace + +void TensorDynamicConstantOp::getCanonicalizationPatterns( + RewritePatternSet &results, MLIRContext *context) { + results.insert<ExpandDynamicShapeConstant>(context); +} + //===----------------------------------------------------------------------===// // flow.tensor.tie_shape //===----------------------------------------------------------------------===// @@ -973,26 +1016,6 @@ return success(); } - // Special handling of flow.tensor.constant which may be acting as a - // dynamically shaped value that we want to remove the tensor.dim of but - // still treat the shape as dynamic. We do this by inserting an optimization - // barrier between the constant and the consumers. Note that this use case - // is very specific and generally only applicable to tests/benchmarks. - if (auto constantOp = dyn_cast_if_present<IREE::Flow::TensorConstantOp>( - op.getShapedValue().getDefiningOp())) { - auto valueType = dyn_cast<ShapedType>(constantOp.getValue().getType()); - if (valueType && valueType != constantOp.getType()) { - // Constant op is acting as a cast. If the dimension being queried was - // static it would have been resolved above so we know it's dynamic - // here. - Value staticValue = rewriter.create<arith::ConstantIndexOp>( - op.getLoc(), valueType.getDimSize(idx)); - rewriter.replaceOpWithNewOp<IREE::Util::OptimizationBarrierOp>( - op, staticValue); - return success(); - } - } - return rewriter.notifyMatchFailure(op, "no dynamic dims found/usable"); } };
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp index 0cb7297..c90bbfd 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp +++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
@@ -1597,35 +1597,18 @@ // flow.tensor.constant //===----------------------------------------------------------------------===// -ParseResult TensorConstantOp::parse(OpAsmParser &parser, - OperationState &result) { - if (parser.parseOptionalAttrDict(result.attributes)) - return failure(); - TypedAttr valueAttr; - if (failed(parser.parseAttribute(valueAttr))) - return failure(); - result.addAttribute("value", valueAttr); - if (succeeded(parser.parseOptionalArrow())) { - Type resultType; - if (failed(parser.parseType(resultType))) - return failure(); - result.addTypes(resultType); - } else { - result.addTypes(valueAttr.getType()); - } - return success(); +// static +bool TensorConstantOp::isBuildableWith(Attribute value, Type type) { + return isa<RankedTensorType>(type); } -void TensorConstantOp::print(OpAsmPrinter &p) { - p << " "; - p.printOptionalAttrDict((*this)->getAttrs(), {"value"}); - p.printAttribute(getValue()); - auto attrType = getValue().getType(); - auto resultType = getType(); - if (attrType != resultType) { - p << " -> "; - p.printType(resultType); - } +//===----------------------------------------------------------------------===// +// flow.tensor.dynamic_constant +//===----------------------------------------------------------------------===// + +// static +bool TensorDynamicConstantOp::isBuildableWith(Attribute value, Type type) { + return TensorConstantOp::isBuildableWith(value, type); } //===----------------------------------------------------------------------===//
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.td b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.td index c797c03..d488b02 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.td +++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.td
@@ -1040,26 +1040,62 @@ let opDocGroup = OpGroupTensorOps in { -def FLOW_TensorConstantOp : FLOW_PureOp<"tensor.constant"> { +def FLOW_TensorConstantOp : FLOW_PureOp<"tensor.constant", [ + ConstantLike, + AllTypesMatch<["value", "result"]>, +]> { let summary = [{tensor constant that can have dynamic dimensions}]; let description = [{ - Allows specifying a tensor constant of IREE-specific types/attributes or - where the return value can erase shape information. + Allows specifying a tensor constant of IREE-specific types/attributes. ```mlir %cst = flow.tensor.constant #something_tensor_like : tensor<2x2xf32> %res = math.absf %cst : tensor<2x2xf32> ``` + }]; + let arguments = (ins TypedAttrInterface:$value); + let results = (outs AnyTensor:$result); + + let assemblyFormat = [{ + attr-dict $value + }]; + + let extraClassDeclaration = [{ + // Returns true if the constant op can be built with the given attribute. + static bool isBuildableWith(Attribute value, Type type); + }]; + + let hasFolder = 1; +} + +def FLOW_TensorDynamicConstantOp : FLOW_Op<"tensor.dynamic_constant"> { + let summary = [{tensor constant that can have dynamic dimensions}]; + let description = [{ + Allows specifying a tensor constant of IREE-specific types/attributes with + a dynamic shape that approximates a value as passed from the user. This + disables many optimizations and should only be used when testing or + benchmarking and wanting to ensure that dynamic dimension behavior is + preserved. ```mlir - %cst = flow.tensor.constant dense<4.0> : tensor<2x2xf32> -> tensor<?x2xf32> + %cst = flow.tensor.dynamic_constant #something_tensor_like : tensor<2x2xf32> -> tensor<?x2xf32> %res = math.absf %cst : tensor<?x2xf32> ``` }]; let arguments = (ins TypedAttrInterface:$value); let results = (outs AnyTensor:$result); + let assemblyFormat = [{ + attr-dict $value `->` type($result) + }]; + + let extraClassDeclaration = [{ + // Returns true if the constant op can be built with the given attribute. + static bool isBuildableWith(Attribute value, Type type); + }]; + let hasFolder = 1; + let hasCanonicalizer = 1; } def FLOW_TensorTieShapeOp : FLOW_PureOp<"tensor.tie_shape", [
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowTypes.cpp b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowTypes.cpp index 4f93ea7..c772d63 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowTypes.cpp +++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowTypes.cpp
@@ -276,4 +276,38 @@ *collectiveElemType); } +//===----------------------------------------------------------------------===// +// custom<ParameterReference>($scope, $key) +//===----------------------------------------------------------------------===// + +ParseResult parseParameterReference(AsmParser &parser, StringAttr &scopeAttr, + StringAttr &keyAttr) { + auto builder = parser.getBuilder(); + StringAttr firstAttr; + if (failed(parser.parseCustomAttributeWithFallback(firstAttr, + builder.getNoneType()))) { + return failure(); + } + if (failed(parser.parseOptionalColon())) { + keyAttr = firstAttr; + return success(); + } + scopeAttr = firstAttr; + if (failed(parser.parseColon()) || + failed(parser.parseCustomAttributeWithFallback(keyAttr, + builder.getNoneType()))) { + return failure(); + } + return success(); +} + +void printParameterReference(AsmPrinter &p, StringAttr scopeAttr, + StringAttr keyAttr) { + if (scopeAttr) { + p << "\"" << scopeAttr.getValue() << "\""; + p << "::"; + } + p << "\"" << keyAttr.getValue() << "\""; +} + } // namespace mlir::iree_compiler::IREE::Flow
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowTypes.h b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowTypes.h index 3cc34ff..a1605c7 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowTypes.h +++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowTypes.h
@@ -183,6 +183,20 @@ std::optional<IREE::Flow::CollectiveElementType> convertToFlowCollectiveElementType(Type type); +//===----------------------------------------------------------------------===// +// custom<ParameterReference>($scope, $key) +//===----------------------------------------------------------------------===// + +ParseResult parseParameterReference(AsmParser &parser, StringAttr &scopeAttr, + StringAttr &keyAttr); +void printParameterReference(AsmPrinter &p, StringAttr scopeAttr, + StringAttr keyAttr); +static inline void printParameterReference(AsmPrinter &p, Operation *op, + StringAttr scopeAttr, + StringAttr keyAttr) { + printParameterReference(p, scopeAttr, keyAttr); +} + } // namespace mlir::iree_compiler::IREE::Flow #endif // IREE_COMPILER_DIALECT_FLOW_IR_FLOWTYPES_H_
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/test/tensor_folding.mlir b/compiler/src/iree/compiler/Dialect/Flow/IR/test/tensor_folding.mlir index de5a8d4..bcb1dbb 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/IR/test/tensor_folding.mlir +++ b/compiler/src/iree/compiler/Dialect/Flow/IR/test/tensor_folding.mlir
@@ -4,8 +4,8 @@ util.func public @expandStaticShapeConstant() -> (tensor<2x4xi32>, index, index) { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index - // CHECK-DAG: %[[CST:.+]] = arith.constant dense<2> : tensor<2x4xi32> - %0 = flow.tensor.constant dense<2> : tensor<2x4xi32> -> tensor<2x4xi32> + // CHECK-DAG: %[[CST:.+]] = flow.tensor.constant dense<2> : tensor<2x4xi32> + %0 = flow.tensor.constant dense<2> : tensor<2x4xi32> // CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index %d0 = tensor.dim %0, %c0 : tensor<2x4xi32> // CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index @@ -16,19 +16,20 @@ // ----- -// CHECK-LABEL: @tensorDimOfDynamicConstant -util.func public @tensorDimOfDynamicConstant() -> (index, index) { +// CHECK-LABEL: @expandDynamicShapeConstant +util.func public @expandDynamicShapeConstant() -> (tensor<2x?xi32>, index, index) { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index - // CHECK-NOT: flow.tensor.constant - %0 = flow.tensor.constant dense<2> : tensor<2x4xi32> -> tensor<2x?xi32> + // CHECK-DAG: %[[CST:.+]] = flow.tensor.constant dense<2> : tensor<2x4xi32> + %0 = flow.tensor.dynamic_constant dense<2> : tensor<2x4xi32> -> tensor<2x?xi32> // CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index - %d0 = tensor.dim %0, %c0 : tensor<2x?xi32> // CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index - // CHECK-DAG: %[[C4_DYNAMIC:.+]] = util.optimization_barrier %[[C4]] + // CHECK-DAG: %[[D1:.+]] = util.optimization_barrier %[[C4]] : index + // CHECK: %[[T:.+]] = flow.tensor.reshape %[[CST]] : tensor<2x4xi32> -> tensor<2x?xi32>{%[[D1]]} + %d0 = tensor.dim %0, %c0 : tensor<2x?xi32> %d1 = tensor.dim %0, %c1 : tensor<2x?xi32> - // CHECK: util.return %[[C2]], %[[C4_DYNAMIC]] - util.return %d0, %d1 : index, index + // CHECK: util.return %[[T]], %[[C2]], %[[D1]] + util.return %0, %d0, %d1 : tensor<2x?xi32>, index, index } // -----
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir b/compiler/src/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir index 9a66945..f727f4e 100644 --- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir +++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir
@@ -794,8 +794,8 @@ %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %cst = arith.constant 0.000000e+00 : f32 - %0 = flow.tensor.constant dense<[[1.500000e+01, 1.400000e+01, 1.300000e+01], [1.200000e+01, 1.100000e+01, 1.000000e+01], [9.000000e+00, 8.000000e+00, 7.000000e+00], [6.000000e+00, 5.000000e+00, 4.000000e+00], [3.000000e+00, 2.000000e+00, 1.000000e+00]]> : tensor<5x3xf32> -> tensor<?x?xf32> - %1 = flow.tensor.constant dense<[[1.500000e+01, 1.400000e+01, 1.300000e+01, 1.200000e+01, 1.100000e+01], [1.000000e+01, 9.000000e+00, 8.000000e+00, 7.000000e+00, 6.000000e+00], [5.000000e+00, 4.000000e+00, 3.000000e+00, 2.000000e+00, 1.000000e+00]]> : tensor<3x5xf32> -> tensor<?x?xf32> + %0 = flow.tensor.dynamic_constant dense<[[1.500000e+01, 1.400000e+01, 1.300000e+01], [1.200000e+01, 1.100000e+01, 1.000000e+01], [9.000000e+00, 8.000000e+00, 7.000000e+00], [6.000000e+00, 5.000000e+00, 4.000000e+00], [3.000000e+00, 2.000000e+00, 1.000000e+00]]> : tensor<5x3xf32> -> tensor<?x?xf32> + %1 = flow.tensor.dynamic_constant dense<[[1.500000e+01, 1.400000e+01, 1.300000e+01, 1.200000e+01, 1.100000e+01], [1.000000e+01, 9.000000e+00, 8.000000e+00, 7.000000e+00, 6.000000e+00], [5.000000e+00, 4.000000e+00, 3.000000e+00, 2.000000e+00, 1.000000e+00]]> : tensor<3x5xf32> -> tensor<?x?xf32> %2 = tensor.dim %0, %c0 : tensor<?x?xf32> %3 = tensor.dim %1, %c1 : tensor<?x?xf32> %4 = tensor.empty(%2, %3) : tensor<?x?xf32>
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 b80b75f..93f1aef 100644 --- a/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/Patterns.cpp +++ b/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/Patterns.cpp
@@ -41,25 +41,51 @@ LogicalResult matchAndRewrite(IREE::Flow::TensorConstantOp constantOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { + // Capture the tensor constant strongly typed with constant lifetime. + Type constantType = IREE::Stream::ResourceType::get( + getContext(), IREE::Stream::Lifetime::Constant); + auto affinityAttr = IREE::Stream::AffinityAttr::lookup(constantOp); + auto newOp = rewriter.create<IREE::Stream::TensorConstantOp>( + constantOp.getLoc(), constantType, constantOp.getValue(), + TypeAttr::get(constantOp.getType()), ValueRange{}, affinityAttr); + + // Transfer to unknown lifetime. + Type unknownType = IREE::Stream::ResourceType::get(getContext()); + auto constantSize = rewriter.createOrFold<IREE::Stream::ResourceSizeOp>( + constantOp.getLoc(), rewriter.getIndexType(), newOp.getResult()); + rewriter.replaceOpWithNewOp<IREE::Stream::AsyncTransferOp>( + constantOp, unknownType, newOp.getResult(), constantSize, constantSize, + /*source_affinity=*/affinityAttr, + /*result_affinity=*/affinityAttr); + return success(); + } +}; + +struct ConvertTensorDynamicConstantOp + : public OpConversionPattern<IREE::Flow::TensorDynamicConstantOp> { +public: + using OpConversionPattern::OpConversionPattern; + LogicalResult + matchAndRewrite(IREE::Flow::TensorDynamicConstantOp constantOp, + OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override { auto attrType = dyn_cast<RankedTensorType>(constantOp.getValue().getType()); if (!attrType) return failure(); auto resultType = constantOp.getType(); // If the op is acting as a dynamic value then preserve that behavior by - // calculation the shape through optimization barriers. + // calculating the shape through optimization barriers. SmallVector<Value> dynamicDims; - if (!resultType.hasStaticShape()) { - for (unsigned i = 0; i < resultType.getRank(); ++i) { - if (resultType.isDynamicDim(i)) { - Value staticDim = rewriter.create<arith::ConstantIndexOp>( - constantOp.getLoc(), attrType.getDimSize(i)); - Value dynamicDim = rewriter - .create<IREE::Util::OptimizationBarrierOp>( - constantOp.getLoc(), staticDim) - .getResult(0); - dynamicDims.push_back(dynamicDim); - } + for (unsigned i = 0; i < resultType.getRank(); ++i) { + if (resultType.isDynamicDim(i)) { + Value staticDim = rewriter.create<arith::ConstantIndexOp>( + constantOp.getLoc(), attrType.getDimSize(i)); + Value dynamicDim = rewriter + .create<IREE::Util::OptimizationBarrierOp>( + constantOp.getLoc(), staticDim) + .getResult(0); + dynamicDims.push_back(dynamicDim); } } @@ -937,7 +963,7 @@ TypeConverter &typeConverter, RewritePatternSet &patterns) { patterns - .insert<ConvertTensorConstantOp, + .insert<ConvertTensorConstantOp, ConvertTensorDynamicConstantOp, ConvertTensorCastLikeOp<IREE::Flow::TensorReshapeOp>, ConvertTensorCastLikeOp<IREE::Flow::TensorBitCastOp>, ConvertTensorAllocaOp, ConvertTensorEmptyOp, ConvertTensorSplatOp,
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/test/tensor_ops.mlir b/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/test/tensor_ops.mlir index 756f319..029423b 100644 --- a/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/test/tensor_ops.mlir +++ b/compiler/src/iree/compiler/Dialect/Stream/Conversion/FlowToStream/test/tensor_ops.mlir
@@ -12,6 +12,18 @@ // ----- +// CHECK-LABEL: @tensorConstantParameter +util.func public @tensorConstantParameter() -> tensor<4x2xi32> { + // CHECK-DAG: %[[CST:.+]] = stream.tensor.constant : tensor<4x2xi32> in !stream.resource<constant> = #stream.parameter.named<"scope"::"key"> : tensor<4x2xi32> + // CHECK-DAG: %[[SIZE:.+]] = stream.resource.size %[[CST]] : !stream.resource<constant> + // CHECK-DAG: %[[TRANSFER:.+]] = stream.async.transfer %[[CST]] : !stream.resource<constant>{%[[SIZE]]} -> !stream.resource<*>{%[[SIZE]]} + %cst = flow.tensor.constant #stream.parameter.named<"scope"::"key"> : tensor<4x2xi32> + // CHECK: util.return %[[TRANSFER]], %[[SIZE]] + util.return %cst : tensor<4x2xi32> +} + +// ----- + // CHECK-LABEL: @tensorConstantDynamic util.func public @tensorConstantDynamic() -> tensor<?x?xi32> { // CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index @@ -21,25 +33,13 @@ // CHECK-DAG: %[[CST:.+]] = stream.tensor.constant : tensor<?x?xi32>{%[[D0]], %[[D1]]} in !stream.resource<constant> = dense<2> : tensor<2x4xi32> // CHECK-DAG: %[[SIZE:.+]] = stream.resource.size %[[CST]] : !stream.resource<constant> // CHECK-DAG: %[[TRANSFER:.+]] = stream.async.transfer %[[CST]] : !stream.resource<constant>{%[[SIZE]]} -> !stream.resource<*>{%[[SIZE]]} - %cst = flow.tensor.constant dense<2> : tensor<2x4xi32> -> tensor<?x?xi32> + %cst = flow.tensor.dynamic_constant dense<2> : tensor<2x4xi32> -> tensor<?x?xi32> // CHECK: util.return %[[TRANSFER]], %[[SIZE]] util.return %cst : tensor<?x?xi32> } // ----- -// CHECK-LABEL: @tensorConstantParameter -util.func public @tensorConstantParameter() -> tensor<4x2xi32> { - // CHECK-DAG: %[[CST:.+]] = stream.tensor.constant : tensor<4x2xi32> in !stream.resource<constant> = #stream.parameter.named<"scope"::"key"> : tensor<4x2xi32> - // CHECK-DAG: %[[SIZE:.+]] = stream.resource.size %[[CST]] : !stream.resource<constant> - // CHECK-DAG: %[[TRANSFER:.+]] = stream.async.transfer %[[CST]] : !stream.resource<constant>{%[[SIZE]]} -> !stream.resource<*>{%[[SIZE]]} - %cst = flow.tensor.constant #stream.parameter.named<"scope"::"key"> : tensor<4x2xi32> - // CHECK: util.return %[[TRANSFER]], %[[SIZE]] - util.return %cst : tensor<4x2xi32> -} - -// ----- - // CHECK-LABEL: @tensorReshapePassThrough // CHECK-SAME: (%[[INPUT:.+]]: !stream.resource<*>, %[[INPUT_SIZE:.+]]: index) util.func public @tensorReshapePassThrough(%input: tensor<5x24x48xf32>) -> tensor<30x2x96xf32> {
diff --git a/tests/e2e/regression/dynamic_abs.mlir b/tests/e2e/regression/dynamic_abs.mlir index 5749a86..95d0720 100644 --- a/tests/e2e/regression/dynamic_abs.mlir +++ b/tests/e2e/regression/dynamic_abs.mlir
@@ -1,5 +1,5 @@ func.func @dynamic_tensor() { - %input = flow.tensor.constant dense<[[-1.0, 2.0, -3.0], [4.0, -5.0, 6.0]]> : tensor<2x3xf32> -> tensor<?x?xf32> + %input = flow.tensor.dynamic_constant dense<[[-1.0, 2.0, -3.0], [4.0, -5.0, 6.0]]> : tensor<2x3xf32> -> tensor<?x?xf32> %res = stablehlo.abs %input : tensor<?x?xf32> %dshape = util.optimization_barrier %res : tensor<?x?xf32> %result = tensor.cast %dshape : tensor<?x?xf32> to tensor<2x3xf32>
diff --git a/tests/e2e/regression/dynamic_add.mlir b/tests/e2e/regression/dynamic_add.mlir index 318513b..a0fad0d 100644 --- a/tests/e2e/regression/dynamic_add.mlir +++ b/tests/e2e/regression/dynamic_add.mlir
@@ -1,6 +1,6 @@ func.func @main() { - %lhs = flow.tensor.constant dense<[[1.0,2.0,3.0,4.0],[-1.0,-2.0,-3.0,-4.0]]> : tensor<2x4xf32> -> tensor<?x4xf32> - %rhs = flow.tensor.constant dense<[[5.0,6.0,7.0,8.0],[-5.0,-6.0,-7.0,-8.0]]> : tensor<2x4xf32> -> tensor<?x4xf32> + %lhs = flow.tensor.dynamic_constant dense<[[1.0,2.0,3.0,4.0],[-1.0,-2.0,-3.0,-4.0]]> : tensor<2x4xf32> -> tensor<?x4xf32> + %rhs = flow.tensor.dynamic_constant dense<[[5.0,6.0,7.0,8.0],[-5.0,-6.0,-7.0,-8.0]]> : tensor<2x4xf32> -> tensor<?x4xf32> %2 = stablehlo.add %lhs, %rhs : tensor<?x4xf32> %3 = util.optimization_barrier %2 : tensor<?x4xf32> %result = tensor.cast %3 : tensor<?x4xf32> to tensor<2x4xf32>
diff --git a/tests/e2e/regression/dynamic_dot.mlir b/tests/e2e/regression/dynamic_dot.mlir index 041993e..6b7bd20 100644 --- a/tests/e2e/regression/dynamic_dot.mlir +++ b/tests/e2e/regression/dynamic_dot.mlir
@@ -1,11 +1,11 @@ func.func @dynamic_dot() { - %lhs = flow.tensor.constant dense<[ + %lhs = flow.tensor.dynamic_constant dense<[ [15.0, 14.0, 13.0], [12.0, 11.0, 10.0], [09.0, 08.0, 07.0], [06.0, 05.0, 04.0], [03.0, 02.0, 01.0]]> : tensor<5x3xf32> -> tensor<?x?xf32> - %rhs = flow.tensor.constant dense<[ + %rhs = flow.tensor.dynamic_constant dense<[ [15.0, 14.0, 13.0, 12.0, 11.0], [10.0, 09.0, 08.0, 07.0, 06.0], [05.0, 04.0, 03.0, 02.0, 01.0]]> : tensor<3x5xf32> -> tensor<?x?xf32>
diff --git a/tests/e2e/regression/dynamic_reduce_min.mlir b/tests/e2e/regression/dynamic_reduce_min.mlir index 382d268..f484cec 100644 --- a/tests/e2e/regression/dynamic_reduce_min.mlir +++ b/tests/e2e/regression/dynamic_reduce_min.mlir
@@ -1,5 +1,5 @@ func.func @reduce_min() { - %input = flow.tensor.constant + %input = flow.tensor.dynamic_constant dense<[[1.0, 2.0, 3.0, 4.0],[-1.0 ,-2.0 ,-3.0 ,-4.0]]> : tensor<2x4xf32> -> tensor<?x?xf32> %0 = stablehlo.constant dense<0x7F800000> : tensor<f32> %1 = "stablehlo.reduce"(%input, %0) ( {
diff --git a/tests/e2e/regression/dynamic_torch_index_select_high_rank.mlir b/tests/e2e/regression/dynamic_torch_index_select_high_rank.mlir index 501b242..ccfed09 100644 --- a/tests/e2e/regression/dynamic_torch_index_select_high_rank.mlir +++ b/tests/e2e/regression/dynamic_torch_index_select_high_rank.mlir
@@ -1,6 +1,6 @@ func.func @torch_index_select1() { - %lhs = flow.tensor.constant dense<[[6,7],[8,9]]> : tensor<2x2xi32> -> tensor<?x?xi32> - %rhs = flow.tensor.constant dense<[[[[0,1],[1,0]],[[0,0],[1,1]]],[[[1,1],[0,0]],[[0,1],[1,0]]]]> : tensor<2x2x2x2xi32> -> tensor<?x?x?x?xi32> + %lhs = flow.tensor.dynamic_constant dense<[[6,7],[8,9]]> : tensor<2x2xi32> -> tensor<?x?xi32> + %rhs = flow.tensor.dynamic_constant dense<[[[[0,1],[1,0]],[[0,0],[1,1]]],[[[1,1],[0,0]],[[0,1],[1,0]]]]> : tensor<2x2x2x2xi32> -> tensor<?x?x?x?xi32> %0 = "stablehlo.torch_index_select"(%lhs, %rhs) {batch_dims = 1 : i64, dim = 1 : i64} : (tensor<?x?xi32>, tensor<?x?x?x?xi32>) -> tensor<?x?x?x?xi32> %dshape = util.optimization_barrier %0 : tensor<?x?x?x?xi32> %result = tensor.cast %dshape : tensor<?x?x?x?xi32> to tensor<2x2x2x2xi32> @@ -13,8 +13,8 @@ } func.func @torch_index_select2() { - %lhs = flow.tensor.constant dense<[[6,7],[8,9]]> : tensor<2x2xi32> -> tensor<?x?xi32> - %rhs = flow.tensor.constant dense<[[[[0,1],[1,0]],[[0,0],[1,1]]],[[[1,1],[0,0]],[[0,1],[1,0]]]]> : tensor<2x2x2x2xi32> -> tensor<?x?x?x?xi32> + %lhs = flow.tensor.dynamic_constant dense<[[6,7],[8,9]]> : tensor<2x2xi32> -> tensor<?x?xi32> + %rhs = flow.tensor.dynamic_constant dense<[[[[0,1],[1,0]],[[0,0],[1,1]]],[[[1,1],[0,0]],[[0,1],[1,0]]]]> : tensor<2x2x2x2xi32> -> tensor<?x?x?x?xi32> %0 = "stablehlo.torch_index_select"(%lhs, %rhs) {batch_dims = 0 : i64, dim = 0 : i64} : (tensor<?x?xi32>, tensor<?x?x?x?xi32>) -> tensor<?x?x?x?x?xi32> %dshape = util.optimization_barrier %0 : tensor<?x?x?x?x?xi32> %result = tensor.cast %dshape : tensor<?x?x?x?x?xi32> to tensor<2x2x2x2x2xi32>
diff --git a/tests/e2e/regression/dynamic_torch_index_select_negative.mlir b/tests/e2e/regression/dynamic_torch_index_select_negative.mlir index c2a4a60..9e4e6d5 100644 --- a/tests/e2e/regression/dynamic_torch_index_select_negative.mlir +++ b/tests/e2e/regression/dynamic_torch_index_select_negative.mlir
@@ -1,6 +1,6 @@ func.func @torch_index_select1() { - %lhs = flow.tensor.constant dense<[[[100, 101],[110, 111]],[[200, 201],[210, 211]]]> : tensor<2x2x2xi32> -> tensor<?x?x?xi32> - %rhs = flow.tensor.constant dense<[[[0, 1],[1, 0]],[[0, 0],[1, 1]]]> : tensor<2x2x2xi32> -> tensor<?x?x?xi32> + %lhs = flow.tensor.dynamic_constant dense<[[[100, 101],[110, 111]],[[200, 201],[210, 211]]]> : tensor<2x2x2xi32> -> tensor<?x?x?xi32> + %rhs = flow.tensor.dynamic_constant dense<[[[0, 1],[1, 0]],[[0, 0],[1, 1]]]> : tensor<2x2x2xi32> -> tensor<?x?x?xi32> %0 = "stablehlo.torch_index_select"(%lhs, %rhs) {batch_dims = -1 : i64, dim = -1 : i64} : (tensor<?x?x?xi32>, tensor<?x?x?xi32>) -> tensor<?x?x?xi32> %dshape = util.optimization_barrier %0 : tensor<?x?x?xi32> %result = tensor.cast %dshape : tensor<?x?x?xi32> to tensor<2x2x2xi32>
diff --git a/tests/e2e/regression/dynamic_torch_index_select_scalar.mlir b/tests/e2e/regression/dynamic_torch_index_select_scalar.mlir index 06c5540..b77cc51 100644 --- a/tests/e2e/regression/dynamic_torch_index_select_scalar.mlir +++ b/tests/e2e/regression/dynamic_torch_index_select_scalar.mlir
@@ -1,5 +1,5 @@ func.func @torch_index_select1() { - %lhs = flow.tensor.constant + %lhs = flow.tensor.dynamic_constant dense<[[[1,2,3,4,5]], [[6,7,8,9,10]], [[11,12,13,14,15]], @@ -15,7 +15,7 @@ } func.func @torch_index_select2() { - %lhs = flow.tensor.constant + %lhs = flow.tensor.dynamic_constant dense<[[[1,2,3,4,5]], [[6,7,8,9,10]], [[11,12,13,14,15]],
diff --git a/tests/e2e/regression/dynamic_torch_index_select_vector.mlir b/tests/e2e/regression/dynamic_torch_index_select_vector.mlir index c0662c4..57cfa28 100644 --- a/tests/e2e/regression/dynamic_torch_index_select_vector.mlir +++ b/tests/e2e/regression/dynamic_torch_index_select_vector.mlir
@@ -1,7 +1,7 @@ func.func @torch_index_select1() { - %lhs = flow.tensor.constant + %lhs = flow.tensor.dynamic_constant dense<[[[1, 2],[3, 4]],[[5, 6],[7, 8]],[[9, 10],[11, 12]]]> : tensor<3x2x2xi32> -> tensor<?x?x?xi32> - %rhs = flow.tensor.constant dense<[0, 1]> : tensor<2xi32> -> tensor<?xi32> + %rhs = flow.tensor.dynamic_constant dense<[0, 1]> : tensor<2xi32> -> tensor<?xi32> %0 = "stablehlo.torch_index_select"(%lhs, %rhs) {batch_dims = 0 : i64, dim = 1 : i64} : (tensor<?x?x?xi32>, tensor<?xi32>) -> tensor<?x?x?xi32> %dshape = util.optimization_barrier %0 : tensor<?x?x?xi32> %result = tensor.cast %dshape : tensor<?x?x?xi32> to tensor<3x2x2xi32> @@ -13,9 +13,9 @@ } func.func @torch_index_select2() { - %lhs = flow.tensor.constant + %lhs = flow.tensor.dynamic_constant dense<[[[1, 2],[3, 4]],[[5, 6],[7, 8]],[[9, 10],[11, 12]]]> : tensor<3x2x2xi32> -> tensor<?x?x?xi32> - %rhs = flow.tensor.constant dense<[0, 1]> : tensor<2xi32> -> tensor<?xi32> + %rhs = flow.tensor.dynamic_constant dense<[0, 1]> : tensor<2xi32> -> tensor<?xi32> %0 = "stablehlo.torch_index_select"(%lhs, %rhs) {batch_dims = 0 : i64, dim = 0 : i64} : (tensor<?x?x?xi32>, tensor<?xi32>) -> tensor<?x?x?xi32> %dshape = util.optimization_barrier %0 : tensor<?x?x?xi32> %result = tensor.cast %dshape : tensor<?x?x?xi32> to tensor<2x2x2xi32>
diff --git a/tests/e2e/regression/layernorm.mlir b/tests/e2e/regression/layernorm.mlir index 5cbc0df..15f02dec 100644 --- a/tests/e2e/regression/layernorm.mlir +++ b/tests/e2e/regression/layernorm.mlir
@@ -76,9 +76,9 @@ %cst = arith.constant 1.000000e+00 : f32 %cst_0 = arith.constant 0.000000e+00 : f32 %cst_1 = arith.constant dense<0.000000e+00> : tensor<128x384xf32> - %cst_2 = flow.tensor.constant dense<9.99999996E-13> : tensor<128x1xf32> -> tensor<?x1xf32> - %cst_3 = flow.tensor.constant dense<3.840000e+02> : tensor<128x1xf32> -> tensor<?x1xf32> - %cst_4 = flow.tensor.constant dense<5.000000e+00> : tensor<128x384xf32> -> tensor<?x?xf32> + %cst_2 = flow.tensor.dynamic_constant dense<9.99999996E-13> : tensor<128x1xf32> -> tensor<?x1xf32> + %cst_3 = flow.tensor.dynamic_constant dense<3.840000e+02> : tensor<128x1xf32> -> tensor<?x1xf32> + %cst_4 = flow.tensor.dynamic_constant dense<5.000000e+00> : tensor<128x384xf32> -> tensor<?x?xf32> %c_0_index = arith.constant 0 : index %c_1_index = arith.constant 1 : index %dim_0 = tensor.dim %cst_4, %c_0_index : tensor<?x?xf32>
diff --git a/tests/e2e/regression/reduction_broadcast_elementwise.mlir b/tests/e2e/regression/reduction_broadcast_elementwise.mlir index 1ce6bd5..38c108c 100644 --- a/tests/e2e/regression/reduction_broadcast_elementwise.mlir +++ b/tests/e2e/regression/reduction_broadcast_elementwise.mlir
@@ -40,7 +40,7 @@ func.func @max_sub_exp_dynamic() { %cst = arith.constant -3.40282347E+38 : f32 %cst_0 = arith.constant dense<1.000000e+00> : tensor<12x128x128xf32> - %cst_1 = flow.tensor.constant dense<5.000000e+00> : tensor<12x128x128xf32> -> tensor<?x?x?xf32> + %cst_1 = flow.tensor.dynamic_constant dense<5.000000e+00> : tensor<12x128x128xf32> -> tensor<?x?x?xf32> %c_0_index = arith.constant 0 : index %c_1_index = arith.constant 1 : index %c_2_index = arith.constant 2 : index
diff --git a/tests/e2e/regression/softmax.mlir b/tests/e2e/regression/softmax.mlir index fb77441..19fb5bc 100644 --- a/tests/e2e/regression/softmax.mlir +++ b/tests/e2e/regression/softmax.mlir
@@ -62,7 +62,7 @@ %cst_0 = arith.constant 0.000000e+00 : f32 %cst_1 = arith.constant -3.40282347E+38 : f32 %cst_2 = arith.constant dense<7.812500e-03> : tensor<12x128x128xf32> - %cst_3 = flow.tensor.constant dense<5.000000e+00> : tensor<12x128x128xf32> -> tensor<?x?x?xf32> + %cst_3 = flow.tensor.dynamic_constant dense<5.000000e+00> : tensor<12x128x128xf32> -> tensor<?x?x?xf32> %c_0_index = arith.constant 0 : index %c_1_index = arith.constant 1 : index %c_2_index = arith.constant 2 : index
diff --git a/tests/e2e/stablehlo_ops/scatter_dynamic.mlir b/tests/e2e/stablehlo_ops/scatter_dynamic.mlir index 50ae47c..0127585 100644 --- a/tests/e2e/stablehlo_ops/scatter_dynamic.mlir +++ b/tests/e2e/stablehlo_ops/scatter_dynamic.mlir
@@ -1,7 +1,7 @@ func.func @scatter_add_slice_2D_dynamic_num_updates() { %arg0 = util.unfoldable_constant dense<1> : tensor<6x3xi32> - %arg1 = flow.tensor.constant dense<[[2], [4]]> : tensor<2x1xi32> -> tensor<?x1xi32> - %arg2 = flow.tensor.constant dense<[[1, 2, 3], + %arg1 = flow.tensor.dynamic_constant dense<[[2], [4]]> : tensor<2x1xi32> -> tensor<?x1xi32> + %arg2 = flow.tensor.dynamic_constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> -> tensor<?x3xi32> %0 = "stablehlo.scatter"(%arg0, %arg1, %arg2) ( { ^bb0(%arg3: tensor<i32>, %arg4: tensor<i32>): // no predecessors
diff --git a/tests/e2e/tensor_ops/pack.mlir b/tests/e2e/tensor_ops/pack.mlir index 3837d2e..c2aa676 100644 --- a/tests/e2e/tensor_ops/pack.mlir +++ b/tests/e2e/tensor_ops/pack.mlir
@@ -28,7 +28,7 @@ } func.func @dynamic_pack_simple() { - %iree_input = flow.tensor.constant dense<[ + %iree_input = flow.tensor.dynamic_constant dense<[ [0, 1, 2, 3], [4, 5, 6, 7], [8, 9, 10, 11], @@ -69,7 +69,7 @@ } func.func @dynamic_pack_simple_pad_mode() { - %iree_input = flow.tensor.constant dense<[ + %iree_input = flow.tensor.dynamic_constant dense<[ [0, 1, 2, 3], [4, 5, 6, 7], [8, 9, 10, 11],
diff --git a/tests/e2e/tensor_ops/pack_dynamic_inner_tiles.mlir b/tests/e2e/tensor_ops/pack_dynamic_inner_tiles.mlir index a2b1b5c..d75eca5 100644 --- a/tests/e2e/tensor_ops/pack_dynamic_inner_tiles.mlir +++ b/tests/e2e/tensor_ops/pack_dynamic_inner_tiles.mlir
@@ -19,7 +19,7 @@ } func.func @fully_dynamic_pack_simple() { - %iree_input = flow.tensor.constant dense<[ + %iree_input = flow.tensor.dynamic_constant dense<[ [0, 1, 2, 3], [4, 5, 6, 7], [8, 9, 10, 11],
diff --git a/tests/e2e/tensor_ops/unpack.mlir b/tests/e2e/tensor_ops/unpack.mlir index b187b4f..e5b0df4 100644 --- a/tests/e2e/tensor_ops/unpack.mlir +++ b/tests/e2e/tensor_ops/unpack.mlir
@@ -31,7 +31,7 @@ } func.func @dynamic_unpack_simple() { - %iree_input = flow.tensor.constant dense<[[[[0, 1], [4, 5]], [[2, 3], [6, 7]]], [[[8, 9], [12, 13]], [[10 ,11], [14, 15]]]]> : tensor<2x2x2x2xi32> -> tensor<?x?x2x2xi32> + %iree_input = flow.tensor.dynamic_constant dense<[[[[0, 1], [4, 5]], [[2, 3], [6, 7]]], [[[8, 9], [12, 13]], [[10 ,11], [14, 15]]]]> : tensor<2x2x2x2xi32> -> tensor<?x?x2x2xi32> %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %c2 = arith.constant 2 : index @@ -60,7 +60,7 @@ } func.func @dynamic_unpack_simple_extract_slice() { - %iree_input = flow.tensor.constant dense<[[[[0, 1, 2], [4, 5, 6], [8, 9, 10]], + %iree_input = flow.tensor.dynamic_constant dense<[[[[0, 1, 2], [4, 5, 6], [8, 9, 10]], [[3, 0, 0], [7, 0, 0], [11, 0, 0]]], [[[12, 13, 14], [0, 0, 0], [0, 0, 0]], [[15, 0, 0], [0, 0, 0], [0, 0, 0]]]]> : tensor<2x2x3x3xi32> -> tensor<?x?x3x3xi32>
diff --git a/tests/microbenchmarks/dynamic_shape_vectorization.mlir b/tests/microbenchmarks/dynamic_shape_vectorization.mlir index f7ce096..be10d3a 100644 --- a/tests/microbenchmarks/dynamic_shape_vectorization.mlir +++ b/tests/microbenchmarks/dynamic_shape_vectorization.mlir
@@ -14,9 +14,9 @@ %dim1 = util.unfoldable_constant 513 : index %dim2 = util.unfoldable_constant 385 : index - %A = flow.tensor.constant dense<1.0> : tensor<513x257xf32> -> tensor<?x?xf32> - %B = flow.tensor.constant dense<2.0> : tensor<257x385xf32> -> tensor<?x?xf32> - %C = flow.tensor.constant dense<0.0> : tensor<513x385xf32> -> tensor<?x?xf32> + %A = flow.tensor.dynamic_constant dense<1.0> : tensor<513x257xf32> -> tensor<?x?xf32> + %B = flow.tensor.dynamic_constant dense<2.0> : tensor<257x385xf32> -> tensor<?x?xf32> + %C = flow.tensor.dynamic_constant dense<0.0> : tensor<513x385xf32> -> tensor<?x?xf32> %gemm = linalg.matmul ins(%A, %B : tensor<?x?xf32>, tensor<?x?xf32>) @@ -26,9 +26,9 @@ func.func @dynamic_elw() -> tensor<?x?xf32> { %c0 = arith.constant 0.000000e+00 : f32 - %A = flow.tensor.constant dense<1.0> : tensor<513x1025xf32> -> tensor<?x?xf32> - %B = flow.tensor.constant dense<2.0> : tensor<513x1025xf32> -> tensor<?x?xf32> - %C = flow.tensor.constant dense<0.0> : tensor<513x1025xf32> -> tensor<?x?xf32> + %A = flow.tensor.dynamic_constant dense<1.0> : tensor<513x1025xf32> -> tensor<?x?xf32> + %B = flow.tensor.dynamic_constant dense<2.0> : tensor<513x1025xf32> -> tensor<?x?xf32> + %C = flow.tensor.dynamic_constant dense<0.0> : tensor<513x1025xf32> -> tensor<?x?xf32> %gen = linalg.generic { indexing_maps = [
diff --git a/tools/test/iree-run-module-outputs.mlir b/tools/test/iree-run-module-outputs.mlir index fc3157f..66765cf 100644 --- a/tools/test/iree-run-module-outputs.mlir +++ b/tools/test/iree-run-module-outputs.mlir
@@ -24,7 +24,7 @@ %1 = arith.constant dense<4.0> : tensor<f32> // OUTPUT-DEFAULT: result[2]: hal.buffer_view // OUTPUT-DEFAULT-NEXT: 2x4xi32=[0 1 2 3][4 5 6 7] - %2 = flow.tensor.constant dense<[[0,1,2,3],[4,5,6,7]]> : tensor<2x4xi32> -> tensor<?x4xi32> + %2 = flow.tensor.dynamic_constant dense<[[0,1,2,3],[4,5,6,7]]> : tensor<2x4xi32> -> tensor<?x4xi32> return %0, %1, %2 : i32, tensor<f32>, tensor<?x4xi32> } @@ -48,7 +48,7 @@ %1 = arith.constant dense<4.0> : tensor<f32> // OUTPUT-NUMPY-NEXT{LITERAL}: [[0 1 2 3] // OUTPUT-NUMPY-NEXT{LITERAL}: [4 5 6 7]] - %2 = flow.tensor.constant dense<[[0,1,2,3],[4,5,6,7]]> : tensor<2x4xi32> -> tensor<?x4xi32> + %2 = flow.tensor.dynamic_constant dense<[[0,1,2,3],[4,5,6,7]]> : tensor<2x4xi32> -> tensor<?x4xi32> return %0, %1, %2 : i32, tensor<f32>, tensor<?x4xi32> } @@ -90,7 +90,7 @@ func.func @write_binary() -> (tensor<f32>, tensor<?x4xi32>) { %0 = arith.constant dense<4.0> : tensor<f32> - %1 = flow.tensor.constant dense<[[0,1,2,3],[4,5,6,7]]> : tensor<2x4xi32> -> tensor<?x4xi32> + %1 = flow.tensor.dynamic_constant dense<[[0,1,2,3],[4,5,6,7]]> : tensor<2x4xi32> -> tensor<?x4xi32> return %0, %1 : tensor<f32>, tensor<?x4xi32> } func.func @echo_binary(%arg0: tensor<f32>, %arg1: tensor<?x4xi32>) -> (tensor<f32>, tensor<?x4xi32>) {