* Compute common type for shape elements in BroadcastHelper
The corresponding dimensions in the input/output tensors in a broadcast
operations may have the same value, but different types (e.g. int32 vs
int64).
When the broadcast helper tries to unify the dimensions it also needs
to compute the common type to hold the dimension.
* Cast and simplify both members of `Range`
Only the `min` member was type-casted, which could lead to ranges with
different types for `min` and `extent`.
Move the casts to the argument of Simplify, so that they can be eliminated
if they aren't needed.
* Type-check iv domain ranges, use cast only if needed in MakeLoopNest
In some cases the domain ranges had the `min` and the `extent` values
be of different types (e.g. [(int64)0, 32)). This is an error, and it
can lead to compilation failures later on. Add a check for equal types
here to catch this early.
Also, only add the cast operation when the desired type differs from
the current one to keep the expressions simpler.
* Check that variable and substituted expression have same types
Add a check to IRSubstitute to detect when the type of a variable and
the type of the expression to replace it with have different types.
* Add testcase
* [TVMScript] Use void for lambda parameters, allow mismatch in Substitute
When the script parser deals with lambdas, it creates Var objects for each
parameter. Their actual types are not known at the time, and the properly
typed variables are subtituted in the body later. Since the default dtype
of a Var is "int32", this could lead to a type mismatch in Substitute.
To deal with this scenario, use "void" for newly created Vars in the
parser, and add an exception to Substitute to allow replacing void Vars
with expressions of any type.
* Fix type error in test_reduce_combiner_simplify
* Restart CI
Co-authored-by: Jiawei Liu <jaway.liu@gmail.com>
* [TIR.Constant] U1 usecase
Constants are now aggregated into one struct and initialized in default_lib0.c
file
Change-Id: I34d61f8139c8a92c06944fe990ba892a660476fd
Unit test fixed
Change-Id: I436e7b6d6b3064b3f8bbfbb048d4296b63a6b69c
* Refactored
Addressed:
* PoolInfo splitted to WorkspacePoolInfo and ConstantPoolInfo
* workspace_byte_alignment moved to ExecutorCodegenMetadata
* getModuleAlignment -> GetModuleAlignment
* GenerateInternalWorkspaceBuffers refactored
* reverted format change of src/tir/transforms/legalize_packed_calls.cc
* addressed comments for src/tir/usmp/analysis/extract_buffer_info.cc
* removed commented code from include/tvm/tir/usmp/utils.h
Change-Id: I7d1b32884b0e5992e2e00c7838c85e425d9c25fd
* more unit test fixes
Change-Id: I573a05fa1cb4037ae83691f7dff2c2724b1d7700
* More refactoring and unit test fixes
Added ConstantMemoryPools
Change-Id: If1e391c631575980564bca790ba33748c82d907f
* bugfix
Change-Id: Iacc7a9d734a505dfa0d8d32d23ea3f57e6de8582
* refactoring. added constant_alignment
added constant_alignment
unit tests updated
Change-Id: I378193cb9e675e352c61d96ff4e09655090053e1
* unit-test bugix
Change-Id: Ia4411d59c4a376c01326fed366cdb196a432899e
* unit test fix
Change-Id: Ia2077bdeb1d2c6c9827eeef90ab410ae31b8c4a4
* Added support for c++ runtime
* refactored
* renamed pools and consts
renamed pools and consts to workspace_pools and constant_pools
* addressed upstream comments
* addressed upstream comments-2
* addressed upstream comments-3
* [TVMSCRIPT] Improve tvmscript type hints
- Change numeric types to classes so they work as function arguments.
- Add var as a class.
- Add floordiv, index, and mod to PrimExpr.
* use Union
* [TVMScript] Allow T.Buffer[] arg annotation to use int as shape
Both the function `tvm.tir.decl_buffer` and the TVMScript
`T.match_buffer` expression allow a `PrimExpr` to be passed as the buffer
shape, which is interpreted as a 1-d buffer of that size. This allows
the same behavior to be used in the `T.Buffer` syntactic sugar.
(e.g. `A: T.Buffer[16, "float32"]` instead of `A: T.Buffer[(16,), "float32"`)
* Fixed round-trip when buffer size contains an expression
* [TVMScript] Support function call to help construct AST
* add test
* update test
* more comment
* fix for avoiding Buffer.vload(...) case
* update parse error msg
* wrap func call with try / catch, emit error msg
* silence pylint
* [TVMScript] Allow `val = buf[index]` without type annotation
Other instances of `var = expr` were previously allowed without
requiring a type annotation, by using the dtype of the expression as
the dtype of `var`. This behavior didn't work for `buf[index]`
expressions, which are internally represented as `BufferSlice` python
objects, and only converted to `BufferLoad` primexprs when used as an
expression.
This commit adds a `dtype` property to `BufferSlice`, allowing
`buf[index]` to be used in a let statement without a type annotation.
* Reverted a wider change
Automatically adding a type annotation to Var if it could be
determined from the dtype let the unit test directly compare the
annotated and unannotated versions of buffer load. Unfortunately, it
also broke 54 unrelated tests, so that change is removed from this PR.
* add get_c_struct_name() method to Metadata to distinguish struct type name in llvm
* add metadata serialization support to llvm codegen
* Organize MetadataQueuer into a separate file.
* Add DiscoverArraysVisitor to metadata_utils
* Fill DLTensor metadata in LegalizePackedCalls.
* Improve error message from Call asserts
* Pass non-String device_context down to codegen.
* this is necessary to allow CodeGenCPU to emit calls that include resource_handle.
* Scope usage of lvalue refs in LowerTVMBuiltin to avoid corrupt memory.
* test fixes
* Also fill preflattened_buffer_map (TODO, maybe don't do this)
* Fix C codegen.
* Set USMP elem_offset to 0.
* Clarify calculation of byte_offset from elem_offset.
* fix tests
* Fix arm compile warning
* Fix hexagon test.
* previously I believe we required interface_api == "c", but
this really means to generate C API bindings, and we are generating
"packed" bindings.
* I think "c" was chosen here because the distinction between
interface-api and use-unpacked-api is confusing. "c" interface-api
means to generate an entrypoint API for microcontrollers that
accepts bare data buffers. "packed" interface-api means to generate
a TVMBackendPackedCFunc entrypoint. use-unpacked-api forms the same
determination for the operator functions.
* A further confusion here is that there are two ways to call
"packed" operator functions: tir.tvm_builtin_call_packed and
tir.tvm_builtin_call_cpacked. This distinction describes whether or
not to late-bind calls via TVMBackendGetFuncFromEnv. Right now, AOT
only ever requires call_cpacked because target_host == target, and
for all suitable target_host, we expect a single DSO-exportable
runtime.Module. When we move away from this by introducing
heterogeneous target support to AOT, we can use this as a condition
to help us choose between call_cpacked and call_packed (and
possibly add a compile-time option to assert it is call_cpacked,
for situations where we really don't want call_packed).
* Document T.preflattened_buffer
* Fix test_aot_legalize_packed_calls
* Address manupa comments
* Fix convert_pool_allocations_to_offsets test.
* lint
* Fix T.preflattened_buffer
* Add preflattened_buffer_map to TIRTextPrinter
* Fix tests
* Fix BYOC
* Fix invoking C device API.
* remove comments
* Address Mousius comments
* lint
* lint
* Fix GMock linking on new CMake
* address masahi comment
Co-authored-by: Masahiro Masuda <masahi129@gmail.com>
* Respect dtype in Scalarize.
* Add unittest.
* Fix lint.
* Promote dtype of IntImm to match loop_var in For.
* Fix dtype mismatches.
* Lint
* Lint.
* jostle ci
* Match dtype in hybrid parser.
This demonstrates how to selectively extract and tune tasks from a whole relay mod, and apply the tuned schedule during the final `relay.build(...)`.
This flow is entirely different from existing tests in `test_meta_schedule_tune_relay.py` where ALL ops are extracted and auto-scheduled by MS. My test extracts only int8 `dense` op, applies a manual TIR schedule on it, and leaves int8 `batch_matmul` to be scheduled by TE.
This also serves as an example of autotvm style manual template + tensorization. The manual TIR schedule is equivalent to TE VNNI `dense` schedule in https://github.com/apache/tvm/blob/ce335c3a74185df6cc1152e53c60695d8a418d8e/python/tvm/topi/x86/dense.py#L366-L375
## Context
When dealing with end-to-end models, we note that some tensors may have large shapes. Thus, when designing graph-level IR, we sometimes use `int64` instead of `int32` for the shape. Below is an dense GeMM example which has `int64` input tensor shape:
```python
@tvm.script.ir_module
class Module:
@T.prim_func
def main(rxplaceholder: T.Buffer[(1, 512), "float32"], rxplaceholder_1: T.Buffer[(T.int64(1000), T.int64(512)), "float32"], T_matmul_NT: T.Buffer[(1, T.int64(1000)), "float32"]) -> None:
# function attr dict
T.func_attr({"global_symbol": "dense", "tir.noalias": True, "op_pattern": 3})
# body
# with T.block("root")
for i0_0, i1_0, i0_1, i1_1, i2_0, i0_2, i1_2, i2_1, i0_3, i1_3 in T.grid(1, 4, 1, 25, 8, 1, 10, 64, 1, 1):
with T.block("T_matmul_NT"):
i = T.axis.spatial(1, 0)
j = T.axis.spatial(T.int64(1000), i1_0 * T.int64(250) + i1_1 * T.int64(10) + i1_2)
k = T.axis.reduce(512, i2_0 * 64 + i2_1)
T.reads(T_matmul_NT[i, j], rxplaceholder[i, k], rxplaceholder_1[j, k])
T.writes(T_matmul_NT[i, j])
T.block_attr({"layout_free_placeholders":[rxplaceholder_1], "meta_schedule.tiling_structure":"SSRSRS"})
with T.init():
T_matmul_NT[i, j] = T.float32(0)
T_matmul_NT[i, j] = T_matmul_NT[i, j] + rxplaceholder[i, k] * rxplaceholder_1[j, k]
```
## Problem
Though our TVMScript printer can easily print `int64` constants, the parser had poor support for `int64`. So this PR introduces some parser support for `int64`, basically about the data type of loop variables, block iterators and block read/write regions.
Besides the parser, most of the TIR schedule primitives didn't take `int64` into account in their implementations. These schedule primitives will be fixed and updated in recent future, in followup PRs.
* [TIR] Added BufferLoadNode::LegalizeDtype
When modifying a BufferLoad object, the return dtype must also be
updated. This exposes the legalization function, so that passes that
use `BufferLoad::CopyOnWrite` to modify the buffer/indices don't need
to repeat the logic to update the dtype returned.
* Replacing Store/Load in Stmt/Expr Visitor/Mutator
* Removing Store/Load from optimization passes
- UpdatePointerStorageScope
- UnrollLoop
- ThreadSync
- LinearAccessPatternFinder
- StoragePlanRewriter
- VectorTypeRewriter
- VectorTypeAccessChecker
- NarrowDataType
- IRConvertSSA
- CompactBufferRegion
* Removing Store/Load from examples
- ConvertAddToSubtract
* Replacing Store/Load in StorageFlatten
Now, outputs BufferLoad/BufferStore with a flattened buffer object.
temp commit, replacing Store/Load, BufferBindUnwrapper
temp commit, replacing Store/Load, StorageFlattener
* Replacing Store/Load in utility passes.
- StmtSimplifier
- IRSubstitute
- BaseInliner
- FeatureVisitor
* Replacing Store/Load in analysis functions
- StorageAccessVisitor
- VarTouchedAnalysis
- MemoryAccessVerifier
- InplaceOpVerifier
- GPUCodeVerifier
- VarTouchVisitor
- LCADetector
- BlockReadWriteDetector
- InstrumentBoundCheckers
* Replacing Store/Load in lowering/legalization passes.
- MakeCrossThreadReduction
- CacheReadRewriter/CacheWriteRewriter
- InjectVirtualThread
- InjectDoubleBuffer
- InjectCopyIntrin
- LowerWarpMemory
- LowerThreadAllreduce
- LowerThreadAllreduce
- LowerCustomDatatypes
- LowerTVMBuiltin
- CoProcSync
- MergeDynamicSharedMemAllocations
- VectorizeLoop
- BF16Legalize
* Replacing Load/Store in codegens.
- Device code generators
- CodegenC
- CodegenLLVM
- CodeGenOpenCL
- Utilities used during codegen
- ArgBinder
- MakePackedAPI
- ReturnRewriter
- SplitHostDevice
- Execution environments
- CodeGenStackVM
- CodeGenHybrid
- AOTExecutorCodegen
* [UnitTest] Add unit tests to test physical layout remapping.
* Updated tvm::address_of() to hold BufferLoad instead of Load.
* [TIR] Added IndexMap class.
Holds a set of variables representing the input indices and
expressions in terms of those input indices.
TODO:
- Add validation, the index mapping should be invertible.
- Add helper function, apply mapping to a set of indices.
- Add helper function, apply mapping to bounds of input indices.
* Updated Buffer::vstore/vload to return BufferLoad/BufferStore objects.
StorageFlatten/FlattenBuffer passes updated to modify the
buffer/indices directly, rather than using vload/vstore.
- Primary purpose of vstore/vload is to allow IR written in python to
define vectorized load/store. This usage is maintained by returning
a BufferLoad/BufferStore node whose index is a Ramp.
- Previously, vstore/vload was also used to compute the 1-d physical
index of a location within a N-d tensor. This usage will no longer
be allowed, as it would not allow layout transformations to be
performed after a schedule definition, but any uses of the buffer
are flattened.
* [TE] Added Stage::transform_layout to the C++ TE implementation.
Adds an `Array<IndexMap>` in the stage to define the transformations
to be applied on the tensor's layout. As of this commit, this mapping
isn't propagated into the TIR graph yet.
* Replace Store/Load with BufferStore/BufferLoad in ir_builder
* [TE] Added Stage.transform_layout to the Python TE interface.
Allows users to specify `s[A].transform_layout(mapping)`, and
propagate into the TE definitions.
* Added pre_flattened_shape/pre_flattened_stride fields to Buffer.
The shape and stride checks performed in ArgBinder::BindDLTensor
(called from MakePackedAPI) require the tensor shape/strides prior to
index flattening. Therefore, though it is no longer used by the
low-level code generators, we must maintain that information for use
in MakePackedAPI.
* [UnitTest] Test N-d indices exposed to low-level codegen
When using te.AXIS_SEPARATOR in the call to .transform_layout, this
should define groups of axes, each of which is flattened to a single
axis, then exposed to the low-level codegen.
* [TIR] Added PrimFunc attribute "layout_transform_map", filled from TE.
Propagated the TE definition of the physical layout into the TIR
graph.
* Added pre_flattened_type.
If a boolean tensor is backed by an int8 buffer, the check on the
argument buffer's type should be against the boolean type.
When rebasing this PR, should be placed after the addition of
pre_flatten_shape/pre_flatten_strides.
* [UnitTest] Added tests for loop iteration order.
After transformation, the iteration order should follow the new
transformed axes. In addition, the loop iteration variables should be
exposed through the TE interface for further manipulation.
* [TIR] Added BufferNode::axis_separators
- Add axis_separators to represent divisions between groups
of tensor axes, where each group is flattened into a single
output axis, to be exposed to the low-level code generators.
- Expose axis_separators to the python interface.
- Update existing C++ calls to the Buffer() constructor.
* [TIR] Added ApplyLayoutTransforms as part of StorageFlatten.
For any buffers that have layout transforms defined in the
"layout_transform_map" attribute of a PrimFunc, rewrite access into
the buffer such that they use the updated ordering.
* Update usage of ir_builder where necessary.
* [TE] Implement te::Transform
Similar to Fuse and Split, this represents a modification to the
existing loop iterations.
* [TE] Added Stage::set_axis_separators.
In C++, this is implemented as an `Array<IntImm>`, specifying
pre-flatteneing axes after which a new post-flattening should be
started. The python interface uses a sentinel value
`te.AXIS_SEPARATOR` in the call to `transform_layout`, which is then
used to define the array of axis separators.
* [TIR] Expose tir.transform.ApplyLayoutTransforms for testing
* [TE] Rewrite loop iteration order
After .transform_layout, rewrite leaf_iter_vars to follow the updated
order. Use the te::Transform iter_var relationship to track use of
the transformed variable.
* [TE] Fill BufferNode::axis_separators from StageNode
During ScheduleOps and SchedulePostprocToPrimfunc, the axis separators
defined in the stage must be passed through to the TIR BufferNode.
* [TE] Return transformed iteration variables
* Moved Buffer's pre-flatten information to PrimFunc.
Since the pre-flatten information is only used for validating user
inputs, it makes much more sense to store it alongside the buffer_map.
* Updated ethos-u C++ unit tests to remove use of Load/Store.
* Bugfix, layout transformation.
Error occured during conversion from TE to IRModule, when layout
transforms were applied to a reader of a `cache_read`.
* In test directory, replacing all instances of T.load.
* Return buffer object from tvm.tir.script.scope_handler.Allocate
Now that the load/store require buffer objects, allocation should also
return a buffer object to be used.
* Added .astype to tvm.script.tir.node.BufferSlice
Since `buf[i]` returns a `BufferSlice`, this lets the TIR examples
that use `buf[i].astype('out_dtype')` continue functioning.
* Replacing all T.store TIR calls.
* Added LOG(FATAL) in constructor of Store/Load nodes.
* Updated tvmscript parser to report error for Store/Load nodes.
* [TVMScript] Added T.preflattened_buffer stmt
Used to specify `PrimFunc::preflattened_buffer_map`. Takes an argument
of the postflattened buffer, so that it will work for both simple
declarations and `T.match_buffer` statements without needing to
introduce a param handle. All other arguments are identical to
`T.match_buffer.`
* [TVMScript] Updated TVMscript for BufferLoad/BufferStore
- Use `T.preflattened_buffer` calls in TVMScript to represent
`PrimFunc::preflattened_buffer_map`.
- Remove `T.buffer_decl` for return value of `T.allocate`, now that
`T.allocate` returns a buffer.
- For buffer access as a different type, make a `T.buffer_decl` for
those accesses.
* Updated test_tvmscript_roundtrip.py for BufferLoad/BufferStore.
* Updated TIR reference in USMP pool allocation unit tests.
Using let var handles as the data pointer in buffers, rather than just
as `T.load`/`T.store` arguments, requires annotation as
`T.Ptr[T.primtype]`, rather than as `T.handle`.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* fixup! Replacing all T.store TIR calls.
* fixup! Replacing all T.store TIR calls.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* fixup! In test directory, replacing all instances of T.load.
* tir.ComputeInline, correct variable count.
Previously, this metaschedule primitive relied on `tir::UndefinedVars`
ignoring the data pointer of BufferLoad/BufferStore nodes. When
`tir::UndefinedVars` was updated to visit the data pointer, similar to
the previous behavior when visiting Load/Store nodes, this caused the
count of undefined variables to be unexpectedly high.
* fixup! Replacing all T.store TIR calls.
* fixup! Updated Buffer::vstore/vload to return BufferLoad/BufferStore objects.
* fixup! In test directory, replacing all instances of T.load.
* fixup! In test directory, replacing all instances of T.load.
* fixup! Replacing all T.store TIR calls.
* Expose Buffer index flattening function to Python.
* Updated test_tir_buffer.py offset tests.
Replacing calls to `Buffer.vload` with `Buffer.offset_of`, when
testing the index calculations.
* fixup! Replacing all T.store TIR calls.
* fixup! Replacing all T.store TIR calls.
* fixup! Updated Buffer::vstore/vload to return BufferLoad/BufferStore objects.
* fixup! Replacing Store/Load in lowering/legalization passes.
* fixup! Replacing all T.store TIR calls.
* fixup! Updated ethos-u C++ unit tests to remove use of Load/Store.
* fixup! Replacing Store/Load in lowering/legalization passes.
Fix linting for inject_double_buffer.cc
* fixup! Updated ethos-u C++ unit tests to remove use of Load/Store.
* fixup! Added .astype to tvm.script.tir.node.BufferSlice
* fixup! In test directory, replacing all instances of T.load.
* fixup! Replacing all T.store TIR calls.
* fixup! Replacing all T.store TIR calls.
* fixup! In test directory, replacing all instances of T.load.
* fixup! Replacing all T.store TIR calls.
* fixup! Replacing Store/Load in lowering/legalization passes.
* [UnitTests] Added T.preflattened_buffer in expected result
* fixup! In test directory, replacing all instances of T.load.
* [UnitTests] Bound checker update, compare against N-d buffer bounds.
* Fixup, bound checker vectorize test.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* [UnitTest] Fixed breakage in InjectRollingBuffer test.
Needed a bit more re-writing than usual, because the test was
explicitly calling lowering passes, then calling `tvm.build`. Fixed
by using the standard lowering flow, with preprocessing steps
inserting with `tir.add_lower_pass`.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* [UnitTest] Fixed breakage in flatten buffer unit tests.
- Updated pass to allow BufferStore/BufferLoad nodes to be visited
before the block's alloc buffer.
- Added `T.preflattened_buffer` annotations.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* [UnitTests] Fixed breakage in test_tir_buffer.py
- Updated vload test for new behavior.
- Added test for offset_of, testing behavior no longer in vload.
- Added null check for buffer visitor.
* fixup! Replacing Load/Store in codegens.
* [UnitTest] ComputeInline, opaque access test updates
* [UnitTest] Fixup, allow unit test to use `ib.pointer()[0]`.
* fixup! Replacing Load/Store in codegens.
The updated CodegenLLVM should use the BufferStore/BufferLoad
convention of indexing by `sizeof(dtype)`, rather than
`sizeof(dtype.element_of())`.
* fixup! Replacing Store/Load in lowering/legalization passes.
BF16Legalize should also update the preflattened_buffer_map, since it
is overwriting the `BufferNode::data` stored in the buffer_map.
* fixup! Replacing all T.store TIR calls.
* Fixed failing codegen c host unit tests.
- Generated functions were making `uint8_t*` parameter arguments for
array handle for return value, rather than the earlier `void*`.
- New parameter type was due to using
`PointerType(PrimType(DataType::UInt(8)))` as the type annotation, to
be usable as `BufferNode::data`.
- Changing to `PointerType(PrimType(DataType::Void()))` still allows
usage as buffer, more appropriately expresses semantics.
- Updated C codegens to allow `void*` types to be generated from
variables with type annotation, in addition to the previous behavior
of `DataType::Handle()` variables without type annotation.
* Fixup, StorageFlatten when applied to post-StorageRewrite functions.
Identified in a test that applied `tvm.lower`, then `tvm.build` on the
result. If the result of an allocate node is used as the backing
buffer for multiple buffers, such as the output of the StorageRewrite
pass, then StorageFlatten would erroneously think that the second
occurrence was an usage without earlier definition.
* fixup, StorageFlatten
When flattening a boolean buffer, the backing buffer should have type
int8, not the preflattened buffer.
* Bugfix, correctly represent void* in LLVM IR.
* Update, replace tir.Load with tir.BufferLoad
* Added TVMScript error check for matching buffer/index dimensionality
Needed for tests/python/unittest/test_tvmscript_error_report.py::test_high_dim_store
* Bugfix, correct return type when lowering custom datatype.
* Bugfix, removed unused primfunc from test_tvmscript_complete.py
* Updated test_meta_schedule_postproc_verify_gpu_code.py TIR
Replaced Load/Store with BufferLoad/BufferStore.
* Allowed ramp nodes with buffer use analysis.
* Updated tests in test_meta_schedule_postproc_verify_gpu_code.py
Needed dummy writes to prevent buffer resizing, in order to trigger
the verification failure due to memory limits.
* Updated TIR examples to be compatible with buffer dimension check.
* Corrected section header in docstring.
* Corrected indices size check in CogeGenC.
* Fixed breakage in LowerThreadAllreduce.
Since the AllocateNode is rewritten, any buffers that refer to those
variables must also be rewritten.
* [UnitTests] Replaced Store/Load in CUDA codegen tests.
* Resolved breakage in C-based codegen for vectorized store/load.
Needed to update to new convention of using the buffer's element type
as the stride.
* Bugfix, incorrect LCA for buffer access in root scope.
This had been present before the BufferLoad/BufferStore changes, but
hadn't triggered on tests using Load/Store nodes.
* Added docstrings for TransformNode member variables.
* Added TODO for future removal of preflattened_buffer_map.
* Fixup, transform layout + cache write tests.
The correct sequence is to first apply any caching as needed, then to
apply layout transformations, and finally to apply thread binds for
the computation step.
* Bugfix, correct element type for scalarized access.
* Bugfix, cuda buffer indexing when declared as different type.
* Cuda codegen, update reference.
* Bugfix, lower allreduce
Loads of the output of the reduction should be replaced for all
buffers sharing a buffer pointer, not just for the buffer object
itself.
* Removed obsolete comment.
* Changed PrimFunc constructor preflattened_buffer_map to Optional
* Removed flatten_buffer argument from T.match_buffer.
* Correct call to VarUseDefAnalysis::VisitBuffer
* Reverted unintentional testing change, lanes=2.
* Updated lower_cross_thread_reduction to use buffer in allreduce
* Updated transform_layout test to disable CSE
* Updated CSE unit tests to use BufferStore
* Replaced Store/Load for vta.transform and unit tests.
* Updated unit tests for lower_cross_thread_reduction.
* Updated arange to use scalar tensors.
The start/stop/step tensors are declared as 0-d scalar tensors, but
were accessed as 1-d tensors.
* Fix breakage in ethosu constant encoding.
Buffers generated by "ethosu_copy" should have their buffer objects
rewritten, but shouldn't have their size updated in ethosu-specific
Call nodes.
* Fix breakage in ethosu call argument checks.
Need to pull out indices from BufferLoad holders, not Load.
* Resolve breakage from mismatched shape/index dimensions
* Split out encoded parameters from preflattened buffer map.
* Updated buffer shape/index dimensions to match in more ethosu tests
* Fixed lint error
* Removed debug code
* Moved arith::Analyzer local variable to class member
* Fixed SSA conversion of allocations.
Can occur if allocation is inside an unrolled loop. Added unit test
to catch this failure mode.
* Ethos-u index/buffer dimension updates.
* Updated ethosu passes to handle buffer load/store.
* Resolved bug in tvmscript printing of duplicate buffers.
* Fix breakage in ethos-u test_assign_addresses, encode constants
* Apply same changes to T.allocate_const as to T.allocate
Return a buffer when used in TVMScript, allow for aliasing buffers.
* Fix lint errors.
* Further updates for ethos-u tests.
* Updated ethos.u buffer sizes in test.
* Updated tir.BindParams to use BufferLoad instead of Load.
* Updated topi.cuda.scan implementation to follow buffer dimensions.
* Resolved breakage when flattening AllocateConst nodes.
* Resolved breakages from latest merge with main.
* Corrected error in merge.
* Use empty indices for rank-0 tensor.
* Added ir_builder workaround for 1-d indexing.
* Consistent buffer access type in LLVM codegen, to match C codegen
* StorageRewrite, update indices of modified buffers.
* Dynamic relay nodes, access 0-d tensors with 0-d indices.
* BFloat16 legalization, update buffer type.
* Updated meshgrid to use 0-d index for 0-d buffer.
* Corrected boolean handling in Allocate nodes.
* Added workaround to unpack 1-d Tensor indices into N-d buffer indices.
* Resolved a few more failures in relay tests on cuda.
* Resolve linting
* CI bump
* Updated renormalize_split_pattern tests to use BufferLoad/BufferStore
* Fixed cuda codegen checks for BufferStore/Ramp.
* Simplify indices further, needed to avoid cuda register limit.
* fixed dyn onehot shape func accessing 1d buffer with ()
* Fixed codegen indexing for int4 scalar types.
* Temporary workaround for incorrect constant folding.
Need to further investigate vectorized LLVM constants
* s/find_allocate_usage/FindAllocateUsage/g
* Added buffer type consistency TODO.
* Improved comment on address_of Op.
* Rename LegalizeDtype to LegalizeDType, made private.
* fix format and lint errors
* Disable vectorization of AllocateConst buffer in StorageRewrite.
* Pass buffer_map through to the PrimFunc in cmsisnn
* try disabling problematic winograd test case
* try different way of buffer mapping in storage_rewrite
* Removed unnecessary ramp node in ir_builder.
* Updated LLVM codegen for buffer indexing.
TVM data arrays are always densely packed. If the LLVM type
corresponding to a vectorized TVM datatype contains padding for
alignment, the array location should be computed based on the
primitive element type.
Co-authored-by: Masahiro Masuda <masahi129@gmail.com>
Co-authored-by: adstraw <astraw@octoml.ai>
A quick fix of the parser issue mentioned in #10327 .
Ranges and loops require `start` and `stop` to be PrimExpr, however, `BufferSlice` is not always scalar so it's not a `PrimExpr`.
This PR performs the transformation.
* [TIR] Introduce tir.allocate_const to TIR
This PR is adding non-scalar constant representation in TIR. This is used to
express constants (i.e., parameters) in the TIR instead of bypassing the
TIR as it's done until now.
Change-Id: Id3afc4d7197260cb43ecde60f05ccbce3fc42430
Co-authored-by: Giuseppe Rossini <giuseppe.rossini@arm.com>
Change-Id: Id4a09a637c9c1fd7d49989c6c10f474a78569e18
* [TIR] Integrate tir constant nodes in compilation pipeline
This PR integrates tir.allocate_const to the compilation pipeline to support --link-params.
Change-Id: Ic8d0cb75d596299fcae7078b304598afbf0c5494
Co-authored-by: Giuseppe Rossini <giuseppe.rossini@arm.com>
Change-Id: Id98cc682bbfacfe75c4d8b260fd41658f1f196b2
* [TIR] tir.const extraction
This commit tries to implement an amendment to tir.constant RFC
with centralized storage of constant data within the IRModule
Please note that data and irmod_storage_idx are not mutual exclisive
further more the irmod_storage_idx is valid only immediatly after
prim func addition to the mod or after update within the mod.
If prim func is out of the the module scope then the index become
meangless. irmod_storage_idx also is not used in calculation of hash
function of the tir.constant node.
Change-Id: I40742ed580468b0252ea3fec02184cba65e20871
* unit test fixed
Change-Id: Ied2186554d4cbad44b2346216c8be92449e55732
* cmsis-nn codegen fix
Now handled case when params of the functions came as constants
Change-Id: I5874e182e34ef94e23048eaf3c61b01a56d91131
* Fixes for unittests
Change-Id: I5b82ee3f80337155706b5470973f494a301b5d90
* Rebasing tests fixes
Change-Id: I94ac87907081bab53c1dd1ab2db106ae057b4b19
* Linter: added method param description
Change-Id: I2f8c4c8d244b74c794abaa6079c46cc593ffcbdb
* Printing removal fix
This patch removes forgotten print in fuse_ops
Change-Id: I4bb5934f3b4cd5fde19d36a8e3319aae136bce8a
* Bugfix
Fixed concurrent map update bug here
Change-Id: Ifec3bf5030086d9079b9e493096f17dfd82297ec
* Reworked logic for not to introduce empty constant list to modue attrs
Change-Id: I082c85b3b4b70c218f0d714f5613ef6e178bd020
* Added support for tir builtin::tvm_access_ptr
This fixed unit tests for tests/python/integration/test_arm_mprofile_dsp.py
Change-Id: I10919f301ef9ddc3fd87f0e1a8414e9a52fc7938
* Unit test fix
Fixes unit tests in torch frontend
Change-Id: I6c179834f93dd202605d1ce5a7f07d987b9dc469
* Addressed requested changes
Addressed changes requested upstream
Change-Id: I741e52b89eb285732c23b1ac7ff277e757a088c3
* Namespace usage changed to conform earlier C++ standard
Change-Id: I1b29238cfe2a6bedb525f4f823a3a540f631d836
* Bugfix
Change-Id: I57a44b714b307278a243817ec2864e53ad31366b
* updated IRModuleNode::ExtractPrimFuncConstants
Updated IRModuleNode::ExtractPrimFuncConstants as per
request upstream.
Change-Id: I35db0145fb5827efd0445ce665d0c99465274016
* Minor changes
typo fixd
renamed ExtractPrimFuncConstants to ExtractConstants
removed getters/setters from FuseMutator and added parametrized
constructor
Change-Id: Ib2326805781779b88c963a8642ff683c8755956e
* Moved LinkedParam/LinkedParamNode
Moved LinkedParam/LinkedParamNode from tvm::tir namespace to tvm
namespace
Change-Id: Ie3f0303bd4f7890c6d680268c91f2051977bc7f4
* Addressed upstream comments
Changed BindParams argument to Array<NDArray>
Removed 'name' argument from te.const
Switched to in-depth comparision of NDArrays in constant de-duplication
Removed extra final comma from NDArrayToTIR
Changed return type of ConstantAllocationSize to int64_t
Made link_param a tvm.testing.parameter for test_fuse_take and test_fuse_gather_nd
Change-Id: I4285099cc63756aa5ebe91a5bd207d4135499b41
* Removed unnecessary forward declaration
+linter
Change-Id: I2a6c0d1f97773aeb1ae3f458da252a22079ccdb1
* Constant extractor now is a separate pass
Change-Id: Ia4adca9d3315b26fbdc006ef7c115900c081e303
* Added forgotten file + unit test fix
Change-Id: Ice305f4fefd13fe95e97574e6d63ffeb664621df
* Changed to IRModule pass
Refactored ExtractPrimFuncConstants to IRModule pass.
deDup -> DeDup
Refactored logic of Applicator supplementary class
Change-Id: I6c120d175eb6790ba90f176c4f856bde8f0c7c94
* bugfix after rebasing
Change-Id: Ie3ee6ea2479476a30f486baef74f20070f117942
* -v -> -vv to have more debug information
Change-Id: I12c63731663b9c9ea574b9ed5cb17311ba3cf701
Co-authored-by: Giuseppe Rossini <giuseppe.rossini@arm.com>
* [TVMScript] Added unit tests demonstrating desired functionality
* [TVMScript] Implemented parsing of T.Ptr[...]
These can be generated when exporting to TVMscript, but were not
parsable after being generated.
* [TVMScript] Updated buffer_var printing
LetStmt and AllocateNode can both be used to generate handles that are
used in Buffer objects. In these cases, the Buffer declarations must
go after the handle declaration, not in the function header.
* Moved printing of var and buffer_decl into separate statements.
* Updated following @shingjan's review comments.
* fix parse strimm value in for annotations
* flatten buffer allow runtime.String attr value
* remove unused import
* rebase and ensure flattened attr order
* fix number of arguments
* make test clear
* Update tests/python/unittest/test_tvmscript_syntax_sugar.py
Co-authored-by: Wuwei Lin <vincentl13x@gmail.com>
* only tuple for now
Co-authored-by: Wuwei Lin <vincentl13x@gmail.com>
* Update doc building instructions and pin dependencies
This pins the dependencies for the docs and adds `pytest` as a
dependency which was missing when I built. Tested out the
requirements.txt with a fresh `ubuntu:focal` Docker image to verify that
the required depedencies work.
* Address comments, add Makefile for docs and add to instructions
* Use Python for running scripts
* Fix lint, add lint command
* Add option for cpu to only run the precheck, address comments
* Fix 'make doc' usage, add some -x's
* Fix bad condition on --cpu, add defaults for envs
* Fix another 'make doc'
* Fix for running on MacOS
Co-authored-by: driazati <driazati@users.noreply.github.com>
* [TIR][USMP] adding the pass to convert to pool offsets
This commit adds a transform pass that consumes
the planned pool allocations using memory planning algorithm
that convertes them to pool offsets.
* adds two test cases for a linear structure with two pools
* adds test case with a single pool for residual structures
Change-Id: I9d31e854461b5c21df72d1452120d286b96791c0
* [TIR][USMP] adding the pass to convert to pool offsets
* Adding a toggle to produce TIR that is TVMScript printable for unit
testing
* Fixing the unit tests
* Ensure deterministic pool variable ordering.
Change-Id: I317675df03327b0ebbf4ca074255384e63f07cd6
* [TIR][USMP] adding the pass to convert to pool offsets
Fixing the references after changes in the memory planning
algorithm.
Change-Id: Id7c22356fd5de43d10a2b4fc70e978af2c6d599d
* [TIR][USMP] adding the pass to convert to pool offsets
* fixing the lint
Change-Id: I7ff920b92d14a9919c930a4b35a2169c77a57dd1
* [TIR][USMP] adding the pass to convert to pool offsets
* removing unnecessary defitinitions
* remove global var map
* adding explaination for let bindings to pointer type
Change-Id: I31bd1a9f3057ee7f06252263565b0f75c51e6d13
* [TIR][USMP] adding the pass to convert to pool offsets
* rebase changes
* making imports absolute
* fixing typos and removing unnecesary lines
Change-Id: I4c94b9955b001513fecb39ca94f81b1ad99c7bfc
* [TIR][USMP] adding the pass to convert to pool offsets
* fixing typos
Change-Id: I42c557fd394aefdf8c2e825c4e88770eb0732f9b
* [TIR][USMP] Added buffer info extraction pass
This commit adds a pass that takes the main (call graph of operators)
TIR PrimFunc and each operators also as TIR PrimFunc. The pass will
traverse through all TIR PrimFunc starting the from main. Thereafter,
it will extract information from tir.allocates. Among the information,
the liveness conflicts are reported.
* Added test for a linear model
* Added test for parallel/serial mixed for loops
* Added test for a substructure of inception-style model.
* Exposed buffer_info creation to python
* Added member functions to update pool info
* Unit tests to cover functionality of buffer_info
Change-Id: I5e163ac3e83c830629a5d34ed4407c9962701c60
* [TIR][USMP] Added buffer info extraction pass
Swap key-value pairs of returned values of the buffer_info
extraction pass.
Change-Id: Ia4f7289592bc776ef6189a41a7891038751bf31f
* [TIR][USMP] Added buffer info extraction pass
Updating the USMP utility tests to include tests
that test creation of PoolInfo and PoolAllocation
Objects.
Change-Id: I5d349d0ffcac6b0160072d832dd9d5418699228e
* [TIR][USMP] Added buffer info extraction pass
* Removing the unnecessary header : include/tvm/tir/usmp/analysis.h
* Some nits and cleanup
Change-Id: Iac3ddd9428c56cd8ef49cf643e797bf6fdf4e97a
* [TIR][USMP] Added buffer info extraction pass
* Change the class data members to have a trailing underscore
Change-Id: I71809b3c73b0bc0cd133fad1392ae8c17c895ee4
* [TIR][USMP] Added buffer info extraction pass
Adding more documentation for data structures
and the approach
Change-Id: Ide2bfffaeff9add86853b6992017264e5d796299
* [TIR][USMP] Added buffer info extraction pass
* Added more documentation
* Added functionality to handle multiple calls
for the same PrimFunc with a test.
Change-Id: Ib7c27b3cf17f415067a224f1e57d8b928f4c7c6f
* [TIR][USMP] Added buffer info extraction pass
* Attaching targets to PrimFuncs in the util test case
Change-Id: I82960512659a346f6242b2b5789ec1120f8ea2cf
* add support for prevously uncovered cases
* remove PrimExpr import
* add exp test and mypy ignore
* disable ling too long
* resolve long line
* nit
* add dtype to unary ops
Krzysztof Parzyszek