These were largely leftover from when MLIR was a google project, and don't really follow LLVM guidelines.
decimal-literal ::= digit+
hexadecimal-literal ::= `0x` hex_digit+
float-literal ::= [-+]?[0-9]+[.][0-9]*([eE][-+]?[0-9]+)?
-string-literal ::= `"` [^"\n\f\v\r]* `"` TODO define escaping rules
+string-literal ::= `"` [^"\n\f\v\r]* `"` TODO: define escaping rules
```
Not listed here, but MLIR does support comments. They use standard BCPL syntax,
replaced by the operand/result's type. E.g., for `F32` in `F32:$operand`, its
`$_self` will be expanded as `getOperand(...).getType()`.
-TODO(b/130663252): Reconsider the leading symbol for special placeholders.
-Eventually we want to allow referencing operand/result $-names; such $-names
-can start with underscore.
+TODO: Reconsider the leading symbol for special placeholders. Eventually we want
+to allow referencing operand/result $-names; such $-names can start with
+underscore.
For example, to write an attribute `attr` is an `IntegerAttr`, in C++ you can
just call `attr.isa<IntegerAttr>()`. The code can be wrapped in a `CPred` as
### Quantized type
-TODO : Flesh this section out.
+TODO: Flesh this section out.
* QuantizedType base class
* UniformQuantizedType
## TFLite native quantization
-TODO : Flesh this out
+TODO: Flesh this out
### General algorithm
}
```
-TODO (Add more examples showing the IR for a variety of interesting cases)
+TODO: (Add more examples showing the IR for a variety of interesting cases)
## Design alternatives and extensions
/// Returns the rank of the `function` argument at position `pos`.
/// If the argument is of MemRefType, this returns the rank of the MemRef.
/// Otherwise returns `0`.
-/// TODO(ntv): support more than MemRefType and scalar Type.
+/// TODO: support more than MemRefType and scalar Type.
unsigned getRankOfFunctionArgument(mlir_func_t function, unsigned pos);
/// Returns an opaque mlir::Type of the `function` argument at position `pos`.
/// used to add appropriate inequalities. Any symbols founds in the bound
/// operands are added as symbols in the system. Returns failure for the yet
/// unimplemented cases.
-// TODO(bondhugula): handle non-unit strides.
+// TODO: handle non-unit strides.
LogicalResult getIndexSet(MutableArrayRef<AffineForOp> forOps,
FlatAffineConstraints *domain);
SmallVector<Value, 4> indices;
/// Constructs a MemRefAccess from a load or store operation.
- // TODO(b/119949820): add accessors to standard op's load, store, DMA op's to
- // return MemRefAccess, i.e., loadOp->getAccess(), dmaOp->getRead/WriteAccess.
+ // TODO: add accessors to standard op's load, store, DMA op's to return
+ // MemRefAccess, i.e., loadOp->getAccess(), dmaOp->getRead/WriteAccess.
explicit MemRefAccess(Operation *opInst);
// Returns the rank of the memref associated with this access.
/// access the same memref element. If 'allowRAR' is true, will consider
/// read-after-read dependences (typically used by applications trying to
/// optimize input reuse).
-// TODO(andydavis) Wrap 'dependenceConstraints' and 'dependenceComponents' into
-// a single struct.
-// TODO(andydavis) Make 'dependenceConstraints' optional arg.
+// TODO: Wrap 'dependenceConstraints' and 'dependenceComponents' into a single
+// struct.
+// TODO: Make 'dependenceConstraints' optional arg.
struct DependenceResult {
enum ResultEnum {
HasDependence, // A dependence exists between 'srcAccess' and 'dstAccess'.
/// Create a flat affine constraint system from an AffineValueMap or a list of
/// these. The constructed system will only include equalities.
- // TODO(bondhugula)
explicit FlatAffineConstraints(const AffineValueMap &avm);
explicit FlatAffineConstraints(ArrayRef<const AffineValueMap *> avmRef);
/// 'affine.for' operation are added as trailing identifiers (either
/// dimensional or symbolic depending on whether the operand is a valid
/// symbol).
- // TODO(bondhugula): add support for non-unit strides.
+ // TODO: add support for non-unit strides.
LogicalResult addAffineForOpDomain(AffineForOp forOp);
/// Adds a lower or an upper bound for the identifier at the specified
/// Projects out (aka eliminates) 'num' identifiers starting at position
/// 'pos'. The resulting constraint system is the shadow along the dimensions
/// that still exist. This method may not always be integer exact.
- // TODO(bondhugula): deal with integer exactness when necessary - can return a
- // value to mark exactness for example.
+ // TODO: deal with integer exactness when necessary - can return a value to
+ // mark exactness for example.
void projectOut(unsigned pos, unsigned num);
inline void projectOut(unsigned pos) { return projectOut(pos, 1); }
/// multi-result map. The trip count expression is simplified before returning.
/// This method only utilizes map composition to construct lower and upper
/// bounds before computing the trip count expressions
-// TODO(mlir-team): this should be moved into 'Transforms/' and be replaced by a
-// pure analysis method relying on FlatAffineConstraints
+// TODO: this should be moved into 'Transforms/' and be replaced by a pure
+// analysis method relying on FlatAffineConstraints
void buildTripCountMapAndOperands(AffineForOp forOp, AffineMap *map,
SmallVectorImpl<Value> *operands);
/// Checks whether the loop is structurally vectorizable; i.e.:
/// 1. no conditionals are nested under the loop;
/// 2. all nested load/stores are to scalar MemRefs.
-/// TODO(ntv): relax the no-conditionals restriction
+/// TODO: relax the no-conditionals restriction
bool isVectorizableLoopBody(AffineForOp loop,
NestedPattern &vectorTransferMatcher);
/// Checks where SSA dominance would be violated if a for op's body
/// operations are shifted by the specified shifts. This method checks if a
/// 'def' and all its uses have the same shift factor.
-// TODO(mlir-team): extend this to check for memory-based dependence
-// violation when we have the support.
+// TODO: extend this to check for memory-based dependence violation when we have
+// the support.
bool isOpwiseShiftValid(AffineForOp forOp, ArrayRef<uint64_t> shifts);
} // end namespace mlir
/// 1. recursively matches a substructure in the tree;
/// 2. uses a filter function to refine matches with extra semantic
/// constraints (passed via a lambda of type FilterFunctionType);
-/// 3. TODO(ntv) optionally applies actions (lambda).
+/// 3. TODO: optionally applies actions (lambda).
///
/// Nested patterns are meant to capture imperfectly nested loops while matching
/// properties over the whole loop nest. For instance, in vectorization we are
/// Populates 'loops' with IVs of the loops surrounding 'op' ordered from
/// the outermost 'affine.for' operation to the innermost one.
-// TODO(bondhugula): handle 'affine.if' ops.
+// TODO: handle 'affine.if' ops.
void getLoopIVs(Operation &op, SmallVectorImpl<AffineForOp> *loops);
/// Returns the nesting depth of this operation, i.e., the number of loops
/// surrounding ops in 'opsB', as a function of IVs and symbols of loop nest
/// surrounding ops in 'opsA' at 'loopDepth'.
/// Returns 'success' if union was computed, 'failure' otherwise.
-// TODO(andydavis) Change this API to take 'forOpA'/'forOpB'.
+// TODO: Change this API to take 'forOpA'/'forOpB'.
LogicalResult computeSliceUnion(ArrayRef<Operation *> opsA,
ArrayRef<Operation *> opsB, unsigned loopDepth,
unsigned numCommonLoops, bool isBackwardSlice,
// Loop depth is a crucial optimization choice that determines where to
// materialize the results of the backward slice - presenting a trade-off b/w
// storage and redundant computation in several cases.
-// TODO(andydavis) Support computation slices with common surrounding loops.
+// TODO: Support computation slices with common surrounding loops.
AffineForOp insertBackwardComputationSlice(Operation *srcOpInst,
Operation *dstOpInst,
unsigned dstLoopDepth,
/// identifiers since getMemRefRegion() is called with a specific loop depth,
/// and thus the region is symbolic in the outer surrounding loops at that
/// depth.
- // TODO(bondhugula): Replace this to exploit HyperRectangularSet.
+ // TODO: Replace this to exploit HyperRectangularSet.
FlatAffineConstraints cst;
};
I32:$k,
VectorOfLengthAndType<[16, 8], [F32, F64]>:$a,
AnyTypeOf<[I16, I8]>:$imm,
- // TODO(ntv): figure rounding out (optional operand?).
+ // TODO: figure rounding out (optional operand?).
I32:$rounding
);
let results = (outs VectorOfLengthAndType<[16, 8], [F32, F64]>:$dst);
VectorOfLengthAndType<[16, 8], [F32, F64]>:$a,
VectorOfLengthAndType<[16, 8], [F32, F64]>:$b,
AnyTypeOf<[I16, I8]>:$k,
- // TODO(ntv): figure rounding out (optional operand?).
+ // TODO: figure rounding out (optional operand?).
I32:$rounding
);
let results = (outs VectorOfLengthAndType<[16, 8], [F32, F64]>:$dst);
// affine.dma_start %src[%i, %j], %dst[%k, %l], %tag[%idx], %num_elements,
// %stride, %num_elt_per_stride : ...
//
-// TODO(mlir-team): add additional operands to allow source and destination
-// striding, and multiple stride levels (possibly using AffineMaps to specify
-// multiple levels of striding).
-// TODO(andydavis) Consider replacing src/dst memref indices with view memrefs.
+// TODO: add additional operands to allow source and destination striding, and
+// multiple stride levels (possibly using AffineMaps to specify multiple levels
+// of striding).
+// TODO: Consider replacing src/dst memref indices with view memrefs.
class AffineDmaStartOp : public Op<AffineDmaStartOp, OpTrait::VariadicOperands,
OpTrait::ZeroResult> {
public:
/// Returns the number of workgroup (thread, block) dimensions supported in
/// the GPU dialect.
- // TODO(zinenko,herhut): consider generalizing this.
+ // TODO: consider generalizing this.
static unsigned getNumWorkgroupDimensions() { return 3; }
/// Returns the numeric value used to identify the workgroup memory address
/// - the argument/result attributes may need an update: if the new type
/// has less parameters we drop the extra attributes, if there are more
/// parameters they won't have any attributes.
- // TODO(b/146349912): consider removing this function thanks to rewrite
- // patterns.
+ // TODO: consider removing this function thanks to rewrite patterns.
void setType(FunctionType newType);
/// Returns the number of buffers located in the workgroup memory.
}
/// Helper function to create a ParallelDimMapperAttr.
-/// TODO(ravishankarm/antiagainst): Replace its uses with an auto-gened method.
+/// TODO: Replace its uses with an auto-gened method.
ParallelLoopDimMapping getParallelLoopDimMappingAttr(Processor processor,
AffineMap map,
AffineMap bound);
/// the current ScopedContext, at the current insert point.
void macRegionBuilder(ValueRange args);
-/// TODO(ntv): In the future we should tie these implementations to something in
+/// TODO: In the future we should tie these implementations to something in
/// Tablegen that generates the proper interfaces and the proper sugared named
/// ops.
StructuredIndexed I2,
StructuredIndexed O);
-// TODO(ntv): Implement more useful pointwise operations on a per-need basis.
+// TODO: Implement more useful pointwise operations on a per-need basis.
using MatmulRegionBuilder = function_ref<void(ValueRange args)>;
///
/// For now `...` must be empty (i.e. only 2-D convolutions are supported).
///
-// TODO(ntv) Extend convolution rank with some template magic.
+// TODO: Extend convolution rank with some template magic.
Operation *linalg_generic_conv_nhwc(Value vI, Value vW, Value vO,
ArrayRef<int> strides = {},
ArrayRef<int> dilations = {});
///
/// For now `...` must be empty (i.e. only 2-D convolutions are supported).
///
-// TODO(ntv) Extend convolution rank with some template magic.
+// TODO: Extend convolution rank with some template magic.
Operation *linalg_generic_dilated_conv_nhwc(Value vI, Value vW, Value vO,
int depth_multiplier = 1,
ArrayRef<int> strides = {},
OptionalAttr<AffineMapAttr>:$inputPermutation,
OptionalAttr<AffineMapAttr>:$outputPermutation);
- // TODO(ntv) this should go away once the usage of OptionalAttr triggers
- // emission of builders with default arguments left unspecified.
+ // TODO: this should go away once the usage of OptionalAttr triggers emission
+ // of builders with default arguments left unspecified.
let builders = [OpBuilder<
"OpBuilder &builder, OperationState &result, Value input, Value output", [{
return build(
OptionalAttr<I64ElementsAttr>:$padding);
let extraClassDeclaration = commonUtils # [{
- // TODO(ntv) extend to support more than 1 dimensions and potentially
- // grouping too.
+ // TODO: extend to support more than 1 dimensions and potentially grouping
+ // too.
unsigned getNumBatchDimensions() { return 1; }
unsigned getNumInputFeatureDimensions() { return 1; }
cast<ConcreteType>(this->getOperation()).referenceIterators();
// If there is no reference, this must be a generic op.
- // TODO(ntv): Traits are used to define ops. Split into cpp to avoid
- // cyclic dependency.
+ // TODO: Traits are used to define ops. Split into cpp to avoid cyclic
+ // dependency.
auto name = this->getOperation()->getName().getStringRef();
if (!maybeReferenceIteratorTypes && name != "generic" &&
name != "indexed_generic") {
return StringAttr::get(str, ctx);
});
auto attr = ArrayAttr::get(llvm::to_vector<4>(attrRange), ctx);
- // TODO(ntv): Need to memoize this. Can't just store as an attribute atm as
- // it will impact parser, printer and tests.
+ // TODO: Need to memoize this. Can't just store as an attribute atm as it
+ // will impact parser, printer and tests.
// this->getOperation()->setAttr("iterator_types", attr);
return attr;
}
});
SmallVector<Attribute, 4> attrs{attrRange.begin(), attrRange.end()};
auto attr = ArrayAttr::get(attrs, ctx);
- // TODO(ntv): Need to memoize this. Can't just store as an attribute atm as
- // it will impact parser, printer and tests.
+ // TODO: Need to memoize this. Can't just store as an attribute atm as it
+ // will impact parser, printer and tests.
// this->getOperation()->setAttr("indexing_maps", attr);
return attr;
}
AND (`Extension::C`) AND (`Extension::D` OR `Extension::E`) is enabled.
}];
- // TODO(antiagainst): Returning SmallVector<ArrayRef<...>> is not recommended.
+ // TODO: Returning SmallVector<ArrayRef<...>> is not recommended.
// Find a better way for this.
let queryFnRetType = "::llvm::SmallVector<::llvm::ArrayRef<"
"::mlir::spirv::Extension>, 1>";
let instance = "ref";
}
-// TODO(antiagainst): the following interfaces definitions are duplicating with
-// the above. Remove them once we are able to support dialect-specific contents
-// in ODS.
+// TODO: the following interfaces definitions are duplicating with the above.
+// Remove them once we are able to support dialect-specific contents in ODS.
def QueryMinVersionInterface : OpInterface<"QueryMinVersionInterface"> {
let methods = [InterfaceMethod<"", "::mlir::spirv::Version", "getMinVersion">];
}
def SPV_IOrUIVec4 : SPV_Vec4<SPV_SignlessOrUnsignedInt>;
def SPV_Int32Vec4 : SPV_Vec4<AnyI32>;
-// TODO(ravishankarm): From 1.4, this should also include Composite type.
+// TODO: From 1.4, this should also include Composite type.
def SPV_SelectType : AnyTypeOf<[SPV_Scalar, SPV_Vector, SPV_AnyPtr]>;
//===----------------------------------------------------------------------===//
// Base class for all SPIR-V ops.
class SPV_Op<string mnemonic, list<OpTrait> traits = []> :
Op<SPIRV_Dialect, mnemonic, !listconcat(traits, [
- // TODO(antiagainst): We don't need all of the following traits for
- // every op; only the suitable ones should be added automatically
- // after ODS supports dialect-specific contents.
+ // TODO: We don't need all of the following traits for every op; only
+ // the suitable ones should be added automatically after ODS supports
+ // dialect-specific contents.
DeclareOpInterfaceMethods<QueryMinVersionInterface>,
DeclareOpInterfaceMethods<QueryMaxVersionInterface>,
DeclareOpInterfaceMethods<QueryExtensionInterface>,
/// satisfy shader interface requirements: shader interface types must be
/// pointers to structs.
///
-/// TODO(antiagainst): We might want to introduce a way to control how
-/// unsupported bitwidth are handled and explicitly fail if wanted.
+/// TODO: We might want to introduce a way to control how unsupported bitwidth
+/// are handled and explicitly fail if wanted.
class SPIRVTypeConverter : public TypeConverter {
public:
explicit SPIRVTypeConverter(spirv::TargetEnvAttr targetAttr);
/// Performs the index computation to get to the element at `indices` of the
/// memory pointed to by `basePtr`, using the layout map of `baseType`.
-// TODO(ravishankarm) : This method assumes that the `baseType` is a MemRefType
-// with AffineMap that has static strides. Extend to handle dynamic strides.
+// TODO: This method assumes that the `baseType` is a MemRefType with AffineMap
+// that has static strides. Extend to handle dynamic strides.
spirv::AccessChainOp getElementPtr(SPIRVTypeConverter &typeConverter,
MemRefType baseType, Value basePtr,
ValueRange indices, Location loc,
SPV_AnyMatrix:$result
);
- // TODO (Hazem): we need just one matrix type given that the input and result
- // are the same and the scalar's type can be deduced from it.
+ // TODO: we need just one matrix type given that the input and result are the
+ // same and the scalar's type can be deduced from it.
let assemblyFormat = [{
operands attr-dict `:` type($matrix) `,` type($scalar) `->` type($result)
}];
// -----
-#endif // SPIRV_MATRIX_OPS
\ No newline at end of file
+#endif // SPIRV_MATRIX_OPS
%2 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
```
- TODO(antiagainst): support constant structs
+ TODO: support constant structs
}];
let arguments = (ins
spv.specConstant @spec_const2 spec_id(5) = 42 : i32
```
- TODO(antiagainst): support composite spec constants with another op
+ TODO: support composite spec constants with another op
}];
let arguments = (ins
ImageSamplingInfo getSamplingInfo() const;
ImageSamplerUseInfo getSamplerUseInfo() const;
ImageFormat getImageFormat() const;
- // TODO(ravishankarm): Add support for Access qualifier
+ // TODO: Add support for Access qualifier
void getExtensions(SPIRVType::ExtensionArrayRefVector &extensions,
Optional<spirv::StorageClass> storage = llvm::None);
// Shape constraint related ops.
//===----------------------------------------------------------------------===//
-//TODO(tpopp): Move the code below and witnesses to a different file.
+// TODO: Move the code below and witnesses to a different file.
def Shape_AnyOp : Shape_Op<"any", [Commutative, NoSideEffect]> {
let summary = "Return any combination of the input shapes";
let description = [{
/// MemRef. It has placeholders for non-contiguous tensors that fit within the
/// Fortran subarray model.
/// At the moment it can only capture a MemRef with an identity layout map.
-// TODO(ntv): Support MemRefs with layoutMaps.
+// TODO: Support MemRefs with layoutMaps.
class MemRefBoundsCapture : public BoundsCapture {
public:
explicit MemRefBoundsCapture(Value v);
// dma_start %src[%i, %j], %dst[%k, %l], %num_elements, %tag[%idx], %stride,
// %num_elt_per_stride :
//
-// TODO(mlir-team): add additional operands to allow source and destination
-// striding, and multiple stride levels.
-// TODO(andydavis) Consider replacing src/dst memref indices with view memrefs.
+// TODO: add additional operands to allow source and destination striding, and
+// multiple stride levels.
+// TODO: Consider replacing src/dst memref indices with view memrefs.
class DmaStartOp
: public Op<DmaStartOp, OpTrait::VariadicOperands, OpTrait::ZeroResult> {
public:
let parser = [{ return ::parse$cppClass(parser, result); }];
}
-// TODO(andydavis, ntv) Add an attribute to specify a different algebra
-// with operators other than the current set: {*, +}.
+// TODO: Add an attribute to specify a different algebra with operators other
+// than the current set: {*, +}.
def Vector_ContractionOp :
Vector_Op<"contract", [
NoSideEffect,
}];
}
-// TODO(andydavis) Add transformation which decomposes ReshapeOp into an
-// optimized sequence of vector rotate/shuffle/select operations.
+// TODO: Add transformation which decomposes ReshapeOp into an optimized
+// sequence of vector rotate/shuffle/select operations.
def Vector_ReshapeOp :
Vector_Op<"reshape", [AttrSizedOperandSegments, NoSideEffect]>,
Arguments<(ins AnyVector:$vector, Variadic<Index>:$input_shape,
attribute and extracts the n-D subvector at the proper offset.
At the moment strides must contain only 1s.
- // TODO(ntv) support non-1 strides.
+ // TODO: support non-1 strides.
Returns an n-D vector where the first k-D dimensions match the `sizes`
attribute. The returned subvector contains the elements starting at offset
{offsets = [0, 2], sizes = [2, 4], strides = [1, 1]}:
vector<4x8x16xf32> to vector<2x4x16xf32>
- // TODO(ntv) Evolve to a range form syntax similar to:
+ // TODO: Evolve to a range form syntax similar to:
%1 = vector.extract_strided_slice %0[0:2:1][2:4:1]
vector<4x8x16xf32> to vector<2x4x16xf32>
```
PredOpTrait<"rhs operand and result have same element type",
TCresVTEtIsSameAsOpBase<0, 1>>]>,
Arguments<(
- // TODO(ntv, fhahn): tighten vector element types that make sense.
+ // TODO: tighten vector element types that make sense.
ins VectorOfRankAndType<[1],
[AnySignlessInteger, AnySignedInteger, AnyFloat]>:$lhs,
VectorOfRankAndType<[1],
PredOpTrait<"source operand and result have same element type",
TCresVTEtIsSameAsOpBase<0, 0>>]>,
Arguments<(
- // TODO(ntv, fhahn, ajcbik): tighten vector element types that make sense.
+ // TODO: tighten vector element types that make sense.
ins VectorOfRankAndType<[1],
[AnySignlessInteger, AnySignedInteger, AnyFloat]>:$matrix,
I32Attr:$rows, I32Attr:$columns)>,
///
/// This only kicks in when VectorTransformsOptions is set to AXPY.
//
-// TODO (ajcbik): this is very similar, but not quite the same as
-// the outerproduct lowering above; merge the two?
+// TODO: this is very similar, but not quite the same as the outerproduct
+// lowering above; merge the two?
class ContractionOpToAXPYLowering
: public OpRewritePattern<vector::ContractionOp> {
public:
/// over other smaller vector types in the function and avoids interfering with
/// operations on those.
/// This is a first approximation, it can easily be extended in the future.
-/// TODO(ntv): this could all be much simpler if we added a bit that a vector
-/// type to mark that a vector is a strict super-vector but it still does not
-/// warrant adding even 1 extra bit in the IR for now.
+/// TODO: this could all be much simpler if we added a bit that a vector type to
+/// mark that a vector is a strict super-vector but it still does not warrant
+/// adding even 1 extra bit in the IR for now.
bool operatesOnSuperVectorsOf(Operation &op, VectorType subVectorType);
} // end namespace matcher
APInt getValue() const;
/// Return the integer value as a 64-bit int. The attribute must be a signless
/// integer.
- // TODO(jpienaar): Change callers to use getValue instead.
+ // TODO: Change callers to use getValue instead.
int64_t getInt() const;
/// Return the integer value as a signed 64-bit int. The attribute must be
/// a signed integer.
auto zeroValue = getZeroValue<T>();
auto valueIt = getValues().getValues<T>().begin();
const std::vector<ptrdiff_t> flatSparseIndices(getFlattenedSparseIndices());
- // TODO(riverriddle): Move-capture flatSparseIndices when c++14 is
- // available.
+ // TODO: Move-capture flatSparseIndices when c++14 is available.
std::function<T(ptrdiff_t)> mapFn = [=](ptrdiff_t index) {
// Try to map the current index to one of the sparse indices.
for (unsigned i = 0, e = flatSparseIndices.size(); i != e; ++i)
llvm::SMLoc convertLocToSMLoc(FileLineColLoc loc);
/// The maximum depth that a call stack will be printed.
- /// TODO(riverriddle) This should be a tunable flag.
+ /// TODO: This should be a tunable flag.
unsigned callStackLimit = 10;
std::unique_ptr<detail::SourceMgrDiagnosticHandlerImpl> impl;
// Constraint definitions
//===----------------------------------------------------------------------===//
-// TODO(b/130064155): Merge Constraints into Pred.
+// TODO: Merge Constraints into Pred.
// Base class for named constraints.
//
// Note that this does not "inherit" from Type because it would require
// duplicating Type subclasses for buildable and non-buildable cases to avoid
// diamond "inheritance".
-// TODO(zinenko): we may extend this to a more general 'Buildable' trait,
-// making some Types and some Attrs buildable.
+// TODO: we may extend this to a more general 'Buildable' trait, making some
+// Types and some Attrs buildable.
class BuildableType<code builder> {
// The builder call to invoke (if specified) to construct the BuildableType.
code builderCall = builder;
ShapedContainerType<[AnyType], And<[IsTensorTypePred, HasRankPred]>,
"ranked tensor">;
-// TODO(b/130064155) Have an easy way to add another constraint to a type.
+// TODO: Have an easy way to add another constraint to a type.
class StaticShapeTensorOf<list<Type> allowedTypes>
: Type<And<[TensorOf<allowedTypes>.predicate, HasStaticShapePred]>,
"statically shaped " # TensorOf<allowedTypes>.description>;
def F32MemRef : MemRefOf<[F32]>;
def F64MemRef : MemRefOf<[F64]>;
-// TODO(b/130064155) Have an easy way to add another constraint to a type.
+// TODO: Have an easy way to add another constraint to a type.
class MemRefRankOf<list<Type> allowedTypes, list<int> ranks> :
Type<And<[MemRefOf<allowedTypes>.predicate, HasAnyRankOfPred<ranks>]>,
StrJoin<!foreach(rank, ranks, rank # "D"), "/">.result # " " #
// use "::" as the delimiter, e.g., given "A::B", generated code will be
// placed in `namespace A { namespace B { ... } }`. To avoid placing in any
// namespace, use "".
- // TODO(b/134741431): use dialect to provide the namespace.
+ // TODO: use dialect to provide the namespace.
string cppNamespace = "";
// The name of the utility function that converts a value of the underlying
// An attribute constraint on FlatSymbolRefAttr that requires that the
// reference point to an op of `opClass` within the closest parent with a symbol
// table.
-// TODO(riverriddle) Add support for nested symbol references.
+// TODO: Add support for nested symbol references.
class ReferToOp<string opClass> : AttrConstraint<
CPred<"isa_and_nonnull<" # opClass # ">("
"::mlir::SymbolTable::lookupNearestSymbolFrom("
def FirstAttrDerivedResultType :
GenInternalOpTrait<"FirstAttrDerivedResultType">;
-// TODO(antiagainst): Turn the following into normal traits and generate
-// verification for them.
+// TODO: Turn the following into normal traits and generate verification for
+// them.
// All variadic operands of the op have the same number of values.
// A variadic operand contains an array of values whose array size is only
code verifier = ?;
// Whether this op has associated canonicalization patterns.
- // TODO(b/120163349): figure out a better way to write canonicalization
- // patterns in TableGen rules directly instead of using this marker
- // and C++ implementations.
+ // TODO: figure out a better way to write canonicalization patterns in
+ // TableGen rules directly instead of using this marker and C++
+ // implementations.
bit hasCanonicalizer = 0;
// Whether this op has a folder.
// These traits often require including "mlir/IR/TypeUtilities.h".
-// TODO(b/135033717): Improve the autogenerated error messages.
+// TODO: Improve the autogenerated error messages.
class Rank<string name> :
StrFunc<"$" # name # ".getType().cast<::mlir::ShapedType>().getRank()">;
class AllMatch<list<string> values, string description> :
PredOpTrait<description, AllMatchPred<values>>;
-// TODO(b/135032064): Only works for non-variadic.
+// TODO: Only works for non-variadic.
class AllMatchSameOperatorPred<list<string> names, string operator> :
AllMatchPred<!foreach(n, names, !subst("$_self", "$" # n, operator))>;
using ImplType = detail::ShapedTypeStorage;
using Type::Type;
- // TODO(ntv): merge these two special values in a single one used everywhere.
+ // TODO: merge these two special values in a single one used everywhere.
// Unfortunately, uses of `-1` have crept deep into the codebase now and are
// hard to track.
static constexpr int64_t kDynamicSize = -1;
/// Returns the memory space in which data referred to by this memref resides.
unsigned getMemorySpace() const;
- // TODO(ntv): merge these two special values in a single one used everywhere.
+ // TODO: merge these two special values in a single one used everywhere.
// Unfortunately, uses of `-1` have crept deep into the codebase now and are
// hard to track.
static constexpr int64_t kDynamicSize = -1;
/// constructed from a new SourceMgr with a single a MemoryBuffer wrapping
/// `attrStr`. If the passed `attrStr` has additional tokens that were not part
/// of the type, an error is emitted.
-// TODO(ntv) Improve diagnostic reporting.
+// TODO: Improve diagnostic reporting.
Attribute parseAttribute(llvm::StringRef attrStr, MLIRContext *context);
Attribute parseAttribute(llvm::StringRef attrStr, Type type);
/// constructed from a new SourceMgr with a single a MemoryBuffer wrapping
/// `typeStr`. If the passed `typeStr` has additional tokens that were not part
/// of the type, an error is emitted.
-// TODO(ntv) Improve diagnostic reporting.
+// TODO: Improve diagnostic reporting.
Type parseType(llvm::StringRef typeStr, MLIRContext *context);
/// This parses a single MLIR type to an MLIR context if it was valid. If not,
/// The specific parser to use depending on llvm::cl parser used. This is only
/// necessary because we need to provide additional methods for certain data
/// type parsers.
- /// TODO(riverriddle) We should upstream the methods in GenericOptionParser to
- /// avoid the need to do this.
+ /// TODO: We should upstream the methods in GenericOptionParser to avoid the
+ /// need to do this.
template <typename DataType>
using OptionParser =
std::conditional_t<std::is_base_of<llvm::cl::generic_parser_base,
StringRef getExtraClassDeclaration() const;
// Returns the Tablegen definition this operator was constructed from.
- // TODO(antiagainst,zinenko): do not expose the TableGen record, this is a
- // temporary solution to OpEmitter requiring a Record because Operator does
- // not provide enough methods.
+ // TODO: do not expose the TableGen record, this is a temporary solution to
+ // OpEmitter requiring a Record because Operator does not provide enough
+ // methods.
const llvm::Record &getDef() const;
// Returns the dialect of the op.
const llvm::Record &def;
// All operators.
- // TODO(antiagainst): we need a proper context manager, like MLIRContext,
- // for managing the lifetime of shared entities.
+ // TODO: we need a proper context manager, like MLIRContext, for managing the
+ // lifetime of shared entities.
RecordOperatorMap *recordOpMap;
};
return matchAndRewrite(cast<SourceOp>(op), operands, rewriter);
}
- // TODO(b/142763075): Use OperandAdaptor when it supports access to unnamed
- // operands.
+ // TODO: Use OperandAdaptor when it supports access to unnamed operands.
/// Rewrite and Match methods that operate on the SourceOp type. These must be
/// overridden by the derived pattern class.
struct ComputationSliceState;
class Operation;
-// TODO(andydavis) Extend this module to include utility functions for querying
-// fusion cost/storage reduction, and for performing the loop fusion
-// transformation.
+// TODO: Extend this module to include utility functions for querying fusion
+// cost/storage reduction, and for performing the loop fusion transformation.
struct FusionResult {
enum ResultEnum {
/// returns a FusionResult explaining why fusion is not feasible.
/// NOTE: This function is not feature complete and should only be used in
/// testing.
-/// TODO(andydavis) Update comments when this function is fully implemented.
+/// TODO: Update comments when this function is fully implemented.
FusionResult canFuseLoops(AffineForOp srcForOp, AffineForOp dstForOp,
unsigned dstLoopDepth,
ComputationSliceState *srcSlice);
/// Collect loop nest statistics (eg. loop trip count and operation count)
/// in 'stats' for loop nest rooted at 'forOp'. Returns true on success,
/// returns false otherwise.
-// TODO(andydavis) Consider moving this to LoopUtils.
+// TODO: Consider moving this to LoopUtils.
bool getLoopNestStats(AffineForOp forOp, LoopNestStats *stats);
/// Computes the total cost of the loop nest rooted at 'forOp' using 'stats'.
/// Currently, the total cost is computed by counting the total operation
/// instance count (i.e. total number of operations in the loop body * loop
/// trip count) for the entire loop nest.
-// TODO(andydavis) Improve this cost model.
+// TODO: Improve this cost model.
int64_t getComputeCost(AffineForOp forOp, LoopNestStats &stats);
/// Computes and returns in 'computeCost', the total compute cost of fusing the
/// (i.e. total number of operations in the loop body * loop trip count) for
/// the entire loop nest.
/// Returns true on success, failure otherwise (e.g. non-constant trip counts).
-// TODO(andydavis) Improve this cost model.
+// TODO: Improve this cost model.
bool getFusionComputeCost(AffineForOp srcForOp, LoopNestStats &srcStats,
AffineForOp dstForOp, LoopNestStats &dstStats,
ComputationSliceState *slice, int64_t *computeCost);
// d1, d2) -> (d0 - d1, d2), and %ii will be the extra operand. Without any
// extra operands, note that 'indexRemap' would just be applied to existing
// indices (%i, %j).
-// TODO(bondhugula): allow extraIndices to be added at any position.
+// TODO: allow extraIndices to be added at any position.
LogicalResult replaceAllMemRefUsesWith(Value oldMemRef, Value newMemRef,
ArrayRef<Value> extraIndices = {},
AffineMap indexRemap = AffineMap(),
/// Returns the sequence of AffineApplyOp Operations operation in
/// 'affineApplyOps', which are reachable via a search starting from 'operands',
/// and ending at operands which are not defined by AffineApplyOps.
-// TODO(andydavis) Add a method to AffineApplyOp which forward substitutes
-// the AffineApplyOp into any user AffineApplyOps.
+// TODO: Add a method to AffineApplyOp which forward substitutes the
+// AffineApplyOp into any user AffineApplyOps.
void mlir::getReachableAffineApplyOps(
ArrayRef<Value> operands, SmallVectorImpl<Operation *> &affineApplyOps) {
struct State {
// the loop IVs of the forOps appearing in that order. Any symbols founds in
// the bound operands are added as symbols in the system. Returns failure for
// the yet unimplemented cases.
-// TODO(andydavis,bondhugula) Handle non-unit steps through local variables or
-// stride information in FlatAffineConstraints. (For eg., by using iv - lb %
-// step = 0 and/or by introducing a method in FlatAffineConstraints
-// setExprStride(ArrayRef<int64_t> expr, int64_t stride)
+// TODO: Handle non-unit steps through local variables or stride information in
+// FlatAffineConstraints. (For eg., by using iv - lb % step = 0 and/or by
+// introducing a method in FlatAffineConstraints setExprStride(ArrayRef<int64_t>
+// expr, int64_t stride)
LogicalResult mlir::getIndexSet(MutableArrayRef<AffineForOp> forOps,
FlatAffineConstraints *domain) {
SmallVector<Value, 4> indices;
// potentially involving any Function symbols. The dimensional identifiers in
// 'indexSet' correspond to the loops surrounding 'op' from outermost to
// innermost.
-// TODO(andydavis) Add support to handle IfInsts surrounding 'op'.
+// TODO: Add support to handle IfInsts surrounding 'op'.
static LogicalResult getInstIndexSet(Operation *op,
FlatAffineConstraints *indexSet) {
- // TODO(andydavis) Extend this to gather enclosing IfInsts and consider
+ // TODO: Extend this to gather enclosing IfInsts and consider
// factoring it out into a utility function.
SmallVector<AffineForOp, 4> loops;
getLoopIVs(*op, &loops);
// an Value in multiple maps are provided (i.e. getSrcDimOrSymPos) to handle
// the common case of resolving positions for all access function operands.
//
-// TODO(andydavis) Generalize this: could take a template parameter for
-// the number of maps (3 in the current case), and lookups could take indices
-// of maps to check. So getSrcDimOrSymPos would be "getPos(value, {0, 2})".
+// TODO: Generalize this: could take a template parameter for the number of maps
+// (3 in the current case), and lookups could take indices of maps to check. So
+// getSrcDimOrSymPos would be "getPos(value, {0, 2})".
class ValuePositionMap {
public:
void addSrcValue(Value value) {
// 0 0 0 -1 0 0 0 50 >= 0
//
//
-// TODO(andydavis) Support AffineExprs mod/floordiv/ceildiv.
+// TODO: Support AffineExprs mod/floordiv/ceildiv.
DependenceResult mlir::checkMemrefAccessDependence(
const MemRefAccess &srcAccess, const MemRefAccess &dstAccess,
unsigned loopDepth, FlatAffineConstraints *dependenceConstraints,
FlatAffineConstraints dependenceConstraints;
SmallVector<DependenceComponent, 2> depComps;
- // TODO(andydavis,bondhugula) Explore whether it would be profitable
- // to pre-compute and store deps instead of repeatedly checking.
+ // TODO: Explore whether it would be profitable to pre-compute and store
+ // deps instead of repeatedly checking.
DependenceResult result = checkMemrefAccessDependence(
srcAccess, dstAccess, d, &dependenceConstraints, &depComps);
if (hasDependence(result))
// We are going to be removing one or more identifiers from the range.
assert(idStart < numIds && "invalid idStart position");
- // TODO(andydavis) Make 'removeIdRange' a lambda called from here.
+ // TODO: Make 'removeIdRange' a lambda called from here.
// Remove eliminated identifiers from equalities.
shiftColumnsToLeft(this, idStart, idLimit, /*isEq=*/true);
if (c == pos)
continue;
// The coefficient of the quotient should be +/-divisor.
- // TODO(bondhugula): could be extended to detect an affine function for
- // the quotient (i.e., the coeff could be a non-zero multiple of divisor).
+ // TODO: could be extended to detect an affine function for the quotient
+ // (i.e., the coeff could be a non-zero multiple of divisor).
int64_t v = cst.atEq(r, c) * cst.atEq(r, pos);
if (v == divisor || v == -divisor) {
seenQuotient++;
quotientSign = v > 0 ? 1 : -1;
}
// The coefficient of the dividend should be +/-1.
- // TODO(bondhugula): could be extended to detect an affine function of
- // the other identifiers as the dividend.
+ // TODO: could be extended to detect an affine function of the other
+ // identifiers as the dividend.
else if (v == -1 || v == 1) {
seenDividend++;
dividendPos = c;
}
// Expression can't be constructed as it depends on a yet unknown
// identifier.
- // TODO(mlir-team): Visit/compute the identifiers in an order so that
- // this doesn't happen. More complex but much more efficient.
+ // TODO: Visit/compute the identifiers in an order so that this doesn't
+ // happen. More complex but much more efficient.
if (c < f)
continue;
// Successfully detected the floordiv.
lbMap = AffineMap::get(numMapDims, numMapSymbols, expr);
ubMap = AffineMap::get(numMapDims, numMapSymbols, expr + 1);
} else {
- // TODO(bondhugula): Whenever there are local identifiers in the
- // dependence constraints, we'll conservatively over-approximate, since we
- // don't always explicitly compute them above (in the while loop).
+ // TODO: Whenever there are local identifiers in the dependence
+ // constraints, we'll conservatively over-approximate, since we don't
+ // always explicitly compute them above (in the while loop).
if (getNumLocalIds() == 0) {
// Work on a copy so that we don't update this constraint system.
if (!tmpClone) {
// If the above fails, we'll just use the constant lower bound and the
// constant upper bound (if they exist) as the slice bounds.
- // TODO(b/126426796): being conservative for the moment in cases that
+ // TODO: being conservative for the moment in cases that
// lead to multiple bounds - until getConstDifference in LoopFusion.cpp is
// fixed (b/126426796).
if (!lbMap || lbMap.getNumResults() > 1) {
}
inequalities.resize(numReservedCols * pos);
- // TODO(bondhugula): consider doing this for equalities as well, but probably
- // not worth the savings.
+ // TODO: consider doing this for equalities as well, but probably not worth
+ // the savings.
}
void FlatAffineConstraints::clearAndCopyFrom(
/// holes/splinters: j = 2
///
/// darkShadow = false, isResultIntegerExact = nullptr are default values.
-// TODO(bondhugula): a slight modification to yield dark shadow version of FM
-// (tightened), which can prove the existence of a solution if there is one.
+// TODO: a slight modification to yield dark shadow version of FM (tightened),
+// which can prove the existence of a solution if there is one.
void FlatAffineConstraints::FourierMotzkinEliminate(
unsigned pos, bool darkShadow, bool *isResultIntegerExact) {
LLVM_DEBUG(llvm::dbgs() << "FM input (eliminate pos " << pos << "):\n");
}
if (r == getNumInequalities()) {
// If it doesn't appear, just remove the column and return.
- // TODO(andydavis,bondhugula): refactor removeColumns to use it from here.
+ // TODO: refactor removeColumns to use it from here.
removeId(pos);
LLVM_DEBUG(llvm::dbgs() << "FM output:\n");
LLVM_DEBUG(dump());
// coefficient in the canonical form as the view taken here is that of the
// term being moved to the other size of '>='.
int64_t ubCoeff = -atIneq(ubPos, pos);
- // TODO(bondhugula): refactor this loop to avoid all branches inside.
+ // TODO: refactor this loop to avoid all branches inside.
for (unsigned l = 0, e = getNumCols(); l < e; l++) {
if (l == pos)
continue;
for (unsigned d = 0, e = getNumDimIds(); d < e; ++d) {
auto extent = getConstantBoundOnDimSize(d, &lb, &lbFloorDivisor, &ub);
if (!extent.hasValue())
- // TODO(bondhugula): symbolic extents when necessary.
- // TODO(bondhugula): handle union if a dimension is unbounded.
+ // TODO: symbolic extents when necessary.
+ // TODO: handle union if a dimension is unbounded.
return failure();
auto otherExtent = otherAligned.getConstantBoundOnDimSize(
d, &otherLb, &otherLbFloorDivisor, &otherUb);
if (!otherExtent.hasValue() || lbFloorDivisor != otherLbFloorDivisor)
- // TODO(bondhugula): symbolic extents when necessary.
+ // TODO: symbolic extents when necessary.
return failure();
assert(lbFloorDivisor > 0 && "divisor always expected to be positive");
append(commonCst);
removeTrivialRedundancy();
- // TODO(mlir-team): copy over pure symbolic constraints from this and 'other'
- // over to the union (since the above are just the union along dimensions); we
- // shouldn't be discarding any other constraints on the symbols.
+ // TODO: copy over pure symbolic constraints from this and 'other' over to the
+ // union (since the above are just the union along dimensions); we shouldn't
+ // be discarding any other constraints on the symbols.
return success();
}
/// Returns false if the MemRef has a non-identity layoutMap or more than 1
/// layoutMap. This is conservative.
///
-// TODO(ntv): check strides.
+// TODO: check strides.
template <typename LoadOrStoreOp>
static bool isContiguousAccess(Value iv, LoadOrStoreOp memoryOp,
int *memRefDim) {
auto memRefType = memoryOp.getMemRefType();
auto layoutMap = memRefType.getAffineMaps();
- // TODO(ntv): remove dependence on Builder once we support non-identity
- // layout map.
+ // TODO: remove dependence on Builder once we support non-identity layout map.
Builder b(memoryOp.getContext());
if (layoutMap.size() >= 2 ||
(layoutMap.size() == 1 &&
auto store = dyn_cast<AffineStoreOp>(op);
// Only scalar types are considered vectorizable, all load/store must be
// vectorizable for a loop to qualify as vectorizable.
- // TODO(ntv): ponder whether we want to be more general here.
+ // TODO: ponder whether we want to be more general here.
bool vector = load ? isVectorElement(load) : isVectorElement(store);
if (vector) {
return false;
/// Checks whether SSA dominance would be violated if a for op's body
/// operations are shifted by the specified shifts. This method checks if a
/// 'def' and all its uses have the same shift factor.
-// TODO(mlir-team): extend this to check for memory-based dependence violation
-// when we have the support.
+// TODO: extend this to check for memory-based dependence violation when we have
+// the support.
bool mlir::isOpwiseShiftValid(AffineForOp forOp, ArrayRef<uint64_t> shifts) {
auto *forBody = forOp.getBody();
assert(shifts.size() == forBody->getOperations().size());
/// 3. if all is good, recursively matches the nested patterns;
/// 4. if all nested match then the single operation matches too and is
/// appended to the list of matches;
-/// 5. TODO(ntv) Optionally applies actions (lambda), in which case we will
-/// want to traverse in post-order DFS to avoid invalidating iterators.
+/// 5. TODO: Optionally applies actions (lambda), in which case we will want
+/// to traverse in post-order DFS to avoid invalidating iterators.
void NestedPattern::matchOne(Operation *op,
SmallVectorImpl<NestedMatch> *matches) {
if (skip == op) {
// region: {memref = %A, write = false, {%i <= m0 <= %i + 7} }
// The last field is a 2-d FlatAffineConstraints symbolic in %i.
//
-// TODO(bondhugula): extend this to any other memref dereferencing ops
+// TODO: extend this to any other memref dereferencing ops
// (dma_start, dma_wait).
LogicalResult MemRefRegion::compute(Operation *op, unsigned loopDepth,
ComputationSliceState *sliceState,
if (auto loop = getForInductionVarOwner(operand)) {
// Note that cst can now have more dimensions than accessMap if the
// bounds expressions involve outer loops or other symbols.
- // TODO(bondhugula): rewrite this to use getInstIndexSet; this way
+ // TODO: rewrite this to use getInstIndexSet; this way
// conditionals will be handled when the latter supports it.
if (failed(cst.addAffineForOpDomain(loop)))
return failure();
// Add upper/lower bounds for each memref dimension with static size
// to guard against potential over-approximation from projection.
- // TODO(andydavis) Support dynamic memref dimensions.
+ // TODO: Support dynamic memref dimensions.
if (addMemRefDimBounds) {
auto memRefType = memref.getType().cast<MemRefType>();
for (unsigned r = 0; r < rank; r++) {
/// Returns the size of memref data in bytes if it's statically shaped, None
/// otherwise. If the element of the memref has vector type, takes into account
/// size of the vector as well.
-// TODO(mlir-team): improve/complete this when we have target data.
+// TODO: improve/complete this when we have target data.
Optional<uint64_t> mlir::getMemRefSizeInBytes(MemRefType memRefType) {
if (!memRefType.hasStaticShape())
return None;
// feasible. If it is, there is at least one point out of bounds.
SmallVector<int64_t, 4> ineq(rank + 1, 0);
int64_t dimSize = loadOrStoreOp.getMemRefType().getDimSize(r);
- // TODO(bondhugula): handle dynamic dim sizes.
+ // TODO: handle dynamic dim sizes.
if (dimSize == -1)
continue;
}
// Returns the innermost common loop depth for the set of operations in 'ops'.
-// TODO(andydavis) Move this to LoopUtils.
+// TODO: Move this to LoopUtils.
static unsigned
getInnermostCommonLoopDepth(ArrayRef<Operation *> ops,
SmallVectorImpl<AffineForOp> &surroundingLoops) {
}
// Clear all sliced loop bounds beginning at the first sequential loop, or
// first loop with a slice fusion barrier attribute..
- // TODO(andydavis, bondhugula) Use MemRef read/write regions instead of
+ // TODO: Use MemRef read/write regions instead of
// using 'kSliceFusionBarrierAttrName'.
auto getSliceLoop = [&](unsigned i) {
return isBackwardSlice ? srcLoopIVs[i] : dstLoopIVs[i];
/// updates the slice loop bounds with any non-null bound maps specified in
/// 'sliceState', and inserts this slice into the loop nest surrounding
/// 'dstOpInst' at loop depth 'dstLoopDepth'.
-// TODO(andydavis,bondhugula): extend the slicing utility to compute slices that
+// TODO: extend the slicing utility to compute slices that
// aren't necessarily a one-to-one relation b/w the source and destination. The
// relation between the source and destination could be many-to-many in general.
-// TODO(andydavis,bondhugula): the slice computation is incorrect in the cases
+// TODO: the slice computation is incorrect in the cases
// where the dependence from the source to the destination does not cover the
// entire destination index set. Subtract out the dependent destination
// iterations from destination index set and check for emptiness --- this is one
// Find the op block positions of 'srcOpInst' within 'srcLoopIVs'.
SmallVector<unsigned, 4> positions;
- // TODO(andydavis): This code is incorrect since srcLoopIVs can be 0-d.
+ // TODO: This code is incorrect since srcLoopIVs can be 0-d.
findInstPosition(srcOpInst, srcLoopIVs[0].getOperation()->getBlock(),
&positions);
return op.src().getType().template cast<VectorType>().getElementType();
}
-// TODO(ntv, zinenko): Code is currently copy-pasted and adapted from the code
+// TODO: Code is currently copy-pasted and adapted from the code
// 1-1 LLVM conversion. It would better if it were properly exposed in core and
// reusable.
/// Basic lowering implementation for one-to-one rewriting from AVX512 Ops to
}
namespace {
-// TODO(ntv): Patterns are too verbose due to the fact that we have 1 op (e.g.
+// TODO: Patterns are too verbose due to the fact that we have 1 op (e.g.
// MaskRndScaleOp) and different possible target ops. It would be better to take
// a Functor so that all these conversions become 1-liners.
struct MaskRndScaleOpPS512Conversion : public ConvertToLLVMPattern {
LLVM::LogOp, LLVM::Log10Op, LLVM::Log2Op>();
target.addIllegalOp<FuncOp>();
target.addLegalDialect<NVVM::NVVMDialect>();
- // TODO(csigg): Remove once we support replacing non-root ops.
+ // TODO: Remove once we support replacing non-root ops.
target.addLegalOp<gpu::YieldOp, gpu::GPUModuleOp, gpu::ModuleEndOp>();
if (failed(applyPartialConversion(m, target, patterns)))
signalPassFailure();
LLVM::LogOp, LLVM::Log10Op, LLVM::Log2Op>();
target.addIllegalOp<FuncOp>();
target.addLegalDialect<ROCDL::ROCDLDialect>();
- // TODO(whchung): Remove once we support replacing non-root ops.
+ // TODO: Remove once we support replacing non-root ops.
target.addLegalOp<gpu::YieldOp, gpu::GPUModuleOp, gpu::ModuleEndOp>();
if (failed(applyPartialConversion(m, target, patterns)))
signalPassFailure();
BaseViewConversionHelper desc(
typeConverter.convertType(sliceOp.getShapedType()));
- // TODO(ntv): extract sizes and emit asserts.
+ // TODO: extract sizes and emit asserts.
SmallVector<Value, 4> strides(memRefType.getRank());
for (int i = 0, e = memRefType.getRank(); i < e; ++i)
strides[i] = baseDesc.stride(i);
if (genericOp.indexing_maps().getValue().size() != 2)
return llvm::None;
- // TODO(nicolasvasilache): create utility functions for these checks in Linalg
+ // TODO: create utility functions for these checks in Linalg
// and use them.
auto inputMap = genericOp.indexing_maps().getValue()[0].cast<AffineMapAttr>();
auto outputMap =
[](const APInt &size) { return !size.isOneValue(); }))
return failure();
- // TODO(antiagainst): Query the target environment to make sure the current
+ // TODO: Query the target environment to make sure the current
// workload fits in a local workgroup.
Value convertedInput = operands[0], convertedOutput = operands[1];
// Get the invocation ID.
Value x = getLocalInvocationDimSize(genericOp, /*dim=*/0, loc, &rewriter);
- // TODO(antiagainst): Load to Workgroup storage class first.
+ // TODO: Load to Workgroup storage class first.
// Get the input element accessed by this invocation.
Value inputElementPtr = spirv::getElementPtr(
if (!libraryCallName)
return failure();
- // TODO(pifon, ntv): Use induction variables values instead of zeros, when
+ // TODO: Use induction variables values instead of zeros, when
// IndexedGenericOp is tiled.
auto zero = rewriter.create<mlir::ConstantOp>(
op.getLoc(), rewriter.getIntegerAttr(rewriter.getIndexType(), 0));
/// Populate the given list with patterns that convert from Linalg to Standard.
void mlir::populateLinalgToStandardConversionPatterns(
OwningRewritePatternList &patterns, MLIRContext *ctx) {
- // TODO(ntv) ConvOp conversion needs to export a descriptor with relevant
+ // TODO: ConvOp conversion needs to export a descriptor with relevant
// attribute values such as kernel striding and dilation.
// clang-format off
patterns.insert<
ParallelOp parallelOp, gpu::LaunchOp launchOp,
BlockAndValueMapping &cloningMap, SmallVectorImpl<Operation *> &worklist,
DenseMap<gpu::Processor, Value> &bounds, PatternRewriter &rewriter) {
- // TODO(herhut): Verify that this is a valid GPU mapping.
+ // TODO: Verify that this is a valid GPU mapping.
// processor ids: 0-2 block [x/y/z], 3-5 -> thread [x/y/z], 6-> sequential
ArrayAttr mapping =
parallelOp.getAttrOfType<ArrayAttr>(gpu::getMappingAttrName());
- // TODO(herhut): Support reductions.
+ // TODO: Support reductions.
if (!mapping || parallelOp.getNumResults() != 0)
return failure();
loc, annotation.map().getValue().compose(lowerAndStep),
ValueRange{operand, step, lowerBound});
// If there was also a bound, insert that, too.
- // TODO(herhut): Check that we do not assign bounds twice.
+ // TODO: Check that we do not assign bounds twice.
if (annotation.bound().getValue()) {
// We pass as the single operand to the bound-map the number of
// iterations, which is (upperBound - lowerBound) ceilDiv step. To
// 4. a second array containing as many `index`-type integers as the rank of
// the MemRef: the second array represents the "stride" (in tensor abstraction
// sense), i.e. the number of consecutive elements of the underlying buffer.
-// TODO(ntv, zinenko): add assertions for the static cases.
+// TODO: add assertions for the static cases.
//
// template <typename Elem, size_t Rank>
// struct {
srcType.getElementType());
dstElementsAttr = dstElementsAttr.reshape(dstAttrType);
} else {
- // TODO(antiagainst): add support for large vectors.
+ // TODO: add support for large vectors.
return failure();
}
}
}
// Helper that returns a subset of `arrayAttr` as a vector of int64_t.
-// TODO(rriddle): Better support for attribute subtype forwarding + slicing.
+// TODO: Better support for attribute subtype forwarding + slicing.
static SmallVector<int64_t, 4> getI64SubArray(ArrayAttr arrayAttr,
unsigned dropFront = 0,
unsigned dropBack = 0) {
loc, toLLVMTy(vectorCmpType), linearIndices);
// 3. Create offsetVector = [ offset + 0 .. offset + vector_length - 1 ].
- // TODO(ntv, ajcbik): when the leaf transfer rank is k > 1 we need the last
+ // TODO: when the leaf transfer rank is k > 1 we need the last
// `k` dimensions here.
unsigned lastIndex = llvm::size(xferOp.indices()) - 1;
Value offsetIndex = *(xferOp.indices().begin() + lastIndex);
// output of any shaped and dimensioned vector. Due to full unrolling,
// this approach is less suited for very large vectors though.
//
- // TODO(ajcbik): rely solely on libc in future? something else?
+ // TODO: rely solely on libc in future? something else?
//
LogicalResult
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
scope(std::make_unique<ScopedContext>(rewriter, loc)), xferOp(xferOp),
op(xferOp.getOperation()) {
vectorType = xferOp.getVectorType();
- // TODO(ntv, ajcbik): when we go to k > 1-D vectors adapt minorRank.
+ // TODO: when we go to k > 1-D vectors adapt minorRank.
minorRank = 1;
majorRank = vectorType.getRank() - minorRank;
leadingRank = xferOp.getMemRefType().getRank() - (majorRank + minorRank);
/// in the presence of data-parallel only operations, we generate code that
/// writes the same value multiple time on the edge locations.
///
-/// TODO(ntv): implement alternatives to clipping.
-/// TODO(ntv): support non-data-parallel operations.
+/// TODO: implement alternatives to clipping.
+/// TODO: support non-data-parallel operations.
/// Performs the rewrite.
template <>
/// See `Important notes about clipping and full-tiles only abstraction` in the
/// description of `readClipped` above.
///
-/// TODO(ntv): implement alternatives to clipping.
-/// TODO(ntv): support non-data-parallel operations.
+/// TODO: implement alternatives to clipping.
+/// TODO: support non-data-parallel operations.
template <>
LogicalResult VectorTransferRewriter<TransferWriteOp>::matchAndRewrite(
Operation *op, PatternRewriter &rewriter) const {
// of allowing mathematical AffineMap composition.
// Returns an AffineMap where symbols that come from an AffineApplyOp have been
// rewritten as dims and are ordered after the original dims.
-// TODO(andydavis,ntv): This promotion makes AffineMap lose track of which
+// TODO: This promotion makes AffineMap lose track of which
// symbols are represented as dims. This loss is static but can still be
// recovered dynamically (with `isValidSymbol`). Still this is annoying for the
// semi-affine map case. A dynamic canonicalization of all dims that are valid
/// composition via symbols, which is ambiguous mathematically. This corner case
/// is handled by locally rewriting such symbols that come from AffineApplyOp
/// into dims and composing through dims.
-/// TODO(andydavis, ntv): Composition via symbols comes at a significant code
+/// TODO: Composition via symbols comes at a significant code
/// complexity. Alternatively we should investigate whether we want to
/// explicitly disallow symbols coming from affine.apply and instead force the
/// user to compose symbols beforehand. The annoyances may be small (i.e. 1 or 2
LLVM_DEBUG(auxiliaryMap.print(dbgs() << "\nWith map: "));
LLVM_DEBUG(map.compose(auxiliaryMap).print(dbgs() << "\nResult: "));
- // TODO(andydavis,ntv): Disabling simplification results in major speed gains.
+ // TODO: Disabling simplification results in major speed gains.
// Another option is to cache the results as it is expected a lot of redundant
// work is performed in practice.
affineMap = simplifyAffineMap(map.compose(auxiliaryMap));
// AffineDmaStartOp
//===----------------------------------------------------------------------===//
-// TODO(b/133776335) Check that map operands are loop IVs or symbols.
+// TODO: Check that map operands are loop IVs or symbols.
void AffineDmaStartOp::build(OpBuilder &builder, OperationState &result,
Value srcMemRef, AffineMap srcMap,
ValueRange srcIndices, Value destMemRef,
// AffineDmaWaitOp
//===----------------------------------------------------------------------===//
-// TODO(b/133776335) Check that map operands are loop IVs or symbols.
+// TODO: Check that map operands are loop IVs or symbols.
void AffineDmaWaitOp::build(OpBuilder &builder, OperationState &result,
Value tagMemRef, AffineMap tagMap,
ValueRange tagIndices, Value numElements) {
"expected affine min or max op");
// Fold the affine map.
- // TODO(andydavis, ntv) Fold more cases:
+ // TODO: Fold more cases:
// min(some_affine, some_affine + constant, ...), etc.
SmallVector<int64_t, 2> results;
auto foldedMap = op.map().partialConstantFold(operands, &results);
return false;
}
auto expr = const_cast<AffineValueMap *>(this)->getAffineMap().getResult(idx);
- // TODO(ntv): this is better implemented on a flattened representation.
+ // TODO: this is better implemented on a flattened representation.
// At least for now it is conservative.
return expr.isFunctionOfDim(index);
}
/// inner levels if necessary to determine at what depth copies need to be
/// placed so that the allocated buffers fit within the memory capacity
/// provided.
-// TODO(bondhugula): We currently can't generate copies correctly when stores
+// TODO: We currently can't generate copies correctly when stores
// are strided. Check for strided stores.
struct AffineDataCopyGeneration
: public AffineDataCopyGenerationBase<AffineDataCopyGeneration> {
/// Generates copies for memref's living in 'slowMemorySpace' into newly created
/// buffers in 'fastMemorySpace', and replaces memory operations to the former
/// by the latter. Only load op's handled for now.
-/// TODO(bondhugula): extend this to store op's.
+/// TODO: extend this to store op's.
std::unique_ptr<OperationPass<FuncOp>> mlir::createAffineDataCopyGenerationPass(
unsigned slowMemorySpace, unsigned fastMemorySpace, unsigned tagMemorySpace,
int minDmaTransferSize, uint64_t fastMemCapacityBytes) {
// operations excluding AffineForOp's) are always assumed to not exhaust
// memory. As a result, this approach is conservative in some cases at the
// moment; we do a check later and report an error with location info.
- // TODO(bondhugula): An 'affine.if' operation is being treated similar to an
+ // TODO: An 'affine.if' operation is being treated similar to an
// operation. 'affine.if''s could have 'affine.for's in them;
// treat them separately.
namespace {
/// Loop invariant code motion (LICM) pass.
-/// TODO(asabne) : The pass is missing zero-trip tests.
-/// TODO(asabne) : Check for the presence of side effects before hoisting.
+/// TODO: The pass is missing zero-trip tests.
+/// TODO: Check for the presence of side effects before hoisting.
/// TODO: This code should be removed once the new LICM pass can handle its
/// uses.
struct LoopInvariantCodeMotion
SmallPtrSetImpl<Operation *> &opsToHoist);
static bool isMemRefDereferencingOp(Operation &op) {
- // TODO(asabne): Support DMA Ops.
+ // TODO: Support DMA Ops.
return isa<AffineLoadOp, AffineStoreOp>(op);
}
// 'affine.if'.
return false;
} else if (isa<AffineDmaStartOp, AffineDmaWaitOp>(op)) {
- // TODO(asabne): Support DMA ops.
+ // TODO: Support DMA ops.
return false;
} else if (!isa<ConstantOp>(op)) {
if (isMemRefDereferencingOp(op)) {
/// Tiles the specified band of perfectly nested loops creating tile-space loops
/// and intra-tile loops. A band is a contiguous set of loops.
-// TODO(bondhugula): handle non hyper-rectangular spaces.
+// TODO: handle non hyper-rectangular spaces.
LogicalResult
mlir::tilePerfectlyNested(MutableArrayRef<AffineForOp> input,
ArrayRef<unsigned> tileSizes,
// based on a simple model that looks at the memory footprint and determines
// tile sizes assuming identity accesses / 1:1 tile size proportional footprint
// along each of the dimensions being tiled.
-// TODO(mlir-team): evolve this model. Tile size determination is a large area
+// TODO: evolve this model. Tile size determination is a large area
// to play with in general.
void LoopTiling::getTileSizes(ArrayRef<AffineForOp> band,
SmallVectorImpl<unsigned> *tileSizes) {
}
// Divide all loops equally in an attempt to reduce footprint.
- // TODO(bondhugula): this is approximate. Ideally, obtain reuse factor /
+ // TODO: this is approximate. Ideally, obtain reuse factor /
// profitability along each dimension and weight tile sizes based on that as
// one possible approach. Or compute a polynomial in tile sizes and solve for
// it.
///
/// Lowering, unrolling, pipelining:
/// ================================
-/// TODO(ntv): point to the proper places.
+/// TODO: point to the proper places.
///
/// Algorithm:
/// ==========
/// reduction, vectorizable, ...) as well as b. all contiguous load/store
/// operations along a specified minor dimension (not necessarily the
/// fastest varying) ;
-/// 2. analyzing those patterns for profitability (TODO(ntv): and
+/// 2. analyzing those patterns for profitability (TODO: and
/// interference);
/// 3. Then, for each pattern in order:
/// a. applying iterative rewriting of the loop and the load operations in
/// c. traversing the use-def chains starting from the roots and iteratively
/// propagating vectorized values. Scalar values that are encountered
/// during this process must come from outside the scope of the current
-/// pattern (TODO(ntv): enforce this and generalize). Such a scalar value
+/// pattern (TODO: enforce this and generalize). Such a scalar value
/// is vectorized only if it is a constant (into a vector splat). The
/// non-constant case is not supported for now and results in the pattern
/// failing to vectorize;
vectorSizes = virtualVectorSize;
}
-/////// TODO(ntv): Hoist to a VectorizationStrategy.cpp when appropriate.
+/////// TODO: Hoist to a VectorizationStrategy.cpp when appropriate.
/////////
namespace {
/// load/store MemRefs, this creates a generic vectorization strategy that works
/// for any loop in a hierarchy (outermost, innermost or intermediate).
///
-/// TODO(ntv): In the future we should additionally increase the power of the
+/// TODO: In the future we should additionally increase the power of the
/// profitability analysis along 3 directions:
/// 1. account for loop extents (both static and parametric + annotations);
/// 2. account for data layout permutations;
return success();
}
-///// end TODO(ntv): Hoist to a VectorizationStrategy.cpp when appropriate /////
+///// end TODO: Hoist to a VectorizationStrategy.cpp when appropriate /////
namespace {
}
}
-////// TODO(ntv): Hoist to a VectorizationMaterialize.cpp when appropriate. ////
+////// TODO: Hoist to a VectorizationMaterialize.cpp when appropriate. ////
/// Handles the vectorization of load and store MLIR operations.
///
auto memRefType = memoryOp.getMemRef().getType().template cast<MemRefType>();
auto elementType = memRefType.getElementType();
- // TODO(ntv): ponder whether we want to further vectorize a vector value.
+ // TODO: ponder whether we want to further vectorize a vector value.
assert(VectorType::isValidElementType(elementType) &&
"Not a valid vector element type");
auto vectorType = VectorType::get(state->strategy->vectorSizes, elementType);
auto *opInst = memoryOp.getOperation();
// For now, vector.transfers must be aligned, operate only on indices with an
// identity subset of AffineMap and do not change layout.
- // TODO(ntv): increase the expressiveness power of vector.transfer operations
+ // TODO: increase the expressiveness power of vector.transfer operations
// as needed by various targets.
if (auto load = dyn_cast<AffineLoadOp>(opInst)) {
OpBuilder b(opInst);
}
return success();
}
-/// end TODO(ntv): Hoist to a VectorizationMaterialize.cpp when appropriate. ///
+/// end TODO: Hoist to a VectorizationMaterialize.cpp when appropriate. ///
/// Coarsens the loops bounds and transforms all remaining load and store
/// operations into the appropriate vector.transfer.
/// Returns an operand that has been vectorized to match `state`'s strategy if
/// vectorization is possible with the above logic. Returns nullptr otherwise.
///
-/// TODO(ntv): handle more complex cases.
+/// TODO: handle more complex cases.
static Value vectorizeOperand(Value operand, Operation *op,
VectorizationState *state) {
LLVM_DEBUG(dbgs() << "\n[early-vect]vectorize operand: " << operand);
LLVM_DEBUG(dbgs() << "-> delayed replacement by: " << res);
return res;
}
- // 2. TODO(ntv): broadcast needed.
+ // 2. TODO: broadcast needed.
if (operand.getType().isa<VectorType>()) {
LLVM_DEBUG(dbgs() << "-> non-vectorizable");
return nullptr;
/// particular operation vectorizes. For now we implement the case distinction
/// here.
/// Returns a vectorized form of an operation or nullptr if vectorization fails.
-// TODO(ntv): consider adding a trait to Op to describe how it gets vectorized.
+// TODO: consider adding a trait to Op to describe how it gets vectorized.
// Maybe some Ops are not vectorizable or require some tricky logic, we cannot
// do one-off logic here; ideally it would be TableGen'd.
static Operation *vectorizeOneOperation(Operation *opInst,
}
// Create a clone of the op with the proper operands and return types.
- // TODO(ntv): The following assumes there is always an op with a fixed
+ // TODO: The following assumes there is always an op with a fixed
// name that works both in scalar mode and vector mode.
- // TODO(ntv): Is it worth considering an Operation.clone operation which
+ // TODO: Is it worth considering an Operation.clone operation which
// changes the type so we can promote an Operation with less boilerplate?
OpBuilder b(opInst);
OperationState newOp(opInst->getLoc(), opInst->getName().getStringRef(),
// Note: we have to exclude terminals because some of their defs may not be
// nested under the vectorization pattern (e.g. constants defined in an
// encompassing scope).
- // TODO(ntv): Use a backward slice for terminals, avoid special casing and
+ // TODO: Use a backward slice for terminals, avoid special casing and
// merge implementations.
for (auto *op : state->roots) {
getForwardSlice(op, &worklist, [state](Operation *op) {
// pattern matching, from profitability analysis, from application.
// As a consequence we must check that each root pattern is still
// vectorizable. If a pattern is not vectorizable anymore, we just skip it.
- // TODO(ntv): implement a non-greedy profitability analysis that keeps only
+ // TODO: implement a non-greedy profitability analysis that keeps only
// non-intersecting patterns.
if (!isVectorizableLoopBody(loop, vectorTransferPattern())) {
LLVM_DEBUG(dbgs() << "\n[early-vect]+++++ loop is not vectorizable");
// 2. Vectorize operations reached by use-def chains from root except the
// terminals (store operations) that need to be post-processed separately.
- // TODO(ntv): add more as we expand.
+ // TODO: add more as we expand.
if (failed(vectorizeNonTerminals(&state))) {
LLVM_DEBUG(dbgs() << "\n[early-vect]+++++ failed vectorizeNonTerminals");
return guard.failure();
// Note: we have to post-process terminals because some of their defs may not
// be nested under the vectorization pattern (e.g. constants defined in an
// encompassing scope).
- // TODO(ntv): Use a backward slice for terminals, avoid special casing and
+ // TODO: Use a backward slice for terminals, avoid special casing and
// merge implementations.
for (auto *op : state.terminals) {
if (!vectorizeOneOperation(op, &state)) { // nullptr == failure
// This automatically prunes intersecting matches.
for (auto m : matches) {
VectorizationStrategy strategy;
- // TODO(ntv): depending on profitability, elect to reduce the vector size.
+ // TODO: depending on profitability, elect to reduce the vector size.
strategy.vectorSizes.assign(vectorSizes.begin(), vectorSizes.end());
if (failed(analyzeProfitability(m.getMatchedChildren(), 1, patternDepth,
&strategy))) {
}
vectorizeLoopIfProfitable(m.getMatchedOperation(), 0, patternDepth,
&strategy);
- // TODO(ntv): if pattern does not apply, report it; alter the
+ // TODO: if pattern does not apply, report it; alter the
// cost/benefit.
vectorizeRootMatch(m, &strategy);
- // TODO(ntv): some diagnostics if failure to vectorize occurs.
+ // TODO: some diagnostics if failure to vectorize occurs.
}
}
LLVM_DEBUG(dbgs() << "\n");
return launchOp.emitOpError("kernel function is missing the '")
<< GPUDialect::getKernelFuncAttrName() << "' attribute";
- // TODO(ntv,zinenko,herhut): if the kernel function has been converted to
+ // TODO: if the kernel function has been converted to
// the LLVM dialect but the caller hasn't (which happens during the
// separate compilation), do not check type correspondence as it would
// require the verifier to be aware of the LLVM type conversion.
map.map(operand.value(), entryBlock.getArgument(operand.index()));
// Clone the region of the gpu.launch operation into the gpu.func operation.
- // TODO(ravishankarm): If cloneInto can be modified such that if a mapping for
+ // TODO: If cloneInto can be modified such that if a mapping for
// a block exists, that block will be used to clone operations into (at the
// end of the block), instead of creating a new block, this would be much
// cleaner.
/// Computed the hardware id to use for a given mapping level. Will
/// assign x,y and z hardware ids for the first 3 dimensions and use
/// sequential after.
-/// TODO(ravishankarm/herhut) : Make this use x for the inner-most loop that is
+/// TODO: Make this use x for the inner-most loop that is
/// distributed to map to x, the next innermost to y and the next innermost to
/// z.
static gpu::Processor getHardwareIdForMapping(MappingLevel level,
// * if we allow memrefs to cast from/to bare pointers, some users might
// alternatively want metadata that only present in the descriptor.
//
- // TODO(timshen): re-evaluate the memref cast design when it's needed.
+ // TODO: re-evaluate the memref cast design when it's needed.
return op.emitOpError("type must be non-index integer types, float types, "
"or vector of mentioned types.");
};
// NVVMDialect initialization, type parsing, and registration.
//===----------------------------------------------------------------------===//
-// TODO(herhut): This should be the llvm.nvvm dialect once this is supported.
+// TODO: This should be the llvm.nvvm dialect once this is supported.
NVVMDialect::NVVMDialect(MLIRContext *context) : Dialect("nvvm", context) {
addOperations<
#define GET_OP_LIST
// ROCDLDialect initialization, type parsing, and registration.
//===----------------------------------------------------------------------===//
-// TODO(herhut): This should be the llvm.rocdl dialect once this is supported.
+// TODO: This should be the llvm.rocdl dialect once this is supported.
ROCDLDialect::ROCDLDialect(MLIRContext *context) : Dialect("rocdl", context) {
addOperations<
#define GET_OP_LIST
SmallVector<Operation *, 8> res;
// Consider an intermediate interleaved `interim` op, look for any dependence
// to an aliasing view on a src -> op -> dst path.
- // TODO(ntv) we are not considering paths yet, just interleaved positions.
+ // TODO: we are not considering paths yet, just interleaved positions.
for (auto dt : types) {
for (auto dependence : getDependencesFrom(src, dt)) {
auto interimPos = linalgOpPositions.lookup(dependence.dependentOpView.op);
ArrayRef<int> strides,
ArrayRef<int> dilations) {
MLIRContext *ctx = ScopedContext::getContext();
- // TODO(ntv) some template magic to make everything rank-polymorphic.
+ // TODO: some template magic to make everything rank-polymorphic.
assert((dilations.empty() || dilations.size() == 2) && "only 2-D conv atm");
assert((strides.empty() || strides.size() == 2) && "only 2-D conv atm");
Value vI, Value vW, Value vO, int depth_multiplier, ArrayRef<int> strides,
ArrayRef<int> dilations) {
MLIRContext *ctx = ScopedContext::getContext();
- // TODO(ntv) some template magic to make everything rank-polymorphic.
+ // TODO: some template magic to make everything rank-polymorphic.
assert((dilations.empty() || dilations.size() == 2) && "only 2-D conv atm");
assert((strides.empty() || strides.size() == 2) && "only 2-D conv atm");
// proper symbol in the AffineExpr of a stride.
if (ShapedType::isDynamic(sizes[dim + 1]))
return false;
- // TODO(ntv) Refine this by passing the proper nDims and nSymbols so we can
+ // TODO: Refine this by passing the proper nDims and nSymbols so we can
// simplify on the fly and catch more reshapable cases.
if (strides[idx] != strides[idx + 1] * sizes[idx + 1])
return false;
/// Helper functions assert Attribute of the proper type in attr and returns the
/// corresponding vector.
-/// TODO(rridle,ntv) this should be evolved into a generic
+/// TODO: this should be evolved into a generic
/// `getRangeOfType<AffineMap>(ArrayAttr attrs)` that does not copy.
static SmallVector<AffineMap, 4> getAffineMaps(ArrayAttr attrs) {
return llvm::to_vector<8>(llvm::map_range(
if (failed(verifyReshapeLikeTypes(op, expandedType, collapsedType)))
return failure();
auto maps = getAffineMaps(op.reassociation());
- // TODO(ntv): expanding a ? with a non-constant is under-specified. Error
+ // TODO: expanding a ? with a non-constant is under-specified. Error
// out.
RankedTensorType expectedType =
computeTensorReshapeCollapsedType(expandedType, maps);
(void)res;
unsigned rank = memRefType.getRank();
- // TODO(ntv): propagate static size and stride information when available.
+ // TODO: propagate static size and stride information when available.
SmallVector<int64_t, 4> sizes(rank, -1); // -1 encodes dynamic size.
result.addTypes({MemRefType::Builder(memRefType)
.setShape(sizes)
SmallVector<AffineExpr, 4> res;
res.reserve(outputDims.size());
for (unsigned i = 0, e = outputDims.size(); i < e; ++i) {
- // TODO(ntv): add a level of indirection to linalg.generic.
+ // TODO: add a level of indirection to linalg.generic.
auto expr = op.getStride(i) * outputDims[i] +
op.getDilation(i) * windowDims[i] - op.getLowPad(i);
res.push_back(expr);
return ss.str();
}
-// TODO(ntv, rriddle): Consider making all this boilerplate easy to autogenerate
+// TODO: Consider making all this boilerplate easy to autogenerate
// with Tablegen. This seems a desirable property in the context of OpInterfaces
// where a Linalg "named" op **isa** LinalgOp.
LogicalResult ConvOp::fold(ArrayRef<Attribute>,
return false;
}
if (auto convOp = dyn_cast<linalg::ConvOp>(producer.getOperation())) {
- // TODO(ntv): add a level of indirection to linalg.generic.
+ // TODO: add a level of indirection to linalg.generic.
if (convOp.padding())
return false;
}
if (auto convOp = dyn_cast<linalg::ConvOp>(consumer.getOperation())) {
- // TODO(ntv): add a level of indirection to linalg.generic.
+ // TODO: add a level of indirection to linalg.generic.
if (convOp.padding())
return false;
}
linalgOps.push_back(op);
});
- // TODO(pifon, ntv): LinalgDependenceGraph should be able to update itself.
+ // TODO: LinalgDependenceGraph should be able to update itself.
// The current naive and expensive reconstruction of the graph should be
// removed.
for (auto *op : llvm::reverse(linalgOps)) {
SmallVector<Value, 4> indexedValues;
indexedValues.reserve(nInputs + nOutputs);
- // TODO(mravishankar): Avoid the loads if the corresponding argument of the
+ // TODO: Avoid the loads if the corresponding argument of the
// region has no uses.
// 1.a. Emit load from input views.
for (unsigned i = 0; i < nInputs; ++i) {
IndexedValueType(linalgOp.getOutputBuffer(i))(indexing));
}
- // TODO(ntv): When a region inliner exists, use it.
+ // TODO: When a region inliner exists, use it.
// 2. Inline region, currently only works for a single basic block.
// 3. Emit store.
SmallVector<SmallVector<Value, 8>, 8> indexing;
template <typename IndexedValueType>
Value getConvOpInput(ConvOp convOp, StdIndexedValue im,
MutableArrayRef<Value> imIdx) {
- // TODO(ntv): add a level of indirection to linalg.generic.
+ // TODO: add a level of indirection to linalg.generic.
if (!convOp.padding())
return im(imIdx);
for (unsigned i = 0; i < nLoops; ++i)
indexedValues.push_back(allIvs[i]);
- // TODO(mravishankar): Avoid the loads if the corresponding argument of the
+ // TODO: Avoid the loads if the corresponding argument of the
// region has no uses.
// 1.a. Emit load from input views.
for (unsigned i = 0; i < nInputs; ++i) {
IndexedValueType(indexedGenericOp.getOutputBuffer(i))(indexing));
}
- // TODO(ntv): When a region inliner exists, use it.
+ // TODO: When a region inliner exists, use it.
// 2. Inline region, currently only works for a single basic block.
// 3. Emit store.
SmallVector<SmallVector<Value, 8>, 8> indexing;
OwningRewritePatternList patterns;
// Canonicalization and folding patterns applied greedily allow cleaning up
// the emitted IR on the fly.
- // TODO(ntv) fold view and subview ops?
+ // TODO: fold view and subview ops?
insertPatterns<LoopType,
#define GET_OP_LIST
#include "mlir/Dialect/Linalg/IR/LinalgStructuredOps.cpp.inc"
assert(op.hasBufferSemantics() && "expected linalg op with buffer semantics");
if (auto convOp = dyn_cast<linalg::ConvOp>(op.getOperation())) {
- // TODO(ntv): add a level of indirection to linalg.generic.
+ // TODO: add a level of indirection to linalg.generic.
if (convOp.padding())
return {};
}
// }
// }
//
-// TODO(pifon, ntv): Investigate whether mixing implicit and explicit indices
+// TODO: Investigate whether mixing implicit and explicit indices
// does not lead to losing information.
static void transformIndexedGenericOpIndices(
OpBuilder &b, LinalgOp op, SmallVectorImpl<Value> &ivs,
// that refers to an existing function symbol. The `fun` function call will be
// inserted in the loop body in that case.
//
- // TODO(pifon): Add support for `linalg.indexed_generic` with `fun` attribute.
+ // TODO: Add support for `linalg.indexed_generic` with `fun` attribute.
auto ®ion = indexedGenericOp.region();
if (region.empty()) {
indexedGenericOp.emitOpError("expected a region");
if (!genericOp || !::isRowMajorMatmul(genericOp))
return failure();
- // TODO(ntv): non-identity layout.
+ // TODO: non-identity layout.
auto isStaticMemRefWithIdentityLayout = [](Value v) {
auto m = v.getType().dyn_cast<MemRefType>();
if (!m || !m.hasStaticShape() || !m.getAffineMaps().empty())
/// of the innermost loop is populated by `bodyBuilderFn` that accepts a range
/// of induction variables for all loops. `ivStorage` is used to store the
/// partial list of induction variables.
-// TODO(zinenko,ntv): this function can be made iterative instead. However, it
+// TODO: this function can be made iterative instead. However, it
// will have at most as many recursive calls as nested loops, which rarely
// exceeds 10.
static void
// point is derived from the shifted range, and the scale isn't changed. As
// a consequence some values, which are supposed in the original [rmin, rmax]
// range will be outside the shifted range and be clamped during quantization.
-// TODO(fengliuai): we should nudge the scale as well, but that requires the
+// TODO: we should nudge the scale as well, but that requires the
// fake quant op used in the training to use the nudged scale as well.
static void getNudgedScaleAndZeroPoint(int64_t qmin, int64_t qmax, double rmin,
double rmax, double &scale,
quantizedElementType.dyn_cast<UniformQuantizedPerAxisType>()) {
UniformQuantizedPerAxisValueConverter converter(uniformQuantizedPerAxis);
auto converted = converter.convert(realValue);
- // TODO(fengliuai): why we need this outConvertedType? remove it?
+ // TODO: why we need this outConvertedType? remove it?
if (converted) {
outConvertedType = converted.getType();
}
if (auto attr = realValue.dyn_cast<DenseFPElementsAttr>()) {
return convert(attr);
}
- // TODO(fengliuai): handles sparse elements attribute
+ // TODO: handles sparse elements attribute
return nullptr;
}
SDBM SDBM::get(ArrayRef<SDBMExpr> inequalities, ArrayRef<SDBMExpr> equalities) {
SDBM result;
- // TODO(zinenko): consider detecting equalities in the list of inequalities.
+ // TODO: consider detecting equalities in the list of inequalities.
// This is potentially expensive and requires to
// - create a list of negated inequalities (may allocate under lock);
// - perform a pairwise comparison of direct and negated inequalities;
// value is positive, the set defined by SDBM is trivially empty. We store
// this value anyway and continue processing to maintain the correspondence
// between the matrix form and the list-of-SDBMExpr form.
- // TODO(zinenko): we may want to reconsider this once we have canonicalization
+ // TODO: we may want to reconsider this once we have canonicalization
// or simplification in place
auto updateMatrix = [](SDBM &sdbm, const SDBMBuilderResult &r) {
for (auto positivePos : r.positivePos) {
if (pattern.match(expr)) {
if (SDBMExpr converted = visit(x.matched())) {
if (auto varConverted = converted.dyn_cast<SDBMTermExpr>())
- // TODO(ntv): return varConverted.stripe(C.getConstantValue());
+ // TODO: return varConverted.stripe(C.getConstantValue());
return SDBMStripeExpr::get(
varConverted,
SDBMConstantExpr::get(dialect,
/// operation registered to the current dialect.
bool isLegalToInline(Operation *op, Region *dest,
BlockAndValueMapping &) const final {
- // TODO(antiagainst): Enable inlining structured control flows with return.
+ // TODO: Enable inlining structured control flows with return.
if ((isa<spirv::SelectionOp, spirv::LoopOp>(op)) &&
containsReturn(op->getRegion(0)))
return false;
- // TODO(antiagainst): we need to filter OpKill here to avoid inlining it to
+ // TODO: we need to filter OpKill here to avoid inlining it to
// a loop continue construct:
// https://github.com/KhronosGroup/SPIRV-Headers/issues/86
// However OpKill is fragment shader specific and we don't support it yet.
return CooperativeMatrixNVType::get(elementTy, scope, dims[0], dims[1]);
}
-// TODO(ravishankarm) : Reorder methods to be utilities first and parse*Type
+// TODO: Reorder methods to be utilities first and parse*Type
// methods in alphabetical order
//
// storage-class ::= `UniformConstant`
template <>
Optional<Type> parseAndVerify<Type>(SPIRVDialect const &dialect,
DialectAsmParser &parser) {
- // TODO(ravishankarm): Further verify that the element type can be sampled
+ // TODO: Further verify that the element type can be sampled
auto ty = parseAndVerifyType(dialect, parser);
if (!ty)
return llvm::None;
StringRef symbol = attribute.first.strref();
Attribute attr = attribute.second;
- // TODO(antiagainst): figure out a way to generate the description from the
+ // TODO: figure out a way to generate the description from the
// StructAttr definition.
if (symbol == spirv::getEntryPointABIAttrName()) {
if (!attr.isa<spirv::EntryPointABIAttr>())
Type SPIRVTypeConverter::getIndexType(MLIRContext *context) {
// Convert to 32-bit integers for now. Might need a way to control this in
// future.
- // TODO(ravishankarm): It is probably better to make it 64-bit integers. To
+ // TODO: It is probably better to make it 64-bit integers. To
// this some support is needed in SPIR-V dialect for Conversion
// instructions. The Vulkan spec requires the builtins like
// GlobalInvocationID, etc. to be 32-bit (unsigned) integers which should be
/// behind the number assignments; we try to follow NVVM conventions and largely
/// give common storage classes a smaller number. The hope is use symbolic
/// memory space representation eventually after memref supports it.
-// TODO(antiagainst): swap Generic and StorageBuffer assignment to be more akin
+// TODO: swap Generic and StorageBuffer assignment to be more akin
// to NVVM.
#define STORAGE_SPACE_MAP_LIST(MAP_FN) \
MAP_FN(spirv::StorageClass::Generic, 1) \
#undef STORAGE_SPACE_MAP_LIST
-// TODO(ravishankarm): This is a utility function that should probably be
+// TODO: This is a utility function that should probably be
// exposed by the SPIR-V dialect. Keeping it local till the use case arises.
static Optional<int64_t> getTypeNumBytes(Type t) {
if (t.isa<spirv::ScalarType>()) {
// Otherwise we need to adjust the type, which really means adjusting the
// bitwidth given this is a scalar type.
- // TODO(antiagainst): We are unconditionally converting the bitwidth here,
+ // TODO: We are unconditionally converting the bitwidth here,
// this might be okay for non-interface types (i.e., types used in
// Private/Function storage classes), but not for interface types (i.e.,
// types used in StorageBuffer/Uniform/PushConstant/etc. storage classes).
convertVectorType(const spirv::TargetEnv &targetEnv, VectorType type,
Optional<spirv::StorageClass> storageClass = {}) {
if (!spirv::CompositeType::isValid(type)) {
- // TODO(antiagainst): One-element vector types can be translated into scalar
+ // TODO: One-element vector types can be translated into scalar
// types. Vector types with more than four elements can be translated into
// array types.
LLVM_DEBUG(llvm::dbgs()
/// manipulate, like what we do for vectors.
static Optional<Type> convertTensorType(const spirv::TargetEnv &targetEnv,
TensorType type) {
- // TODO(ravishankarm) : Handle dynamic shapes.
+ // TODO: Handle dynamic shapes.
if (!type.hasStaticShape()) {
LLVM_DEBUG(llvm::dbgs()
<< type << " illegal: dynamic shape unimplemented\n");
// adopted in the SPIR-V dialect (i.e., IntegerType, FloatType, VectorType)
// were tried before.
//
- // TODO(antiagainst): this assumes that the SPIR-V types are valid to use in
+ // TODO: this assumes that the SPIR-V types are valid to use in
// the given target environment, which should be the case if the whole
// pipeline is driven by the same target environment. Still, we probably still
// want to validate and convert to be safe.
FuncOpConversion::matchAndRewrite(FuncOp funcOp, ArrayRef<Value> operands,
ConversionPatternRewriter &rewriter) const {
auto fnType = funcOp.getType();
- // TODO(antiagainst): support converting functions with one result.
+ // TODO: support converting functions with one result.
if (fnType.getNumResults())
return failure();
using namespace mlir;
-// TODO(antiagainst): generate these strings using ODS.
+// TODO: generate these strings using ODS.
static constexpr const char kAlignmentAttrName[] = "alignment";
static constexpr const char kBranchWeightAttrName[] = "branch_weights";
static constexpr const char kCallee[] = "callee";
// ODS already checks ptr is spirv::PointerType. Just check that the pointee
// type of the pointer and the type of the value are the same
//
- // TODO(ravishankarm): Check that the value type satisfies restrictions of
+ // TODO: Check that the value type satisfies restrictions of
// SPIR-V OpLoad/OpStore operations
if (val.getType() !=
ptr.getType().cast<spirv::PointerType>().getPointeeType()) {
Operation *sizeOp = groupOp->getOperand(1).getDefiningOp();
int32_t clusterSize = 0;
- // TODO(antiagainst): support specialization constant here.
+ // TODO: support specialization constant here.
if (failed(extractValueFromConstOp(sizeOp, clusterSize)))
return groupOp->emitOpError(
"cluster size operand must come from a constant op");
return nullptr;
}
- // TODO(denis0x0D): this should be relaxed to allow
+ // TODO: this should be relaxed to allow
// integer literals of other bitwidths.
if (failed(extractValueFromConstOp(op, index))) {
emitError(baseLoc,
"as the op result type, but found ")
<< pointeeType << " vs " << atomOp.getType();
- // TODO(antiagainst): Unequal cannot be set to Release or Acquire and Release.
+ // TODO: Unequal cannot be set to Release or Acquire and Release.
// In addition, Unequal cannot be set to a stronger memory-order then Equal.
return success();
if (type.getKind() >= Type::FIRST_SPIRV_TYPE &&
type.getKind() <= spirv::TypeKind::LAST_SPIRV_TYPE) {
- // TODO(antiagainst): support constant struct
+ // TODO: support constant struct
return type.isa<spirv::ArrayType>();
}
return WalkResult::advance();
});
- // TODO(antiagainst): verify other bits like linkage type.
+ // TODO: verify other bits like linkage type.
return failure(walkResult.wasInterrupted());
}
}
static ParseResult parseLoopOp(OpAsmParser &parser, OperationState &state) {
- // TODO(antiagainst): support loop control properly
+ // TODO: support loop control properly
Builder builder = parser.getBuilder();
state.addAttribute("loop_control",
builder.getI32IntegerAttr(
if (funcOp.isExternal())
return op.emitError("'spv.module' cannot contain external functions");
- // TODO(antiagainst): move this check to spv.func.
+ // TODO: move this check to spv.func.
for (auto &block : funcOp)
for (auto &op : block) {
if (op.getDialect() != dialect)
static ParseResult parseSelectionOp(OpAsmParser &parser,
OperationState &state) {
- // TODO(antiagainst): support selection control properly
+ // TODO: support selection control properly
Builder builder = parser.getBuilder();
state.addAttribute("selection_control",
builder.getI32IntegerAttr(
if (block->hasNoPredecessors())
return success();
- // TODO(antiagainst): further verification needs to analyze reachability from
+ // TODO: further verification needs to analyze reachability from
// the entry block.
return success();
"constant or spv.globalVariable op");
}
- // TODO(antiagainst): generate these strings using ODS.
+ // TODO: generate these strings using ODS.
auto *op = varOp.getOperation();
auto descriptorSetName = llvm::convertToSnakeFromCamelCase(
stringifyDecoration(spirv::Decoration::DescriptorSet));
llvm::SetVector<Capability, SmallVector<Capability, 0>> allCaps(
directCaps.begin(), directCaps.end());
- // TODO(antiagainst): This is insufficient; find a better way to handle this
+ // TODO: This is insufficient; find a better way to handle this
// (e.g., using static lists) if this turns out to be a bottleneck.
for (unsigned i = 0; i < allCaps.size(); ++i)
for (Capability c : getDirectImpliedCapabilities(allCaps[i]))
/// higher-order bits. So this deserializer uses that to get instruction
/// boundary and parse instructions and build a SPIR-V ModuleOp gradually.
///
-// TODO(antiagainst): clean up created ops on errors
+// TODO: clean up created ops on errors
class Deserializer {
public:
/// Creates a deserializer for the given SPIR-V `binary` module.
/// MLIRContext to create SPIR-V ModuleOp into.
MLIRContext *context;
- // TODO(antiagainst): create Location subclass for binary blob
+ // TODO: create Location subclass for binary blob
Location unknownLoc;
/// The SPIR-V ModuleOp.
<< majorVersion;
}
- // TODO(antiagainst): generator number, bound, schema
+ // TODO: generator number, bound, schema
curOffset = spirv::kHeaderWordCount;
return success();
}
}
LogicalResult Deserializer::processDecoration(ArrayRef<uint32_t> words) {
- // TODO : This function should also be auto-generated. For now, since only a
+ // TODO: This function should also be auto-generated. For now, since only a
// few decorations are processed/handled in a meaningful manner, going with a
// manual implementation.
if (words.size() < 2) {
return emitError(unknownLoc, "unknown Function Control: ") << operands[2];
}
if (functionControl.getValue() != spirv::FunctionControl::None) {
- /// TODO : Handle different function controls
+ /// TODO: Handle different function controls
return emitError(unknownLoc, "unhandled Function Control: '")
<< spirv::stringifyFunctionControl(functionControl.getValue())
<< "'";
}
unsigned count = 0;
- // TODO(antiagainst): The count can also come frome a specialization constant.
+ // TODO: The count can also come frome a specialization constant.
auto countInfo = getConstant(operands[2]);
if (!countInfo) {
return emitError(unknownLoc, "OpTypeArray count <id> ")
}
typeMap[operands[0]] =
spirv::StructType::get(memberTypes, offsetInfo, memberDecorationsInfo);
- // TODO(ravishankarm): Update StructType to have member name as attribute as
+ // TODO: Update StructType to have member name as attribute as
// well.
return success();
}
// merge block so that the newly created LoopOp will be inserted there.
OpBuilder builder(&mergeBlock->front());
- // TODO(antiagainst): handle loop control properly
+ // TODO: handle loop control properly
auto loopOp = builder.create<spirv::LoopOp>(location);
loopOp.addEntryAndMergeBlock();
// selection/loop. If so, they will be recorded within blockMergeInfo.
// We need to update the pointers there to the newly remapped ones so we can
// continue structurizing them later.
- // TODO(antiagainst): The asserts in the following assumes input SPIR-V blob
+ // TODO: The asserts in the following assumes input SPIR-V blob
// forms correctly nested selection/loop constructs. We should relax this
// and support error cases better.
auto it = blockMergeInfo.find(block);
/// iterates the DenseElementsAttr to construct the constant array, and
/// returns the result <id> associated with it. Returns 0 if failed. Note
/// that the size of `index` must match the rank.
- /// TODO(hanchung): Consider to enhance splat elements cases. For splat cases,
+ /// TODO: Consider to enhance splat elements cases. For splat cases,
/// we don't need to loop over all elements, especially when the splat value
/// is zero. We can use OpConstantNull when the value is zero.
uint32_t prepareDenseElementsConstant(Location loc, Type constType,
if (failed(module.verify()))
return failure();
- // TODO(antiagainst): handle the other sections
+ // TODO: handle the other sections
processCapability();
processExtension();
processMemoryModel();
operands.push_back(resTypeID);
auto funcID = getOrCreateFunctionID(op.getName());
operands.push_back(funcID);
- // TODO : Support other function control options.
+ // TODO: Support other function control options.
operands.push_back(static_cast<uint32_t>(spirv::FunctionControl::None));
operands.push_back(fnTypeID);
encodeInstructionInto(functionHeader, spirv::Opcode::OpFunction, operands);
return success();
}
- // TODO(ravishankarm) : Handle other types.
+ // TODO: Handle other types.
return emitError(loc, "unhandled type in serialization: ") << type;
}
return resultID;
}
-// TODO(hanchung): Turn the below function into iterative function, instead of
+// TODO: Turn the below function into iterative function, instead of
// recursive function.
uint32_t
Serializer::prepareDenseElementsConstant(Location loc, Type constType,
auto emitSelectionMerge = [&]() {
emitDebugLine(functionBody, loc);
lastProcessedWasMergeInst = true;
- // TODO(antiagainst): properly support selection control here
+ // TODO: properly support selection control here
encodeInstructionInto(
functionBody, spirv::Opcode::OpSelectionMerge,
{mergeID, static_cast<uint32_t>(spirv::SelectionControl::None)});
auto emitLoopMerge = [&]() {
emitDebugLine(functionBody, loc);
lastProcessedWasMergeInst = true;
- // TODO(antiagainst): properly support loop control here
+ // TODO: properly support loop control here
encodeInstructionInto(
functionBody, spirv::Opcode::OpLoopMerge,
{mergeID, continueID, static_cast<uint32_t>(spirv::LoopControl::None)});
}
llvm::SetVector<Operation *> interfaceVarSet;
- // TODO(ravishankarm) : This should in reality traverse the entry function
+ // TODO: This should in reality traverse the entry function
// call graph and collect all the interfaces. For now, just traverse the
// instructions in this function.
funcOp.walk([&](spirv::AddressOfOp addressOfOp) {
auto var =
module.lookupSymbol<spirv::GlobalVariableOp>(addressOfOp.variable());
- // TODO(antiagainst): Per SPIR-V spec: "Before version 1.4, the interface’s
+ // TODO: Per SPIR-V spec: "Before version 1.4, the interface’s
// storage classes are limited to the Input and Output storage classes.
// Starting with version 1.4, the interface’s storage classes are all
// storage classes used in declaring all global variables referenced by the
ConversionPatternRewriter &rewriter) const {
if (!funcOp.getAttrOfType<spirv::EntryPointABIAttr>(
spirv::getEntryPointABIAttrName())) {
- // TODO(ravishankarm) : Non-entry point functions are not handled.
+ // TODO: Non-entry point functions are not handled.
return failure();
}
TypeConverter::SignatureConversion signatureConverter(
auto abiInfo = funcOp.getArgAttrOfType<spirv::InterfaceVarABIAttr>(
argType.index(), attrName);
if (!abiInfo) {
- // TODO(ravishankarm) : For non-entry point functions, it should be legal
+ // TODO: For non-entry point functions, it should be legal
// to pass around scalar/vector values and return a scalar/vector. For now
// non-entry point functions are not handled in this ABI lowering and will
// produce an error.
rewriter.create<spirv::AddressOfOp>(funcOp.getLoc(), var);
// Check if the arg is a scalar or vector type. In that case, the value
// needs to be loaded into registers.
- // TODO(ravishankarm) : This is loading value of the scalar into registers
+ // TODO: This is loading value of the scalar into registers
// at the start of the function. It is probably better to do the load just
// before the use. There might be multiple loads and currently there is no
// easy way to replace all uses with a sequence of operations.
if (walkResult.wasInterrupted())
return signalPassFailure();
- // TODO(antiagainst): verify that the deduced version is consistent with
+ // TODO: verify that the deduced version is consistent with
// SPIR-V ops' maximal version requirements.
auto triple = spirv::VerCapExtAttr::get(
auto lhs = operands.front().dyn_cast_or_null<FloatAttr>();
auto rhs = operands.back().dyn_cast_or_null<FloatAttr>();
- // TODO(gcmn) We could actually do some intelligent things if we know only one
+ // TODO: We could actually do some intelligent things if we know only one
// of the operands, but it's inf or nan.
if (!lhs || !rhs)
return {};
Type resultType;
auto loc = parser.getCurrentLocation();
DictionaryAttr dictAttr;
- // TODO(andydavis, ntv) Unify linalg op attribute parsing.
+ // TODO: Unify linalg op attribute parsing.
if (parser.parseAttribute(dictAttr, "_", result.attributes) ||
parser.parseOperand(lhsInfo) || parser.parseComma() ||
parser.parseOperand(rhsInfo) || parser.parseComma() ||
}
static void print(OpAsmPrinter &p, ContractionOp op) {
- // TODO(andydavis, ntv) Unify printing code with linalg ops.
+ // TODO: Unify printing code with linalg ops.
auto attrNames = op.getTraitAttrNames();
llvm::StringSet<> traitAttrsSet;
traitAttrsSet.insert(attrNames.begin(), attrNames.end());
TupleType tupleType, ArrayRef<int64_t> sizes,
ArrayRef<int64_t> strides) {
// Check for non-unit strides.
- // TODO(b/144845578) Support non-1 strides.
+ // TODO: Support non-1 strides.
if (llvm::any_of(strides, [](int64_t s) { return s != 1; }))
return op->emitError("requires unit strides");
// Check that 'vectorType' rank matches rank of tuple element vectors.
result.addAttribute(getStridesAttrName(), stridesAttr);
}
-// TODO(ntv) Should be moved to Tablegen Confined attributes.
+// TODO: Should be moved to Tablegen Confined attributes.
template <typename OpType>
static LogicalResult isIntegerArrayAttrSmallerThanShape(OpType op,
ArrayAttr arrayAttr,
if (memrefVecEltRank > resultVecRank)
return op->emitOpError(
"requires memref vector element and vector result ranks to match.");
- // TODO(b/146516564) Move this to isSuffix in Vector/Utils.h.
+ // TODO: Move this to isSuffix in Vector/Utils.h.
unsigned rankOffset = resultVecRank - memrefVecEltRank;
auto memrefVecEltShape = memrefVectorElementType.getShape();
auto resultVecShape = vectorType.getShape();
}
// Helper method to possibly drop a dimension in a load.
-// TODO(ajcbik): use a reshaping vector load (and share lowering code)
+// TODO
static Value reshapeLoad(Location loc, Value val, VectorType type,
int64_t index, int64_t pos,
PatternRewriter &rewriter) {
}
// Helper method to possibly drop a dimension in a store.
-// TODO(ajcbik): use a reshaping vector store (and share lowering code)
+// TODO
static Value reshapeStore(Location loc, Value val, Value result,
VectorType type, int64_t index, int64_t pos,
PatternRewriter &rewriter) {
// Returns a tuple type with vector element types for each resulting slice
// of 'vectorType' unrolled by 'sizes' and 'strides'.
-// TODO(andydavis) Move this to a utility function and share it with
+// TODO: Move this to a utility function and share it with
// Extract/InsertSlicesOp verification.
static TupleType generateExtractSlicesOpResultType(VectorType vectorType,
ArrayRef<int64_t> sizes,
// Compute slice offsets.
SmallVector<int64_t, 4> sliceOffsets(state.unrolledShape.size());
getMappedElements(indexMap, offsets, sliceOffsets);
- // TODO(b/144845578) Support non-1 strides.
+ // TODO: Support non-1 strides.
SmallVector<int64_t, 4> sliceStrides(state.unrolledShape.size(), 1);
// Compute linear index of 'sliceOffsets' w.r.t 'state.basis'.
int64_t sliceLinearIndex =
// insertslice
// |
-// TODO(andydavis) Add the following canonicalization/simplification patterns:
+// TODO: Add the following canonicalization/simplification patterns:
// *) Add pattern which matches InsertStridedSlice -> StridedSlice and forwards
// InsertStridedSlice operand to StridedSlice.
// *) Add pattern which matches SourceOp -> StridedSlice -> UserOp which checks
// operation, and leave the duplicate StridedSlice ops with no users
// (removable with DCE).
-// TODO(andydavis) Generalize this to support structured ops beyond
+// TODO: Generalize this to support structured ops beyond
// vector ContractionOp, and merge it with 'unrollSingleResultVectorOp'
static Value unrollSingleResultStructuredOp(Operation *op,
ArrayRef<int64_t> iterationBounds,
vectors.push_back({contractionOp.getRHSVectorMaskType(),
vectors[1].indexMap, accOperandIndex + 2, false});
}
- // TODO(andydavis) Use linalg style 'args_in'/'args_out' to partition
+ // TODO: Use linalg style 'args_in'/'args_out' to partition
// 'vectors' instead of 'resultIndex'.
resultIndex = accOperandIndex;
}
LogicalResult matchAndRewrite(vector::TransferReadOp xferReadOp,
PatternRewriter &rewriter) const override {
- // TODO(andydavis, ntv) Support splitting TransferReadOp with non-identity
+ // TODO: Support splitting TransferReadOp with non-identity
// permutation maps. Repurpose code from MaterializeVectors transformation.
if (!isIdentitySuffix(xferReadOp.permutation_map()))
return failure();
LogicalResult matchAndRewrite(vector::TransferWriteOp xferWriteOp,
PatternRewriter &rewriter) const override {
- // TODO(andydavis, ntv) Support splitting TransferWriteOp with non-identity
+ // TODO: Support splitting TransferWriteOp with non-identity
// permutation maps. Repurpose code from MaterializeVectors transformation.
if (!isIdentitySuffix(xferWriteOp.permutation_map()))
return failure();
/// the vector.contract op is a row-major matrix multiply.
LogicalResult
ContractionOpToMatmulOpLowering::match(vector::ContractionOp op) const {
- // TODO(ajcbik): implement masks
+ // TODO: implement masks
if (llvm::size(op.masks()) != 0)
return failure();
/// otherwise supports any layout permutation of the matrix-multiply.
LogicalResult
ContractionOpToOuterProductOpLowering ::match(vector::ContractionOp op) const {
- // TODO(ajcbik): implement masks
+ // TODO: implement masks
if (llvm::size(op.masks()) != 0)
return failure();
///
/// This only kicks in when VectorTransformsOptions is set to AXPY.
//
-// TODO (ajcbik): this is very similar, but not quite the same as
-// the outerproduct lowering above; merge the two?
+// TODO: this is very similar, but not quite the same as the outerproduct
+// lowering above; merge the two?
LogicalResult
ContractionOpToAXPYLowering::match(vector::ContractionOp op) const {
- // TODO(ajcbik): implement masks
+ // TODO: implement masks
if (llvm::size(op.masks()) != 0)
return failure();
/// This only kicks in when either VectorTransformsOptions is set
/// to DOT or when other contraction patterns fail.
//
-// TODO(ajcbik): break down into transpose/reshape/cast ops
+// TODO: break down into transpose/reshape/cast ops
// when they become available to avoid code dup
-// TODO(ajcbik): investigate lowering order impact on performance
+// TODO: investigate lowering order impact on performance
LogicalResult
ContractionOpLowering::matchAndRewrite(vector::ContractionOp op,
PatternRewriter &rewriter) const {
- // TODO(ajcbik): implement masks.
+ // TODO: implement masks.
if (llvm::size(op.masks()) != 0)
return failure();
- // TODO(thomasraoux): support mixed mode contract lowering.
+ // TODO: support mixed mode contract lowering.
if (op.getLhsType().getElementType() !=
getElementTypeOrSelf(op.getAccType()) ||
op.getRhsType().getElementType() != getElementTypeOrSelf(op.getAccType()))
return failure();
- // TODO(ntv, ajcbik): implement benefits, cost models.
+ // TODO: implement benefits, cost models.
MLIRContext *ctx = op.getContext();
ContractionOpToMatmulOpLowering pat1(vectorTransformsOptions, ctx);
if (succeeded(pat1.match(op)))
}
// Lower one parallel dimension.
-// TODO(ajcbik): consider reusing existing contract unrolling
+// TODO: consider reusing existing contract unrolling
Value ContractionOpLowering::lowerParallel(vector::ContractionOp op,
int64_t lhsIndex, int64_t rhsIndex,
PatternRewriter &rewriter) const {
} // namespace mlir
-// TODO(andydavis) Add pattern to rewrite ExtractSlices(ConstantMaskOp).
-// TODO(andydavis) Add this as DRR pattern.
+// TODO: Add pattern to rewrite ExtractSlices(ConstantMaskOp).
+// TODO: Add this as DRR pattern.
void mlir::vector::populateVectorToVectorTransformationPatterns(
OwningRewritePatternList &patterns, MLIRContext *context) {
// clang-format off
/// Implementation detail that walks up the parents and records the ones with
/// the specified type.
-/// TODO(ntv): could also be implemented as a collect parents followed by a
+/// TODO: could also be implemented as a collect parents followed by a
/// filter and made available outside this file.
template <typename T>
static SetVector<Operation *> getParentsOfType(Operation *op) {
// The ops that *may* lower a super-vector only do so if the super-vector to
// sub-vector ratio exists. The ops that *must* lower a super-vector are
// explicitly checked for this property.
- /// TODO(ntv): there should be a single function for all ops to do this so we
+ /// TODO: there should be a single function for all ops to do this so we
/// do not have to special case. Maybe a trait, or just a method, unclear atm.
bool mustDivide = false;
(void)mustDivide;
// Clone module in a new LLVMContext since translateModuleToLLVMIR buries
// ownership too deeply.
- // TODO(zinenko): Reevaluate model of ownership of LLVMContext in LLVMDialect.
+ // TODO: Reevaluate model of ownership of LLVMContext in LLVMDialect.
std::unique_ptr<Module> deserModule =
LLVM::cloneModuleIntoNewContext(ctx.get(), llvmModule.get());
auto dataLayout = deserModule->getDataLayout();
auto rhsConst = operandExprStack.back()[getConstantIndex()];
operandExprStack.pop_back();
auto &lhs = operandExprStack.back();
- // TODO(bondhugula): handle modulo by zero case when this issue is fixed
+ // TODO: handle modulo by zero case when this issue is fixed
// at the other places in the IR.
assert(rhsConst > 0 && "RHS constant has to be positive");
// This is a pure affine expr; the RHS is a positive constant.
int64_t rhsConst = operandExprStack.back()[getConstantIndex()];
- // TODO(bondhugula): handle division by zero at the same time the issue is
+ // TODO: handle division by zero at the same time the issue is
// fixed at other places.
assert(rhsConst > 0 && "RHS constant has to be positive");
operandExprStack.pop_back();
/// Simplify the affine expression by flattening it and reconstructing it.
AffineExpr mlir::simplifyAffineExpr(AffineExpr expr, unsigned numDims,
unsigned numSymbols) {
- // TODO(bondhugula): only pure affine for now. The simplification here can
+ // TODO: only pure affine for now. The simplification here can
// be extended to semi-affine maps in the future.
if (!expr.isPureAffine())
return expr;
if (results[idx].isMultipleOf(factor))
return true;
- // TODO(bondhugula): use simplifyAffineExpr and FlatAffineConstraints to
+ // TODO: use simplifyAffineExpr and FlatAffineConstraints to
// complete this (for a more powerful analysis).
return false;
}
// be pure for the simplification implemented.
void MutableAffineMap::simplify() {
// Simplify each of the results if possible.
- // TODO(ntv): functional-style map
+ // TODO: functional-style map
for (unsigned i = 0, e = getNumResults(); i < e; i++) {
results[i] = simplifyAffineExpr(getResult(i), numDims, numSymbols);
}
// Utility to generate a function to register a symbol alias.
static bool canRegisterAlias(StringRef name, llvm::StringSet<> &usedAliases) {
assert(!name.empty() && "expected alias name to be non-empty");
- // TODO(riverriddle) Assert that the provided alias name can be lexed as
+ // TODO: Assert that the provided alias name can be lexed as
// an identifier.
// Check that the alias doesn't contain a '.' character and the name is not
/// Record a reference to the given type.
void AliasState::recordTypeReference(Type ty) { usedTypes.insert(ty); }
-// TODO Support visiting other types/operations when implemented.
+// TODO: Support visiting other types/operations when implemented.
void AliasState::visitType(Type type) {
recordTypeReference(type);
}
void Value::printAsOperand(raw_ostream &os, AsmState &state) {
- // TODO(riverriddle) This doesn't necessarily capture all potential cases.
+ // TODO: This doesn't necessarily capture all potential cases.
// Currently, region arguments can be shadowed when printing the main
// operation. If the IR hasn't been printed, this will produce the old SSA
// name and not the shadowed name.
/// Construct a new storage instance.
static OpaqueElementsAttributeStorage *
construct(AttributeStorageAllocator &allocator, KeyTy key) {
- // TODO(b/131468830): Provide a way to avoid copying content of large opaque
+ // TODO: Provide a way to avoid copying content of large opaque
// tensors This will likely require a new reference attribute kind.
return new (allocator.allocate<OpaqueElementsAttributeStorage>())
OpaqueElementsAttributeStorage(std::get<0>(key), std::get<1>(key),
if (getContext()->shouldPrintOpOnDiagnostic()) {
// Print out the operation explicitly here so that we can print the generic
// form.
- // TODO(riverriddle) It would be nice if we could instead provide the
+ // TODO: It would be nice if we could instead provide the
// specific printing flags when adding the operation as an argument to the
// diagnostic.
std::string printedOp;
// For now strides are only computed on a single affine map with a single
// result (i.e. the closed subset of linearization maps that are compatible
// with striding semantics).
- // TODO(ntv): support more forms on a per-need basis.
+ // TODO: support more forms on a per-need basis.
if (affineMaps.size() > 1)
return failure();
if (affineMaps.size() == 1 && affineMaps[0].getNumResults() != 1)
/// In practice, a strided memref must be internally non-aliasing. Test
/// against 0 as a proxy.
- /// TODO(ntv) static cases can have more advanced checks.
- /// TODO(ntv) dynamic cases would require a way to compare symbolic
+ /// TODO: static cases can have more advanced checks.
+ /// TODO: dynamic cases would require a way to compare symbolic
/// expressions and would probably need an affine set context propagated
/// everywhere.
if (llvm::any_of(strides, [](AffineExpr e) {
/// Walk all of the operations nested under and including the given operations.
void detail::walkOperations(Operation *op,
function_ref<void(Operation *op)> callback) {
- // TODO(b/140235992) This walk should be iterative over the operations.
+ // TODO: This walk should be iterative over the operations.
for (auto ®ion : op->getRegions())
for (auto &block : region)
// Early increment here in the case where the operation is erased.
WalkResult
detail::walkOperations(Operation *op,
function_ref<WalkResult(Operation *op)> callback) {
- // TODO(b/140235992) This walk should be iterative over the operations.
+ // TODO: This walk should be iterative over the operations.
for (auto ®ion : op->getRegions()) {
for (auto &block : region) {
// Early increment here in the case where the operation is erased.
auto name = getToken().getStringValue();
auto *dialect = builder.getContext()->getRegisteredDialect(name);
- // TODO(shpeisman): Allow for having an unknown dialect on an opaque
+ // TODO: Allow for having an unknown dialect on an opaque
// attribute. Otherwise, it can't be roundtripped without having the dialect
// registered.
if (!dialect)
return success();
}
- // TODO(riverriddle) support hex floating point values.
+ // TODO: support hex floating point values.
return emitError(getCurrentLocation(), "expected floating point literal");
}
llvm::SourceMgr &sourceMgr,
MLIRContext *context) {
if (sourceMgr.getNumBuffers() != 0) {
- // TODO(b/136086478): Extend to support multiple buffers.
+ // TODO: Extend to support multiple buffers.
emitError(mlir::UnknownLoc::get(context),
"only main buffer parsed at the moment");
return nullptr;
}
LogicalResult detail::PassOptions::parseFromString(StringRef options) {
- // TODO(parkers): Handle escaping strings.
+ // TODO: Handle escaping strings.
// NOTE: `options` is modified in place to always refer to the unprocessed
// part of the string.
while (!options.empty()) {
};
llvm::array_pod_sort(orderedOps.begin(), orderedOps.end(), compareOptionArgs);
for (OptionBase *option : orderedOps) {
- // TODO(riverriddle) printOptionInfo assumes a specific indent and will
+ // TODO: printOptionInfo assumes a specific indent and will
// print options with values with incorrect indentation. We should add
// support to llvm::cl::Option for passing in a base indent to use when
// printing.
// Skip over everything until the closing '}' and store as options.
size_t close = text.find('}');
- // TODO(parkers): Handle skipping over quoted sub-strings.
+ // TODO: Handle skipping over quoted sub-strings.
if (close == StringRef::npos) {
return errorHandler(
/*rawLoc=*/text.data() - 1,
void tblgen::OpMethodSignature::writeDefTo(raw_ostream &os,
StringRef namePrefix) const {
// We need to remove the default values for parameters in method definition.
- // TODO(antiagainst): We are using '=' and ',' as delimiters for parameter
+ // TODO: We are using '=' and ',' as delimiters for parameter
// initializers. This is incorrect for initializer list with more than one
// element. Change to a more robust approach.
auto removeParamDefaultValue = [](StringRef params) {
// still point to the original predicate records. While the original
// predicate may be known to be true or false, it is not necessarily the case
// after rewriting.
- // TODO(zinenko,jpienaar): we can support ground truth for rewritten
+ // TODO: we can support ground truth for rewritten
// predicates by either (a) having our own unique'ing of the predicates
// instead of relying on TableGen record pointers or (b) taking ground truth
// values optionally prefixed with a list of substitutions to apply, e.g.
if (!module->walk(interruptIfValidLocation).wasInterrupted())
return;
- // TODO(riverriddle) Several parts of this are incorrect. Different source
+ // TODO: Several parts of this are incorrect. Different source
// languages may interpret different parts of the debug information
// differently. Frontends will also want to pipe in various information, like
// flags. This is fine for now as we only emit line-table information and not
auto *file = translateFile(fileLoc ? fileLoc.getFilename() : "<unknown>");
unsigned line = fileLoc ? fileLoc.getLine() : 0;
- // TODO(riverriddle) This is the bare essentials for now. We will likely end
+ // TODO: This is the bare essentials for now. We will likely end
// up with wrapper metadata around LLVMs metadata in the future, so this
// doesn't need to be smart until then.
llvm::DISubroutineType *type =
locationToLoc;
/// A mapping between filename and llvm debug file.
- /// TODO(riverriddle) Change this to DenseMap<Identifier, ...> when we can
+ /// TODO: Change this to DenseMap<Identifier, ...> when we can
/// access the Identifier filename in FileLineColLoc.
llvm::StringMap<llvm::DIFile *> fileMap;
}
}
-// TODO(mlir-team): implement an iterative version
+// TODO: implement an iterative version
static void topologicalSortImpl(llvm::SetVector<Block *> &blocks, Block *b) {
blocks.insert(b);
for (Block *bb : b->getSuccessors()) {
if (op->getNumRegions() != 0)
return failure();
- // TODO(riverriddle) We currently only eliminate non side-effecting
+ // TODO: We currently only eliminate non side-effecting
// operations.
if (!MemoryEffectOpInterface::hasNoEffect(op))
return failure();
/// PatternRewriter hook for merging a block into another.
void ConversionPatternRewriter::mergeBlocks(Block *source, Block *dest,
ValueRange argValues) {
- // TODO(riverriddle) This requires fixing the implementation of
+ // TODO: This requires fixing the implementation of
// 'replaceUsesOfBlockArgument', which currently isn't undoable.
llvm_unreachable("block merging updates are currently not supported");
}
}
// If the operation isn't legal, try to fold it in-place.
- // TODO(riverriddle) Should we always try to do this, even if the op is
+ // TODO: Should we always try to do this, even if the op is
// already legal?
if (succeeded(legalizeWithFold(op, rewriter))) {
LLVM_DEBUG({
for (Operation &op : *block) {
if (auto call = dyn_cast<CallOpInterface>(op)) {
- // TODO(riverriddle) Support inlining nested call references.
+ // TODO: Support inlining nested call references.
CallInterfaceCallable callable = call.getCallableForCallee();
if (SymbolRefAttr symRef = callable.dyn_cast<SymbolRefAttr>()) {
if (!symRef.isa<FlatSymbolRefAttr>())
/// which fuses loop nests with single-writer/single-reader memref dependences
/// with the goal of improving locality.
-// TODO(andydavis) Support fusion of source loop nests which write to multiple
+// TODO: Support fusion of source loop nests which write to multiple
// memrefs, where each memref can have multiple users (if profitable).
-// TODO(andydavis) Extend this pass to check for fusion preventing dependences,
+// TODO: Extend this pass to check for fusion preventing dependences,
// and add support for more general loop fusion algorithms.
struct LoopFusion : public AffineLoopFusionBase<LoopFusion> {
maximalFusion);
}
-// TODO(b/117228571) Replace when this is modeled through side-effects/op traits
+// TODO: Replace when this is modeled through side-effects/op traits
static bool isMemRefDereferencingOp(Operation &op) {
return isa<AffineReadOpInterface, AffineWriteOpInterface, AffineDmaStartOp,
AffineDmaWaitOp>(op);
// MemRefDependenceGraph is a graph data structure where graph nodes are
// top-level operations in a FuncOp which contain load/store ops, and edges
// are memref dependences between the nodes.
-// TODO(andydavis) Add a more flexible dependence graph representation.
-// TODO(andydavis) Add a depth parameter to dependence graph construction.
+// TODO: Add a more flexible dependence graph representation.
+// TODO: Add a depth parameter to dependence graph construction.
struct MemRefDependenceGraph {
public:
// Node represents a node in the graph. A Node is either an entire loop nest
// Initializes the data dependence graph by walking operations in 'f'.
// Assigns each node in the graph a node id based on program order in 'f'.
-// TODO(andydavis) Add support for taking a Block arg to construct the
+// TODO: Add support for taking a Block arg to construct the
// dependence graph at a different depth.
bool MemRefDependenceGraph::init(FuncOp f) {
DenseMap<Value, SetVector<unsigned>> memrefAccesses;
getNumCommonSurroundingLoops(*srcOpInst, *dstOpInst);
for (unsigned d = 1; d <= numCommonLoops + 1; ++d) {
FlatAffineConstraints dependenceConstraints;
- // TODO(andydavis) Cache dependence analysis results, check cache here.
+ // TODO: Cache dependence analysis results, check cache here.
DependenceResult result = checkMemrefAccessDependence(
srcAccess, dstAccess, d, &dependenceConstraints,
/*dependenceComponents=*/nullptr);
node->op = newRootForOp.getOperation();
}
-// TODO(mlir-team): improve/complete this when we have target data.
+// TODO: improve/complete this when we have target data.
static unsigned getMemRefEltSizeInBytes(MemRefType memRefType) {
auto elementType = memRefType.getElementType();
// Creates and returns a private (single-user) memref for fused loop rooted
// at 'forOp', with (potentially reduced) memref size based on the
// MemRefRegion written to by 'srcStoreOpInst' at depth 'dstLoopDepth'.
-// TODO(bondhugula): consider refactoring the common code from generateDma and
+// TODO: consider refactoring the common code from generateDma and
// this one.
static Value createPrivateMemRef(AffineForOp forOp, Operation *srcStoreOpInst,
unsigned dstLoopDepth,
// Create new private memref for fused loop 'forOp'. 'newShape' is always
// a constant shape.
- // TODO(andydavis) Create/move alloc ops for private memrefs closer to their
+ // TODO: Create/move alloc ops for private memrefs closer to their
// consumer loop nests to reduce their live range. Currently they are added
// at the beginning of the function, because loop nests can be reordered
// during the fusion pass.
// 'srcLiveOutStoreOp', has output edges.
// Returns true if 'dstNode's read/write region to 'memref' is a super set of
// 'srcNode's write region to 'memref' and 'srcId' has only one output edge.
-// TODO(andydavis) Generalize this to handle more live in/out cases.
+// TODO: Generalize this to handle more live in/out cases.
static bool
canFuseSrcWhichWritesToLiveOut(unsigned srcId, unsigned dstId,
AffineWriteOpInterface srcLiveOutStoreOp,
return false;
// Compute MemRefRegion 'dstRegion' for 'dstStore/LoadOpInst' on 'memref'.
- // TODO(andydavis) Compute 'unionboundingbox' of all write regions (one for
+ // TODO: Compute 'unionboundingbox' of all write regions (one for
// each store op in 'dstStoreOps').
SmallVector<Operation *, 2> dstStoreOps;
dstNode->getStoreOpsForMemref(memref, &dstStoreOps);
// Return false if write region is not a superset of 'srcNodes' write
// region to 'memref'.
- // TODO(andydavis) Check the shape and lower bounds here too.
+ // TODO: Check the shape and lower bounds here too.
if (srcNumElements != dstNumElements)
return false;
maybeSliceWriteRegionSizeBytes.getValue();
// If we are fusing for reuse, check that write regions remain the same.
- // TODO(andydavis) Write region check should check sizes and offsets in
+ // TODO: Write region check should check sizes and offsets in
// each dimension, so that we are sure they are covering the same memref
// region. Also, move this out to a isMemRefRegionSuperSet helper function.
if (srcOpInst != srcStoreOpInst &&
llvm::dbgs() << msg.str();
});
- // TODO(b/123247369): This is a placeholder cost model.
+ // TODO: This is a placeholder cost model.
// Among all choices that add an acceptable amount of redundant computation
// (as per computeToleranceThreshold), we will simply pick the one that
// reduces the intermediary size the most.
// takes O(V) time for initialization, and has runtime O(V + E).
//
// This greedy algorithm is not 'maximal' due to the current restriction of
-// fusing along single producer consumer edges, but there is a TODO to fix this.
+// fusing along single producer consumer edges, but there is a TODO: to fix
+// this.
//
-// TODO(andydavis) Experiment with other fusion policies.
+// TODO: Experiment with other fusion policies.
struct GreedyFusion {
public:
// The data dependence graph to traverse during fusion.
// Initializes 'worklist' with nodes from 'mdg'
void init() {
- // TODO(andydavis) Add a priority queue for prioritizing nodes by different
+ // TODO: Add a priority queue for prioritizing nodes by different
// metrics (e.g. arithmetic intensity/flops-to-bytes ratio).
worklist.clear();
worklistSet.clear();
// *) Second pass fuses sibling nodes which share no dependence edges.
// *) Third pass fuses any remaining producer nodes into their users.
void run() {
- // TODO(andydavis) Run this repeatedly until a fixed-point is reached.
+ // TODO: Run this repeatedly until a fixed-point is reached.
fuseProducerConsumerNodes(/*maxSrcUserCount=*/1);
fuseSiblingNodes();
fuseProducerConsumerNodes(
continue;
// Skip if 'srcNode' has more than one live-out store to a
// function-local memref.
- // TODO(andydavis) Support more generic multi-output src loop nests
+ // TODO: Support more generic multi-output src loop nests
// fusion.
auto srcStoreOp = mdg->getUniqueOutgoingStore(srcNode);
if (!srcStoreOp) {
unsigned dstLoopDepthTest = getInnermostCommonLoopDepth(dstOps);
// Check the feasibility of fusing src loop nest into dst loop nest
// at loop depths in range [1, dstLoopDepthTest].
- // TODO(andydavis) Use slice union computation and union of memref
+ // TODO: Use slice union computation and union of memref
// read/write regions to cost model and fusion.
bool canFuse = false;
for (unsigned i = 1; i <= dstLoopDepthTest; ++i) {
memref)
storesForMemref.push_back(storeOpInst);
}
- // TODO(andydavis) Use union of memref write regions to compute
+ // TODO: Use union of memref write regions to compute
// private memref footprint.
auto newMemRef = createPrivateMemRef(
dstAffineForOp, storesForMemref[0], bestDstLoopDepth,
while (findSiblingNodeToFuse(dstNode, &visitedSibNodeIds, &idAndMemref)) {
unsigned sibId = idAndMemref.first;
Value memref = idAndMemref.second;
- // TODO(andydavis) Check that 'sibStoreOpInst' post-dominates all other
+ // TODO: Check that 'sibStoreOpInst' post-dominates all other
// stores to the same memref in 'sibNode' loop nest.
auto *sibNode = mdg->getNode(sibId);
// Compute an operation list insertion point for the fused loop
assert(sibLoadOpInsts.size() == 1);
Operation *sibLoadOpInst = sibLoadOpInsts[0];
assert(!sibNode->stores.empty());
- // TODO(andydavis) Choose the store which postdominates all other stores.
+ // TODO: Choose the store which postdominates all other stores.
auto *sibStoreOpInst = sibNode->stores.back();
// Gather 'dstNode' load ops to 'memref'.
// on 'memref'.
auto canFuseWithSibNode = [&](Node *sibNode, Value memref) {
// Skip if 'outEdge' is not a read-after-write dependence.
- // TODO(andydavis) Remove restrict to single load op restriction.
+ // TODO: Remove restrict to single load op restriction.
if (sibNode->getLoadOpCount(memref) != 1)
return false;
// Skip if there exists a path of dependent edges between
//
// This file implements a pass to forward memref stores to loads, thereby
// potentially getting rid of intermediate memref's entirely.
-// TODO(mlir-team): In the future, similar techniques could be used to eliminate
+// TODO: In the future, similar techniques could be used to eliminate
// dead memref store's and perform more complex forwarding when support for
// SSA scalars live out of 'affine.for'/'affine.if' statements is available.
//===----------------------------------------------------------------------===//
// don't reason about loops that are guaranteed to execute at least once or
// multiple sources to forward from.
//
-// TODO(mlir-team): more forwarding can be done when support for
+// TODO: more forwarding can be done when support for
// loop/conditional live-out SSA values is available.
-// TODO(mlir-team): do general dead store elimination for memref's. This pass
+// TODO: do general dead store elimination for memref's. This pass
// currently only eliminates the stores only if no other loads/uses (other
// than dealloc) remain.
//
// If the memref hasn't been alloc'ed in this function, skip.
Operation *defOp = memref.getDefiningOp();
if (!defOp || !isa<AllocOp>(defOp))
- // TODO(mlir-team): if the memref was returned by a 'call' operation, we
+ // TODO: if the memref was returned by a 'call' operation, we
// could still erase it if the call had no side-effects.
continue;
if (llvm::any_of(memref.getUsers(), [&](Operation *ownerOp) {
// Returns the position of the tag memref operand given a DMA operation.
// Temporary utility: will be replaced when DmaStart/DmaFinish abstract op's are
-// added. TODO(b/117228571)
+// added. TODO
static unsigned getTagMemRefPos(Operation &dmaOp) {
assert((isa<AffineDmaStartOp, AffineDmaWaitOp>(dmaOp)));
if (auto dmaStartOp = dyn_cast<AffineDmaStartOp>(dmaOp)) {
e = startIndices.end();
it != e; ++it, ++wIt) {
// Keep it simple for now, just checking if indices match.
- // TODO(mlir-team): this would in general need to check if there is no
+ // TODO: this would in general need to check if there is no
// intervening write writing to the same tag location, i.e., memory last
// write/data flow analysis. This is however sufficient/powerful enough for
// now since the DMA generation pass or the input for it will always have
continue;
// Only DMAs incoming into higher memory spaces are pipelined for now.
- // TODO(bondhugula): handle outgoing DMA pipelining.
+ // TODO: handle outgoing DMA pipelining.
if (!dmaStartOp.isDestMemorySpaceFaster())
continue;
// Check for dependence with outgoing DMAs. Doing this conservatively.
- // TODO(andydavis,bondhugula): use the dependence analysis to check for
+ // TODO: use the dependence analysis to check for
// dependences between an incoming and outgoing DMA in the same iteration.
auto it = outgoingDmaOps.begin();
for (; it != outgoingDmaOps.end(); ++it) {
// Identify memref's to replace by scanning through all DMA start
// operations. A DMA start operation has two memref's - the one from the
// higher level of memory hierarchy is the one to double buffer.
- // TODO(bondhugula): check whether double-buffering is even necessary.
- // TODO(bondhugula): make this work with different layouts: assuming here that
+ // TODO: check whether double-buffering is even necessary.
+ // TODO: make this work with different layouts: assuming here that
// the dimension we are adding here for the double buffering is the outermost
// dimension.
for (auto &pair : startWaitPairs) {
for (Value operand : operands) {
// If the use count of this operand is now < 2, we re-add the defining
// operation to the worklist.
- // TODO(riverriddle) This is based on the fact that zero use operations
+ // TODO: This is based on the fact that zero use operations
// may be deleted, and that single use values often have more
// canonicalization opportunities.
if (!operand.use_empty() && !operand.hasOneUse())
return nullptr;
}
// Return insertion point in valid range closest to 'opB'.
- // TODO(andydavis) Consider other insertion points in valid range.
+ // TODO: Consider other insertion points in valid range.
return firstDepOpA;
}
// No dependences from 'opA' to operation in range ('opA', 'opB'), return
return !hasIfOp;
}
-// TODO(andydavis) Prevent fusion of loop nests with side-effecting operations.
+// TODO: Prevent fusion of loop nests with side-effecting operations.
FusionResult mlir::canFuseLoops(AffineForOp srcForOp, AffineForOp dstForOp,
unsigned dstLoopDepth,
ComputationSliceState *srcSlice) {
return tripCount * opCount;
}
-// TODO(andydavis,b/126426796): extend this to handle multiple result maps.
+// TODO: extend this to handle multiple result maps.
static Optional<uint64_t> getConstDifference(AffineMap lbMap, AffineMap ubMap) {
assert(lbMap.getNumResults() == 1 && "expected single result bound map");
assert(ubMap.getNumResults() == 1 && "expected single result bound map");
// nest surrounding represented by slice loop bounds in 'slice'.
// Returns true on success, false otherwise (if a non-constant trip count
// was encountered).
-// TODO(andydavis) Make this work with non-unit step loops.
+// TODO: Make this work with non-unit step loops.
static bool buildSliceTripCountMap(
ComputationSliceState *slice,
llvm::SmallDenseMap<Operation *, uint64_t, 8> *tripCountMap) {
auto *insertPointParent = slice->insertPoint->getParentOp();
// The store and loads to this memref will disappear.
- // TODO(andydavis) Add load coalescing to memref data flow opt pass.
+ // TODO: Add load coalescing to memref data flow opt pass.
if (storeLoadFwdGuaranteed) {
// Subtract from operation count the loads/store we expect load/store
// forwarding to remove.
/// Promotes the loop body of a forOp to its containing block if the forOp
/// was known to have a single iteration.
-// TODO(bondhugula): extend this for arbitrary affine bounds.
+// TODO: extend this for arbitrary affine bounds.
LogicalResult mlir::promoteIfSingleIteration(AffineForOp forOp) {
Optional<uint64_t> tripCount = getConstantTripCount(forOp);
if (!tripCount || tripCount.getValue() != 1)
return failure();
// If the trip count is lower than the unroll factor, no unrolled body.
- // TODO(bondhugula): option to specify cleanup loop unrolling.
+ // TODO: option to specify cleanup loop unrolling.
Optional<uint64_t> mayBeConstantTripCount = getConstantTripCount(forOp);
if (mayBeConstantTripCount.hasValue() &&
mayBeConstantTripCount.getValue() < unrollFactor)
: boundsBuilder.create<ConstantIndexOp>(loc, stepUnrolledCst);
} else {
// Dynamic loop bounds computation.
- // TODO(andydavis) Add dynamic asserts for negative lb/ub/step, or
+ // TODO: Add dynamic asserts for negative lb/ub/step, or
// consider using ceilDiv from AffineApplyExpander.
auto lowerBound = forOp.lowerBound();
auto upperBound = forOp.upperBound();
// Loops where both lower and upper bounds are multi-result maps won't be
// unrolled (since the trip can't be expressed as an affine function in
// general).
- // TODO(mlir-team): this may not be common, but we could support the case
+ // TODO: this may not be common, but we could support the case
// where the lower bound is a multi-result map and the ub is a single result
// one.
if (forOp.getLowerBoundMap().getNumResults() != 1)
// Returns the new AffineForOps, nested immediately under `target`.
template <typename ForType, typename SizeType>
static ForType stripmineSink(ForType forOp, SizeType factor, ForType target) {
- // TODO(ntv): Use cheap structural assertions that targets are nested under
+ // TODO: Use cheap structural assertions that targets are nested under
// forOp and that targets are not nested under each other when DominanceInfo
// exposes the capability. It seems overkill to construct a whole function
// dominance tree at this point.
continue;
}
// Skip if op has side effects.
- // TODO(ntv): loads to immutable memory regions are ok.
+ // TODO: loads to immutable memory regions are ok.
if (!MemoryEffectOpInterface::hasNoEffect(&op)) {
status = failure();
continue;
auto intraTile = tile(forOps, tileSizes, forOps.back());
TileLoops tileLoops = std::make_pair(forOps, intraTile);
- // TODO(ntv, zinenko) for now we just ignore the result of band isolation.
+ // TODO: for now we just ignore the result of band isolation.
// In the future, mapping decisions may be impacted by the ability to
// isolate perfectly nested bands.
tryIsolateBands(tileLoops);
// Compute the number of iterations the loop executes: ceildiv(ub - lb, step)
// assuming the step is strictly positive. Update the bounds and the step
// of the loop to go from 0 to the number of iterations, if necessary.
- // TODO(zinenko): introduce support for negative steps or emit dynamic asserts
+ // TODO: introduce support for negative steps or emit dynamic asserts
// on step positivity, whatever gets implemented first.
if (isZeroBased && isStepOne)
return {/*lowerBound=*/lowerBound, /*upperBound=*/upperBound,
// symbolic/variant.
auto it = enclosingFors.rbegin();
for (auto e = enclosingFors.rend(); it != e; ++it) {
- // TODO(bondhugula): also need to be checking this for regions symbols that
+ // TODO: also need to be checking this for regions symbols that
// aren't loop IVs, whether we are within their resp. defs' dominance scope.
if (llvm::is_contained(symbols, it->getInductionVar()))
break;
/// potentially multiple striding levels from outermost to innermost. For an
/// n-dimensional region, there can be at most n-1 levels of striding
/// successively nested.
-// TODO(bondhugula): make this work with non-identity layout maps.
+// TODO: make this work with non-identity layout maps.
static void getMultiLevelStrides(const MemRefRegion ®ion,
ArrayRef<int64_t> bufferShape,
SmallVectorImpl<StrideInfo> *strideInfos) {
SmallVector<StrideInfo, 4> dmaStrideInfos;
getMultiLevelStrides(region, fastBufferShape, &dmaStrideInfos);
- // TODO(bondhugula): use all stride levels once DmaStartOp is extended for
+ // TODO: use all stride levels once DmaStartOp is extended for
// multi-level strides.
if (dmaStrideInfos.size() > 1) {
LLVM_DEBUG(llvm::dbgs() << "Only up to one level of stride supported\n");
// Each memref has a single buffer associated with it irrespective of how
// many load's and store's happen on it.
- // TODO(bondhugula): in the future, when regions don't intersect and satisfy
+ // TODO: in the future, when regions don't intersect and satisfy
// other properties (based on load/store regions), we could consider
// multiple buffers per memref.
/// Return true if this operation dereferences one or more memref's.
// Temporary utility: will be replaced when this is modeled through
-// side-effects/op traits. TODO(b/117228571)
+// side-effects/op traits. TODO
static bool isMemRefDereferencingOp(Operation &op) {
return isa<AffineReadOpInterface, AffineWriteOpInterface, AffineDmaStartOp,
AffineDmaWaitOp>(op);
return success();
if (usePositions.size() > 1) {
- // TODO(mlir-team): extend it for this case when needed (rare).
+ // TODO: extend it for this case when needed (rare).
assert(false && "multiple dereferencing uses in a single op not supported");
return failure();
}
// Create new fully composed AffineMap for new op to be created.
assert(newMapOperands.size() == newMemRefRank);
auto newMap = builder.getMultiDimIdentityMap(newMemRefRank);
- // TODO(b/136262594) Avoid creating/deleting temporary AffineApplyOps here.
+ // TODO: Avoid creating/deleting temporary AffineApplyOps here.
fullyComposeAffineMapAndOperands(&newMap, &newMapOperands);
newMap = simplifyAffineMap(newMap);
canonicalizeMapAndOperands(&newMap, &newMapOperands);
// CHECK-LABEL: func @gpu_all_reduce_op()
gpu.func @gpu_all_reduce_op() {
%arg0 = constant 1.0 : f32
- // TODO(csigg): Check full IR expansion once lowering has settled.
+ // TODO: Check full IR expansion once lowering has settled.
// CHECK: nvvm.shfl.sync.bfly
// CHECK: nvvm.barrier0
// CHECK: llvm.fadd
// CHECK-LABEL: func @gpu_all_reduce_region()
gpu.func @gpu_all_reduce_region() {
%arg0 = constant 1 : i32
- // TODO(csigg): Check full IR expansion once lowering has settled.
+ // TODO: Check full IR expansion once lowering has settled.
// CHECK: nvvm.shfl.sync.bfly
// CHECK: nvvm.barrier0
%result = "gpu.all_reduce"(%arg0) ({
store %0#1, %arg2[%j] : memref<10xf32>
gpu.return
}
- // TODO(thomasraoux): The transformation should only be legal if
+ // TODO: The transformation should only be legal if
// VariablePointer capability is supported. This test is still useful to
// make sure we can handle scf op result with type change.
// CHECK-LABEL: @simple_if_yield_type_change
// RUN: mlir-opt -split-input-file %s | FileCheck %s
// | mlir-opt | FileCheck %s
-// TODO(pifon): Re-enable LLVM lowering test after IndexedGenericOp is lowered.
+// TODO: Re-enable LLVM lowering test after IndexedGenericOp is lowered.
//
// Test that we can lower all the way to LLVM without crashing, don't check results here.
// DISABLED: mlir-opt %s --convert-linalg-to-llvm -o=/dev/null 2>&1
// CHECK-LABEL: @double_const
spv.func @double_const() -> () "None" {
- // TODO(antiagainst): test range boundary values
+ // TODO: test range boundary values
// CHECK: spv.constant 1.024000e+03 : f64
%0 = spv.constant 1024. : f64
// CHECK: spv.constant -1.024000e+03 : f64
// spv.constant
//===----------------------------------------------------------------------===//
-// TODO(antiagainst): test constants in different blocks
+// TODO: test constants in different blocks
func @deduplicate_scalar_constant() -> (i32, i32) {
// CHECK: %[[CST:.*]] = spv.constant 42 : i32
spv.EntryPoint "GLCompute" @do_something
}
-/// TODO(ravishankarm) : Add a test that verifies an error is thrown
+/// TODO: Add a test that verifies an error is thrown
/// when interface entries of EntryPointOp are not
/// spv.Variables. There is currently no other op that has a spv.ptr
/// return type
// RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s
-// TODO(b/133530217): Add more tests after switching to the generic parser.
+// TODO: Add more tests after switching to the generic parser.
//===----------------------------------------------------------------------===//
// ArrayType
// Tests for lowering 2-D vector.transpose into vector.flat_transpose.
//
-// TODO(ajcbik,ntv): having ShapeCastOp2DDownCastRewritePattern and
-// ShapeCastOp2DUpCastRewritePattern too early in
-// the greedy rewriting patterns misses opportunities
-// to fold shape casts!
+// TODO: having ShapeCastOp2DDownCastRewritePattern and
+// ShapeCastOp2DUpCastRewritePattern too early in the greedy rewriting
+// patterns misses opportunities to fold shape casts!
// No shape cast folding expected.
//
return
}
-// TODO(andydavis) Update test with VTR split transform.
+// TODO: Update test with VTR split transform.
// CHECK-LABEL: func @vector_transfers
// CHECK-COUNT-8: vector.transfer_read
// CHECK-COUNT-4: addf
// -----
#hello_world = affine_map<(i, j) -> (i, 3*d0 + )> // expected-error {{use of undeclared identifier}}
-// TODO(bondhugula): Add more tests; coverage of error messages emitted not complete
+// TODO: Add more tests; coverage of error messages emitted not complete
// -----
#ABC = affine_map<(i,j) -> (i+j)>
}
// -----
-// TODO (support affine.if (1)?
+// TODO: support affine.if (1)?
func @invalid_if_conditional7() {
affine.for %i = 1 to 10 {
affine.if affine_set<(i) : (1)> // expected-error {{expected '== 0' or '>= 0' at end of affine constraint}}
// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-fusion -split-input-file | FileCheck %s
// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-fusion="fusion-maximal" -split-input-file | FileCheck %s --check-prefix=MAXIMAL
-// TODO(andydavis) Add more tests:
+// TODO: Add more tests:
// *) Add nested fusion test cases when non-constant loop bound support is
// added to iteration domain in dependence check.
// *) Add a test w/ floordiv/ceildiv/mod when supported in dependence check.
// *) Add tests which check fused computation slice indexing and loop bounds.
-// TODO(andydavis) Test clean up: move memref allocs to func args.
+// TODO: Test clean up: move memref allocs to func args.
// -----
}
// Fusing loop %i0 to %i2 would violate the WAW dependence between %i0 and
// %i1, but OK to fuse %i1 into %i2.
- // TODO(andydavis) When the fusion pass is run to a fixed-point, it should
+ // TODO: When the fusion pass is run to a fixed-point, it should
// fuse all three of these loop nests.
// CHECK: alloc() : memref<1xf32>
// CHECK: affine.for %{{.*}} = 0 to 10 {
// expected-remark@above {{dependence from 1 to 2 at depth 1 = false}}
// expected-remark@above {{dependence from 1 to 2 at depth 2 = false}}
// expected-remark@above {{dependence from 1 to 2 at depth 3 = false}}
-// TODO(andydavis): the dep tester shouldn't be printing out these messages
+// TODO: the dep tester shouldn't be printing out these messages
// below; they are redundant.
affine.store %v0, %out[%ii, %jj] : memref<64x9xi32>
// expected-remark@above {{dependence from 2 to 0 at depth 1 = false}}
return
}
-// TODO(bondhugula): add more test cases involving mod's/div's.
+// TODO: add more test cases involving mod's/div's.
// -----
(UnrollVectorOp<[2, 2]> $op_results, $a, $b),
[(Constraint<HasShape<[4, 4]>> $a)]>;
-// TODO(andydavis) Add Constraints on lhs/rhs shapes.
+// TODO: Add Constraints on lhs/rhs shapes.
def : Pat<(Vector_ContractionOp:$op_results $a, $b, $c, $masks, $attr0, $attr1),
(UnrollVectorOp<[2, 2, 2]> $op_results, $a, $b, $c),
[(Constraint<HasShape<[4, 4]>> $c)]>;
context, llvm::None, values, op->getAttrDictionary(),
op->getRegions(), inferredReturnTypes))) {
OperationState state(location, OpTy::getOperationName());
- // TODO(jpienaar): Expand to regions.
+ // TODO: Expand to regions.
OpTy::build(b, state, values, op->getAttrs());
(void)b.createOperation(state);
}
//
//===----------------------------------------------------------------------===//
//
-// TODO(riverriddle) This pass is only necessary because the main inlining pass
+// TODO: This pass is only necessary because the main inlining pass
// has no abstracted away the call+callee relationship. When the inlining
// interface has this support, this pass should be removed.
//
.Case<AffineReadOpInterface, AffineWriteOpInterface>(
[](auto op) { boundCheckLoadOrStoreOp(op); });
- // TODO(bondhugula): do this for DMA ops as well.
+ // TODO: do this for DMA ops as well.
});
}
namespace {
-// TODO(andydavis) Add common surrounding loop depth-wise dependence checks.
+// TODO: Add common surrounding loop depth-wise dependence checks.
/// Checks dependences between all pairs of memref accesses in a Function.
struct TestMemRefDependenceCheck
: public PassWrapper<TestMemRefDependenceCheck, FunctionPass> {
&dependenceComponents);
assert(result.value != DependenceResult::Failure);
bool ret = hasDependence(result);
- // TODO(andydavis) Print dependence type (i.e. RAW, etc) and print
+ // TODO: Print dependence type (i.e. RAW, etc) and print
// distance vectors as: ([2, 3], [0, 10]). Also, shorten distance
// vectors from ([1, 1], [3, 3]) to (1, 3).
srcOpInst->emitRemark("dependence from ")
static std::unique_ptr<llvm::Module> compileModuleToROCDLIR(Operation *m) {
auto llvmModule = translateModuleToROCDLIR(m);
- // TODO(whchung): Link with ROCm-Device-Libs in case needed (ex: the Module
+ // TODO: Link with ROCm-Device-Libs in case needed (ex: the Module
// depends on math functions).
return llvmModule;
}
private:
// The TableGen record for this op.
- // TODO(antiagainst,zinenko): OpEmitter should not have a Record directly,
+ // TODO: OpEmitter should not have a Record directly,
// it should rather go through the Operator for better abstraction.
const Record &def;
if (inferType) {
// Generate builder that infers type too.
- // TODO(jpienaar): Subsume this with general checking if type can be
+ // TODO: Subsume this with general checking if type can be
// inferred automatically.
- // TODO(jpienaar): Expand to handle regions.
+ // TODO: Expand to handle regions.
body << formatv(R"(
::llvm::SmallVector<::mlir::Type, 2> inferredReturnTypes;
if (succeeded({0}::inferReturnTypes(odsBuilder.getContext(),
}
void OpEmitter::genInferredTypeCollectiveParamBuilder() {
- // TODO(jpienaar): Expand to support regions.
+ // TODO: Expand to support regions.
const char *params =
"::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &{0}, "
"::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> "
void OpEmitter::genBuilder() {
// Handle custom builders if provided.
- // TODO(antiagainst): Create wrapper class for OpBuilder to hide the native
+ // TODO: Create wrapper class for OpBuilder to hide the native
// TableGen API calls here.
{
auto *listInit = dyn_cast_or_null<ListInit>(def.getValueInit("builders"));
body << " " << builderOpState << ".addTypes(resultTypes);\n";
// Generate builder that infers type too.
- // TODO(jpienaar): Expand to handle regions and successors.
+ // TODO: Expand to handle regions and successors.
if (canInferType(op) && op.getNumSuccessors() == 0)
genInferredTypeCollectiveParamBuilder();
}
// Creating an APInt requires us to provide bitwidth, value, and
// signedness, which is complicated compared to others. Similarly
// for APFloat.
- // TODO(b/144412160) Adjust the 'returnType' field of such attributes
+ // TODO: Adjust the 'returnType' field of such attributes
// to support them.
StringRef retType = namedAttr->attr.getReturnType();
if (retType == "::llvm::APInt" || retType == "::llvm::APFloat")
os << "[TOC]\n\n";
- // TODO(antiagainst): Add link between use and def for types
+ // TODO: Add link between use and def for types
if (!types.empty()) {
os << "## Type definition\n\n";
for (const Type &type : types)
//===----------------------------------------------------------------------===//
void mlir::tblgen::generateOpFormat(const Operator &constOp, OpClass &opClass) {
- // TODO(riverriddle) Operator doesn't expose all necessary functionality via
+ // TODO: Operator doesn't expose all necessary functionality via
// the const interface.
Operator &op = const_cast<Operator &>(constOp);
if (!op.hasAssemblyFormat())
PrintFatalError(loc, "Attribute " + attr.getAttrDefName() +
" does not have the 'constBuilderCall' field");
- // TODO(jpienaar): Verify the constants here
+ // TODO: Verify the constants here
return std::string(tgfmt(attr.getConstBuilderTemplate(), &fmtCtx, value));
}
"(void)tblgen_attr;\n",
depth, attr.getStorageType(), namedAttr->name);
- // TODO(antiagainst): This should use getter method to avoid duplication.
+ // TODO: This should use getter method to avoid duplication.
if (attr.hasDefaultValue()) {
os.indent(indent) << "if (!tblgen_attr) tblgen_attr = "
<< std::string(tgfmt(attr.getConstBuilderTemplate(),
PrintFatalError(
loc, "cannot use AttrConstraint in Pattern multi-entity constraints");
} else {
- // TODO(b/138794486): replace formatv arguments with the exact specified
+ // TODO: replace formatv arguments with the exact specified
// args.
if (entities.size() > 4) {
PrintFatalError(loc, "only support up to 4-entity constraints now");
auto &info = symbolInfoPair.getValue();
os.indent(4) << info.getVarDecl(symbol);
}
- // TODO(jpienaar): capture ops with consistent numbering so that it can be
+ // TODO: capture ops with consistent numbering so that it can be
// reused for fused loc.
os.indent(4) << formatv("::mlir::Operation *tblgen_ops[{0}];\n\n",
pattern.getSourcePattern().getNumOps());
// `{0}` resolves to an `Operation::result_range` as well as cases that
// are not iterable (e.g. vector that gets wrapped in additional braces by
// RewriterGen).
- // TODO(b/147096809): Revisit the need for materializing a vector.
+ // TODO: Revisit the need for materializing a vector.
os << symbolInfoMap.getAllRangeUse(
val,
" for (auto v : ::llvm::SmallVector<::mlir::Value, 4>{ {0} }) {{ "
LLVM_DEBUG(llvm::dbgs() << '\n');
auto fmt = tree.getNativeCodeTemplate();
- // TODO(b/138794486): replace formatv arguments with the exact specified args.
+ // TODO: replace formatv arguments with the exact specified args.
SmallVector<std::string, 8> attrs(8);
if (tree.getNumArgs() > 8) {
PrintFatalError(loc, "unsupported NativeCodeCall argument numbers: " +
// Otherwise this is an unbound op; we will use all its results.
return pattern.getDialectOp(node).getNumResults();
}
- // TODO(antiagainst): This considers all NativeCodeCall as returning one
+ // TODO: This considers all NativeCodeCall as returning one
// value. Enhance if multi-value ones are needed.
return 1;
}
// The argument in the result DAG pattern.
auto patArgName = node.getArgName(argIndex);
if (leaf.isConstantAttr() || leaf.isEnumAttrCase()) {
- // TODO(jpienaar): Refactor out into map to avoid recomputing these.
+ // TODO: Refactor out into map to avoid recomputing these.
if (!opArg.is<NamedAttribute *>())
PrintFatalError(loc, Twine("expected attribute ") + Twine(argIndex));
if (!patArgName.empty())
os << formatv(" auto tblgen_instance = {0}::{1}(tblgen_attrVal);\n",
enumAttr->getCppNamespace(), avail.getQueryFnName());
os << " if (tblgen_instance) "
- // TODO(antiagainst): use `avail.getMergeCode()` here once ODS supports
+ // TODO` here once ODS supports
// dialect-specific contents so that we can use not implementing the
// availability interface as indication of no requirements.
<< std::string(tgfmt(caseSpecs.front().second.getMergeActionCode(),
#include <chrono>
#include <cstring>
-// TODO(antiagainst): It's generally bad to access stdout/stderr in a library.
+// TODO: It's generally bad to access stdout/stderr in a library.
// Figure out a better way for error reporting.
#include <iomanip>
#include <iostream>
RETURN_ON_VULKAN_ERROR(physicalDeviceCount ? VK_SUCCESS : VK_INCOMPLETE,
"physicalDeviceCount");
- // TODO(denis0x0D): find the best device.
+ // TODO: find the best device.
physicalDevice = physicalDevices.front();
if (failed(getBestComputeQueue()))
return failure();
void expectDiagnostic(StringRef errorMessage) {
ASSERT_NE(nullptr, diagnostic.get());
- // TODO(antiagainst): check error location too.
+ // TODO: check error location too.
EXPECT_THAT(diagnostic->str(), StrEq(std::string(errorMessage)));
}
min_version = ''
elif min_version:
min_version = 'MinVersion<SPV_V_{}>'.format(min_version.replace('.', '_'))
- # TODO(antiagainst): delete this once ODS can support dialect-specific content
+ # TODO: delete this once ODS can support dialect-specific content
# and we can use omission to mean no requirements.
if for_op and not min_version:
min_version = DEFAULT_MIN_VERSION
max_version = enum_case.get('lastVersion', '')
if max_version:
max_version = 'MaxVersion<SPV_V_{}>'.format(max_version.replace('.', '_'))
- # TODO(antiagainst): delete this once ODS can support dialect-specific content
+ # TODO: delete this once ODS can support dialect-specific content
# and we can use omission to mean no requirements.
if for_op and not max_version:
max_version = DEFAULT_MAX_VERSION
# a core symbol since the given version, rather than a minimal version
# requirement.
min_version = DEFAULT_MIN_VERSION if for_op else ''
- # TODO(antiagainst): delete this once ODS can support dialect-specific content
+ # TODO: delete this once ODS can support dialect-specific content
# and we can use omission to mean no requirements.
if for_op and not exts:
exts = DEFAULT_EXT
else:
caps = 'Capability<[{}]>'.format(', '.join(prefixed_caps))
implies = ''
- # TODO(antiagainst): delete this once ODS can support dialect-specific content
+ # TODO: delete this once ODS can support dialect-specific content
# and we can use omission to mean no requirements.
if for_op and not caps:
caps = DEFAULT_CAP
else:
arg_type = 'Variadic<SPV_Type>'
elif kind == 'IdMemorySemantics' or kind == 'IdScope':
- # TODO(antiagainst): Need to further constrain 'IdMemorySemantics'
+ # TODO: Need to further constrain 'IdMemorySemantics'
# and 'IdScope' given that they should be generated from OpConstant.
assert quantifier == '', ('unexpected to have optional/variadic memory '
'semantics or scope <id>')