[WIP][DYN] Support Tuning for Dynamic Workloads #394

rebased [TIR][Schedule] fix reorder/buffer_flatten & finish CPU demo (#59) [CPU DEMO] Update cpu gemm demo and fix bug (#58) * [TIR][Schedule] introduce parallel and fix bugs for cpu demo * [TIR][Schedule] update cpu demo * [TIR][Schedule] fix lint * [TIR][Schedule] fix rebased [TIR][Schedule] introduce reduction block and CPU demo (#53) * [TIR] reduction : split_reduction * [TIR] reduction : split_reduction * [TIR] reduction : fuse_reduction * [TIR] reduction : cpu demo * [TIR] reduction : fix * [TIR] reduction : pattern detect remains * [TIR] reduction : pattern detect remains * [TIR] reduction : pattern match done * [TIR] reduction : fix lint * [TIR] reduction : fix * [TIR] reduction : fix * [TIR] reduction : fix * [TIR] reduction : fix * [TIR] reduction : rebased * [TIR] reduction : rebased [TIR][Schedule] introduce cache_read cache_write (#54) * [TIR][Schedule] introduce cache_read cache_write * [TIR][Schedule] add more comments * [TIR][Schedule] fix problem and add comments * [TIR][Schedule] address comments [TIR] schedule: introduce vectorize, unroll, loop validation (#47) * [TIR] vectorize : basically complete * [TIR] vectorize&unroll : update comments&unroll * [TIR] vectorize&unroll : rebased * [TIR] vectorize, unroll, cpu_demo: done * [TIR] vectorize, unroll, cpu_demo: simplify * [TIR] vectorize, unroll, cpu_demo: fix * [TIR] reduction : rebased * [TIR] reduction : fix [TIR][Schedule] fix sref and scopes problem during replace and compute_at (#50) * [TIR][Schedule] fix sref and scopes problem during replace and compute_at * [TIR][Schedule] fix * [TIR][Schedule] fix [TIR][Refactor] move function to ScheduleNode [TIR] Schedule: introduce primitive compute_at (#36) * [TIR] Schedule: introduce primitive compute_at * [TIR] Schedule: address comments * [TIR] Schedule: address comments * [TIR] Schedule: address comments * [TIR] Schedule: add check to compute_at * [TIR] Schedule: address comments * [TIR] Schedule: address comments [TIR] Schedule: introduce primitive reorder (#37) * [Schedule] debug * [TIR] Schedule: reorder, loop type detect remains * [TIR] reorder complete * [TIR] reorder complete * [TIR] fix * [TIR] reorder : rebased complete * [TIR] reorder : fix container.h * [TIR] reorder : fix * [TIR] reorder : fix * [TIR] reorder : fix * [TIR] reorder : simplify * [TIR] reorder : simplify * [TIR] reorder : simplify * [TIR] reorder : fix * [TIR] reorder : fix * [TIR] reorder : rebased * [TIR] reorder : rebased rebase [TIR] Schedule: introduce BlockRealize and Block SRef reuse(#39) * [TIR] BlockRealize: schedule refactor * [TIR] BlockRealize: debug * [TIR] BlockRealize finish * [TIR] BlockRealize finish * [TIR] BlockRealize fix * [TIR] BlockRealize update test * [TIR] BlockRealize: add loop var reuse * [TIR] BlockRealize: add loop var reuse * [TIR] BlockRealize: fix * [TIR] BlockRealize: fix * [TIR] BlockRealize: fix * [TIR] BlockRealize: fix * [TIR] BlockRealize: fix * [TIR] BlockRealize: fix * [TIR] BlockRealize: fix * [TIR] BlockRealize: fix * [TIR] BlockRealize: fix * [TIR] BlockRealize: fix [TIR] compare for module (#38) * [TIR] compare for module * [TIR] fix * [TIR] fix * [TIR] fix * [TIR] fix * [TIR] fix * [TIR] fix [Hybrid] Module init [Hybrid] Module print [Hybrid] Module print with meta [Hybrid] adjust [Hybrid] finished but without lint and comment check [Hybrid] fix lint [Hybrid] comments [Hybrid] fix script decoration API [Hybrid] using IRModule [Hybrid] fix [Hybrid] adjust API [Hybrid] fix [Hybrid] fix [Hybrid] fix [Hybrid] fix symbol table, adjust API, introduce meta_mutator and resolve import issue [Hybrid] fix lint [TIR] introduce pass BufferFlatten (#32) * [TIR] introduce pass BufferFlatten * [Tir] add comments & remove old TeLower * [TIR] split GatherRegion and BufferFlatten to two Visitor/Mutator * [TIR] address comments: Only consider stmt scope * [TIR] BufferFlatten: address comments * [TIR] BufferFlatten: fold BlockFlattener into BufferFlattener * [TIR] BufferFlatten: add asserts * [TIR] BufferFlatten: use Equal in testcase * [TIR] Equal Pass: Enhanced the pass * [TIR] Equal Pass: add comments [Hybrid] refactor using Doc, introduce annotation, enhance parser (#28) * [Hybrid] refactor printer, enhance parser * [Hybrid] refactor * [Hybrid] fix * [Hybrid] fix * [Hybrid] fix namespace issue * [Hybrid] compare using Equal [TIR] rebased [TE] fix replace again and add primitive fuse and split (#27) * [TE] add: schedule primitive fuse * [TE] add: schedule primitive split * [TE] address comments: add IRSubstitueInScope and other minor fix * [TE] address comments: Enhance Equal api and fix split by nparts * [TE] address comments [Hybrid] introduce printer (#25) * [Hybrid] substitute Block with SeqStmt, change block() syntax * [Hybrid] add printer, type declare intrin * [Hybrid] refactor * [Hybrid] meta * [Hybrid] refactor * [Hybrid] macro [TE] fix replace (#23) * [TE] fix replace * [TE] fix replace: add more tests * [TE] fix replace: add more tests [TE] rebased [Hybrid] python syntax parser (#20) * [Hybrid] python syntax parser * [Hybrid] add a testcase * [Hybrid] improve comments and fix bugs * [Hybrid] improve comments, refactor __internal_assert, add new testcases * [Hybrid] improve error report message, refactor intrin * [Hybrid] separate ScopeEmitter from parser * [Hybrid] refactor type check * [Hybrid] refactor intrin * [Hybrid] refactor intrin, allow register external functions with argument type checking, add a testcase * [Hybrid] address comments, fix a bug in te/ir.h * [Hybrid] remove type check * [Hybrid] python syntax parser * [Hybrid] add a testcase * [Hybrid] improve comments and fix bugs * [Hybrid] improve comments, refactor __internal_assert, add new testcases * [Hybrid] improve error report message, refactor intrin * [Hybrid] separate ScopeEmitter from parser * [Hybrid] refactor type check * [Hybrid] refactor intrin * [Hybrid] refactor intrin, allow register external functions with argument type checking, add a testcase * [Hybrid] address comments, fix a bug in te/ir.h * [Hybrid] remove type check * [Hybrid] refactor intrin, scope_handler, special_stmt * [Hybrid] address comments * [Hybrid] clean code, improve error reporting & testcase * [Hybrid] clean code * [Hybrid] clean code [IR] introduce dependency graph and write map [TE] refactor and clean codebase [TE] refactor IR [TE] introduce schedule, dependency graph and support fuse and split (#17) * fix lint * introduce dependency graph * enable create schedule * support get axes * fix lint * revert Set * add schedule primitive fuse * address comment * support split [IR] Introduce SeqStmt add TeLower pass and enable to run Te IR (#15) * add function data structure add TeLower pass to transform Te to current IR enable to run Te IR * address comments * unify terminology TensorIR data structure init (#14) * init te data structure * finish printer and enhanced ir_builder * address the comments Co-authored-by: Bohan Hou <[email protected]>

* Re-implement: schedule creation, loop validation; Refactor: move creation code to schedule_create.cc, move validation code to schedule_validate.cc * Update * Add StmtSRef::GetStmt<SomeNode>() in preparation of removing DowncastPtr

* [AutoTIR] Add LoopTree * Add ReprPrinter for IteratorNode * Loops whose body isnt BlockRealize; More testcases * doc

…(#78) * [TIR][Hybrid] Extensible Parser & Life of Lowering & Concise Scoping * [TIR][HYBRID] fix several issues * [TIR][HYBRID] fix assert and let * [TIR][HYBRID] fix store * [TIR][Hybrid] improve hybrid type hierarchy * [TIR][Hybrid] improve comments * [TIR][Hybrid] update all the tests * [TIR][Hybrid] improve typing class comments * [TIR][Hybrid] improve comments * [TIR][Hybrid] improve comments * [TIR][Hybrid] move hybrid from tvm.tir to tvm namespace * [TIR][Hybrid] support single point TensorRegion sugar * [TIR][Hybrid] move string key to enum * [TIR][Hybrid] fix enum

* [TIR][HYBRID] Grid Iteration * [TIR][HYBRID] Grid Iteration fix * [TIR][HYBRID] fix bugs in schedule::replace

* [TIR][Schedule] finish blockize * [TIR][Schedule] finish native tensorize * [TIR][Schedule] finish tensorize * address comments * address comments

* [TIR][Hybrid] concise block declaration * [TIR][Hybrid] concise block declaration * [TIR][Hybrid] fix

* [Schedule] GPU support and part of execution scope * [Schedule] address comments * [Schedule] add scope validate

[TIR][HYBRID] fix [TIR][HYBRID] fix [TIR][HYBRID] fix [TIR][HYBRID] clean up (#108) * 【TIR][HYBRID] clean up * [TIR][Hybrid] allow to use expr.dtype

[TIR][Hybrid] fix [TIR][Hybrid] fix [TIR][Hybrid] fold attr&realize [TIR][Hybrid] concise thread env (#113) * [TIR][Hybrid] concise thread env * [TIR][Hybrid] concise thread env * [TIR][Hybrid] fix

* [Schedule] support compute inline * [Schedule] support higher-dimension tensorize

* [TIR] fix auto complete * [TIR] fix auto complete * [TIR] fix auto complete

* [Schedule] support compute inline * [Schedule] support higher-dimension tensorize * TensorCore support * TensorCore support * TensorCore support * Add testcase for exec_scope

* [TIR][Schedule] reverse_compute_at * [TIR][Schedule] reverse_compute_at: fixed * [TIR][Schedule] reverse_compute_at: fix

…efinement (#142)

* [TIR][AutoTensorize] init * [TIR][MetaSchedule] AutoTensorize * [TIR][MetaSchedule] AutoTensorize: dot product * [MetaSchedule] AutoTensorize: refactor * [MetaSchedule] embed auto tensorization into SearchRule&Analysis * [MetaSchedule] rebased * [MetaSchedule] rebased&fixed * [MetaSchedule] rebased * [MetaSchedule] rebased * [MetaSchedule] comment

* [TIR][Schedule] reverse_compute_inline * [TIR][Schedule] add reverse_compute_at into MetaSchedule * [TIR][Schedule] add reverse_compute_inline into MetaSchedule

* [TE] Generate Tir from TE compute * [TE] rebased * [TE] change create_tir to create_func * [TE] change create_tir to create_func Co-authored-by: spectrometerHBH <[email protected]>

* [Topi][Relay] Op Strategy for tir * [TE] Generate Tir from TE compute

* [TVMScript] metaprogramming: tensor shape invariant code * [TVMScript] metaprogramming: tensor shape invariant code * [TVMScript] metaprogramming: fix * [TVMScript] metaprogramming: fix * [TVMScript] metaprogramming: specialize_buffer * [TVMScript] adjust API * [TVMScript] adjust API * [TVMScript] move to global function * [TVMScript] refactor

* Add InstAttrs and make every inst_attr inherit from it * ... * use virtual

- AlwaysInline => InlinePureSpatial - Merge the following rules into Multi-Level-Tiling-And-Fusion - Multi-Level-Tiling - Fusion - Add-Cache-Write

* [TIR][Schedule] iter_var_map integrate & enhance [TIR][Schedule] iter_var_map integrate: bug fix [TIR][Schedule] iter_var_map integrate: bug fix [TIR][Schedule] predicate: init [TIR][Schedule] stride generalize [TIR][Schedule] iter_var_map: split/fuse simplify WIP WIP WIP WIP pass simple predicate test pass predicate tests pass compound predicate tests pass compound predicate tests pass compound predicate tests pass compound predicate tests subspace division init WIP WIP pass simple division test done, do some refactoring to make logic clear done done refactor refactor complete DONE. adress comments * adress comments * rebase

* WIP * basically done * done. * done.

* [TIR][Schedule] fix compute_at * [TIR][Schedule] fix compute_at

* [TIR][Schedule] rfactor : init * [TIR][Schedule] rfactor * [TIR][Schedule] rfactor * WIP * WIP * WIP * WIP * basically done * DONE. * DONE.

* [TIR][Schedule] sparse_dense init * [TIR][Schedule] sparse_dense: x86 * [TIR][Schedule] sparse_dense: x86, generic shape * set_store_predicate remaining * update

* add RecordToFile * integration test * record-to-file on integration tests * auto record loading

* [Bugfix] Cooperative fetching * minor refactors

* [MetaSchedule] Landing AutoTensorize: adjust cache_read/cache_write * [MetaSchedule] Landing AutoTensorize: fix postorderapply * fix get block * fix get block * add get_consumers/get_producers * remove get_only_consumer * pass cpu non-tensorized tests * fix the signature of cache_write * tir side refactor for cache_read * meta schedule side refactor for cache_read * search_rule refactor for cache_read * minor * make things running * configurable storage scope for cache_read/write Co-authored-by: spectrometerHBH <[email protected]>

* Change func => workload * filename => log_file * num_iterations => num_trials * Define Init method for SearchSpace, SearchStrategy, ProgramMeasurer * add SearchTask as an argument of postproc * remove target in verify_gpu_code * remove warp_size in rewrite_cuda_thread_bind * mark_parallelize_out's max_extent => max_jobs_per_core * minor

* Add printer * Match Ansor's FLOP counter * refactor record loading * refactor again * Check structural equality * print flops * minor

* [TIR][Schedule] pragma * [TIR][Schedule] pragma

* loop_type => block_type for a block * add SampleCategorical to schedule * add sample categorical * add unittest for sample categorical * add auto unroll search rule * change signature of mark_loop_type and mark_block_type * refactor annotation checking * prepare for pragma * post processor is done * add postproc on python side * add instruction: auto unroll * allow pragma to be assigned multiple times

…ead of Loop (#204)

…vectorize (#206)

* rename: trace => insts * Add TraceNode and Trace * pass the tests * remove trace-related functionality in the schedule class * remove serialization from schedule class * change the type of decision * Tweak the log format * minor name/sig changes * add AsPython * bugfix * fix them all

…on CUDA (#212)

…(#225)

* add simplify * remove simplify in auto complete

… Database (#241)

…eration (#244)

* find a bug for fused_nn_global_avg_pool2d * temp save * change tag to str * remove simplify in auto complete * correctly run resnet * change ds * fix * fix * change test doc * add decompose reduction in postproc and add docs

* [TIR][Schedule] reduction block * [TIR][Schedule] reduction: decompose * [TIR][Schedule] reduction: merge * [TIR][Schedule] reduction: blockize * [TIR][Schedule] reduction: remove ReduceStep * [TIR][Schedule] reduction: lowering * [TIR][Schedule] reduction: fix tests except rfactor * [TIR][Schedule] rfactor * [TIR][Schedule] rebased * [TIR][Schedule] lowering & allreduce init * [TIR][Schedule] allreduce done * [TIR][Schedule] fix annotation condition * [MetaSchedule][BugFix] remove the uses of ReduceStepNode * [TIR][Schedule] use `value.substr(0, 9) != "threadIdx"` * [TIR][Schedule] Add `const` and `&`. Co-authored-by: Siyuan Feng <[email protected]> * [TIR][Schedule] Change the way of string conversion. Co-authored-by: Siyuan Feng <[email protected]> * [MetaSchedule] Rewrite rules for decompose reduction. * [TIR][Schedule] remove redundant macros/calls * [MetaSchedule] only care the first appearance of reduction blocks Co-authored-by: Junru Shao <[email protected]> * [MetaSchedule][TIR][Schedule] Style change * [TIR][Schedule] Update signatures of "ReducerMatched" and "FromInitUpdate" * [TIR][Schedule] Make "ReducerMatched" and "FromInitUpdate" to be static class methods. * [TIR][Schedule] Refactor. Set AllreduceTransform to be a new pass. * [TIR][Schedule] Remove default argument * [TIR][Schedule] Fix comments * [TIR][Schedule] Update "IsReduction()" * Rename a variable & remove a TODO item * [MetaSchedule] Visit init in `VisitStmt(const BlockNode*)` * [TIR][Schedule] Allreduce style Co-authored-by: Ruihang Lai <[email protected]> Co-authored-by: Siyuan Feng <[email protected]> Co-authored-by: Junru Shao <[email protected]>

Co-authored-by: Wuwei Lin <[email protected]> Co-authored-by: Wuwei Lin <[email protected]>

…#260)

* [TIR] Add 'None' loop support in 'decompose_reduction' * [TIR] Add simplification and a unit test

…m) (#264)

* [Schedule] fix gpu demo * [Schedule] add double_buffer * [Schedule] fix script printer and parser * [Schedule] add test for double_buffer

* [MetaSchedule] Tensorization with wmma_load & wmma_store intrin * rebase * fix wmma_symc, wmma_fill, wmma_store * [ARITH] allow inject mapping for all data par block vars * [ARITH] fix bug * [ARITH] fix iter_affine_map unittest * [ARITH] fix unittests * fix * fix bugs * fix * fix Co-authored-by: spectrometerHBH <[email protected]>

…(#276)

(cherry picked from commit 47c6cae89d7d717a9cb1b4db783c0c8c1cd24fb5)

… with VerifyGPUCode in a better way (#289) (cherry picked from commit 86c0c51efcc3d77cb01032c63ada199a4d6dc86a)

(cherry picked from commit c63c4d561c99ed2bf9e8360883eee3c6b1dd7122)

(cherry picked from commit 3e0523f772d8a08239e4854fce720d58e9a2a6fb)

(cherry picked from commit 9727a85cdbd167da04a14ab2c241baabc4d51396)

(cherry picked from commit dcb54e6619a82d80770bbf9f53f85d0f8f2a9e37)

* [MetaSchedule] add MutateAutoUnroll and MutateParallel * [MetaSchedule] Fix comments * [MetaSchedule] Remove including "thread" head file * [MetaSchedule] Fix bugs in `MutateAutoUnroll` * [MetaSchedule] Fix comments. * [MetaSchedule] Rewrite Mutate * [MetaSchedule] use bool * [MetaSchedule] use std::atomic warned_num_cores_missing * [MetaSchedule] fix * [MetaSchedule] fix the use of atomic<bool> * [MetaSchedule] done * [MetaSchedule] change after refactor (cherry picked from commit 7c984b8cd0053ad21a9fd57b44c5e45ca8086902)

* TE workload for conv2d winograd * Update workload_te.py (cherry picked from commit ba48d5bf8dd2a05c6d9d272fa863357195722fd0)

rename: BufferAllocate => AllocBuffer rename: tag => name_hint remove: BufferAllocate move: exec_scope from block realize to block

* remove Loop and use For in tir * rename allocations => alloc_buffers rename Scope => BlockScope

* fix decompose reduction * rebased * fixed * fixed * fixed * rebased

* [TIR][Schedule] Replace VarCollector with PreOrderVisit * [TIR][Schedule] rfactor: fix one of the bugs * [TIR][Schedule] rfactor: add a new unittest * [TIR][Schedule] rfactor: reorganize * [TIR][Schedule] rfactor: reformat * [TIR][Schedule] add name of rfactor block * [TIR][Schedule] rewrite part of rfactor, and modify the tests * [TIR][Schedule] review write back part * [TIR][Schedule] reformat * [TIR][Schedule] perhaps have fixed the last bug * [TIR][Schedule] add test * [TIR][Schedule] style

* RuleSimplifyComputeWithConstTensor * tiling * fix * fix * new search rule * fix

…lowed by reduction block (#320) * Fix cache write with reduction consumer * debug * fix * case1 * comments * fix

…(#331) This reverts commit 34f89082116460d3073327c04ba559d367500602.

* [TIR][Schedule] Fix rfactor primitive again * [MetaSchedule] rewrite a reduction block only when it has reduction block var * [TIR][Schedule] fix

* [MetaSchedule] change after refactor * [MetaSchedule] fix typo * [MetaSchedule] add rule "AddRfactor" * [MetaSchedule] Move "NeedsInline" * [MetaSchedule] Update "NeedsRfactor" * [MetaSchedule] add reorder * [MetaSchedule] LoopNode -> ForNode after rebasing * [MetaSchedule] simply fix some ce after rebasing * [MetaSchedule] little refactor * [MetaSchedule] use `Schedule::Copy` to make copy * fix ce after rebasing * [MetaSchedule] use SamplePerfectTile. Successfully align! * [MetaSchedule] rename * [MetaSchedule] add the original schedule * [MetaSchedule] move helper functions * sorry... * [MetaSchedule] optimize imports * [MetaSchedule] refactor & document & use TVM_SREF_TO_BLOCK/FOR * [MetaSchedule] fix comments * [MetaSchedule] fix comments * [MetaSchedule] fix comments * [MetaSchedule] rebase

…(#339)

Co-authored-by: Junru Shao <[email protected]>

* [BugFix] Fix bug of compute_at * [BugFix] update doc * [BugFix] remove redundant log output * fix a tiny typo in tir/schedule/analysis.cc

…#351) * [BugFix][MetaSchedule] Add preserve_unit_loop to ComputeAtAttr & ReverseComputeAtAttr * [BugFix][MetaSchedule] Fix * [BugFix][MetaSchedule] Add Serialize/Deserialize

Co-authored-by: Ruihang Lai <[email protected]>

Co-authored-by: Ruihang Lai <[email protected]> Co-authored-by: Junru Shao <[email protected]>

* [BugFix] Fix `LCACollector` in current BufferFlatten * [BugFix] Update comment

* [TIR] Add primitive `SetScope` on C++ side * [TIR] Add primitive `SetScope` on Python side * [MetaSchedule] Add primitive `SetScope` on C++ side * [MetaSchedule] Add primitive `SetScope` on Python side * [MetaSchedule] Fix * [TIR][Schedule] Fix primitive bug * [TIR] Replace blocks as few as possible * [TIR] Add unit test * [TIR][MetaSchedule] Self-review and update. * [TIR] Remove `VisitStmt_(const BufferRealize* node)` * [TIR] Fuse both source&target of MatchBufferRegion

* add feature * add first version layout rewrite * add dense * add extract_tasks api and modify tests * add layout rewrite for batch matmul and conv2d winograd * final commit before rebase * pass compile * fix compile warning * fix * fix * remove useless fields and add doc * fix doc * change name * update doc for dispatcher.py * refactor relay_integration.py * refactor the rest python files * refactor * refactor * refactor * refactor * remove debug output * pass unittest * fix * remove unused fields * fix rebase bugs * style issues * apply black * disable vm compiler * fix bug * address comment

Co-authored-by: Ruihang Lai <[email protected]> Co-authored-by: Junru Shao <[email protected]>

Commits on Apr 23, 2021

Rebase.

ZihengJiang committed Apr 23, 2021

Configuration menu

View commit details

Copy full SHA for e51eab8

Browse repository at this point

Copy the full SHA

e51eab8 View commit details

Browse the repository at this point in the history

Commits on May 5, 2021

Fix bugs.

ZihengJiang committed May 5, 2021

Configuration menu

View commit details

Copy full SHA for b62a767

Browse repository at this point

Copy the full SHA

b62a767 View commit details

Browse the repository at this point in the history

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[WIP][DYN] Support Tuning for Dynamic Workloads #394

[WIP][DYN] Support Tuning for Dynamic Workloads #394

Commits on Mar 26, 2021

Commits on Mar 31, 2021

Commits on Apr 1, 2021

Commits on Apr 2, 2021

Commits on Apr 7, 2021

Commits on Apr 9, 2021

Commits on Apr 11, 2021

Commits on Apr 14, 2021

Commits on Apr 20, 2021

Commits on Apr 23, 2021

Commits on May 5, 2021

Commits on May 25, 2021

[WIP][DYN] Support Tuning for Dynamic Workloads #394

Are you sure you want to change the base?

[WIP][DYN] Support Tuning for Dynamic Workloads #394

Commits on Mar 26, 2021

Commits on Mar 31, 2021

Commits on Apr 1, 2021

Commits on Apr 2, 2021

Commits on Apr 7, 2021

Commits on Apr 9, 2021

Commits on Apr 11, 2021

Commits on Apr 14, 2021

Commits on Apr 20, 2021

Commits on Apr 23, 2021

Commits on May 5, 2021

Commits on May 25, 2021