[tvm] branch aluo/wombo-rebase-beforedyn-change-opaque-wuweiunetswitch created (now 7876fafd0b)

[an...@apache.org]

This is an automated email from the ASF dual-hosted git repository.

andrewzhaoluo pushed a change to branch aluo/wombo-rebase-beforedyn-change-opaque-wuweiunetswitch
in repository https://gitbox.apache.org/repos/asf/tvm.git


      at 7876fafd0b [TIR, Schedule] Update block flags and simplify predicate in Reverse-Compute-Inline

This branch includes the following new commits:

     new 4977f91b30 [VM][DMLC] Lower memory usage when loading and dumping weights (#13877)
     new c061e92695 --> opaque fusion
     new 7876fafd0b [TIR, Schedule] Update block flags and simplify predicate in Reverse-Compute-Inline

The 3 revisions listed above as "new" are entirely new to this
repository and will be described in separate emails.  The revisions
listed as "add" were already present in the repository and have only
been added to this reference.

[tvm] 02/03: --> opaque fusion

[an...@apache.org]

This is an automated email from the ASF dual-hosted git repository.

andrewzhaoluo pushed a commit to branch aluo/wombo-rebase-beforedyn-change-opaque-wuweiunetswitch
in repository https://gitbox.apache.org/repos/asf/tvm.git

commit c061e92695802b3d7c4ce444807470d0f54e19b0
Author: Andrew Zhao Luo <an...@gmail.com>
AuthorDate: Thu Feb 9 15:31:22 2023 -0800

    --> opaque fusion
---
 python/tvm/relay/op/_transform.py | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/python/tvm/relay/op/_transform.py b/python/tvm/relay/op/_transform.py
index d4e4a52783..64b701e9c5 100644
--- a/python/tvm/relay/op/_transform.py
+++ b/python/tvm/relay/op/_transform.py
@@ -95,11 +95,11 @@ _reg.register_injective_schedule("strided_set")
 
 # layout_transform
 _reg.register_injective_schedule("layout_transform")
-_reg.register_pattern("layout_transform", OpPattern.INJECTIVE)
+_reg.register_pattern("layout_transform", OpPattern.OPAQUE)
 _reg.register_injective_schedule("auto_scheduler_layout_transform")
-_reg.register_pattern("auto_scheduler_layout_transform", OpPattern.INJECTIVE)
+_reg.register_pattern("auto_scheduler_layout_transform", OpPattern.OPAQUE)
 _reg.register_injective_schedule("meta_schedule_layout_transform")
-_reg.register_pattern("meta_schedule_layout_transform", OpPattern.INJECTIVE)
+_reg.register_pattern("meta_schedule_layout_transform", OpPattern.OPAQUE)
 
 # argwhere
 _reg.register_strategy("argwhere", strategy.argwhere_strategy)

[tvm] 03/03: [TIR, Schedule] Update block flags and simplify predicate in Reverse-Compute-Inline

[an...@apache.org]

This is an automated email from the ASF dual-hosted git repository.

andrewzhaoluo pushed a commit to branch aluo/wombo-rebase-beforedyn-change-opaque-wuweiunetswitch
in repository https://gitbox.apache.org/repos/asf/tvm.git

commit 7876fafd0b6d5039d5b3261149f8e1c4118e0160
Author: Wuwei Lin <wu...@apache.org>
AuthorDate: Fri Feb 17 12:39:52 2023 -0800

    [TIR, Schedule] Update block flags and simplify predicate in Reverse-Compute-Inline
---
 src/tir/schedule/primitive/compute_inline.cc       |  8 +++++
 .../unittest/test_tir_schedule_compute_inline.py   | 37 ++++++++++++++++++++++
 2 files changed, 45 insertions(+)

diff --git a/src/tir/schedule/primitive/compute_inline.cc b/src/tir/schedule/primitive/compute_inline.cc
index d21149437f..d09042c3f5 100644
--- a/src/tir/schedule/primitive/compute_inline.cc
+++ b/src/tir/schedule/primitive/compute_inline.cc
@@ -651,6 +651,7 @@ class ReverseComputeInliner : public BaseInliner {
     // Substitute the producer block iters with the its bindings since the predicate in BlockRealize
     // should not contain the block iters
     predicate = Substitute(predicate, subst_map);
+    predicate = analyzer_.Simplify(predicate);
     return predicate;
   }
 
@@ -865,6 +866,13 @@ void ReverseComputeInlineImpl(ScheduleState self, const StmtSRef& consumer_block
     return;
   }
   self->Replace(scope_root_sref, tgt_stmt, inliner.block_reuse);
+  // Step 8. Update the cached flags
+  arith::Analyzer analyzer;
+  BlockInfo& block_info = self->block_info[producer_block_sref];
+  block_info.affine_binding = IsAffineBinding(
+      /*realize=*/GetBlockRealize(self, producer_block_sref),
+      /*loop_var_ranges=*/LoopDomainOfSRefTreePath(GetRef<StmtSRef>(producer_block_sref->parent)),
+      /*analyzer=*/&analyzer);
 }
 
 bool CanReverseComputeInline(const ScheduleState& self, const StmtSRef& block_sref) {
diff --git a/tests/python/unittest/test_tir_schedule_compute_inline.py b/tests/python/unittest/test_tir_schedule_compute_inline.py
index f9c5e22e97..b976deb3cd 100644
--- a/tests/python/unittest/test_tir_schedule_compute_inline.py
+++ b/tests/python/unittest/test_tir_schedule_compute_inline.py
@@ -611,6 +611,34 @@ def elementwise_overcomputed_producer_reverse_inlined(
             C[vi, vj] = A[vi, vj] * 2.0 + 1.0
 
 
+@T.prim_func
+def elementwise_overcomputed_producer_simplify_predicate(
+    A: T.Buffer((128, 128), "float32"), C: T.Buffer((127, 127), "float32")
+) -> None:
+    B = T.alloc_buffer((128, 128))
+    for i in T.grid(16384):
+        with T.block("B"):
+            vi = T.axis.spatial(128, i // 128)
+            vj = T.axis.spatial(128, i % 128)
+            B[vi, vj] = A[vi, vj] * 2.0
+    for i, j in T.grid(127, 127):
+        with T.block("C"):
+            cvi, cvj = T.axis.remap("SS", [i, j])
+            C[cvi, cvj] = B[cvi, cvj] + 1.0
+
+
+@T.prim_func
+def elementwise_overcomputed_producer_simplify_predicate_reverse_inlined(
+    A: T.Buffer((128, 128), "float32"), C: T.Buffer((127, 127), "float32")
+) -> None:
+    for i in T.grid(16384):
+        with T.block("B"):
+            vi = T.axis.spatial(128, i // 128)
+            vj = T.axis.spatial(128, i % 128)
+            T.where(i < 16255 and i % 128 < 127)
+            C[vi, vj] = A[vi, vj] * 2.0 + 1.0
+
+
 @T.prim_func
 def elementwise_producer_not_cover_consumer(
     A: T.Buffer[(128, 128), "float32"], D: T.Buffer[(256, 128), "float32"]
@@ -1025,6 +1053,15 @@ def test_reverse_compute_inline_overcomputed_producer(use_block_name):
     )
 
 
+def test_reverse_compute_inline_overcomputed_producer_simplify_predicate(use_block_name):
+    """Test reverse compute inline overcomputed producer where the predicate should be simplified"""
+    sch = tir.Schedule(elementwise_overcomputed_producer_simplify_predicate, debug_mask="all")
+    compute = "C" if use_block_name else sch.get_block("C")
+    sch.reverse_compute_inline(compute)
+    tvm.ir.assert_structural_equal(
+        elementwise_overcomputed_producer_simplify_predicate_reverse_inlined, sch.mod["main"]
+    )
+
 def test_reverse_compute_inline_error_producer_not_cover_consumer(use_block_name):
     """Test reverse compute inline failure when the inlined block iter domains are not covered by
     its producer

[tvm] 01/03: [VM][DMLC] Lower memory usage when loading and dumping weights (#13877)

[an...@apache.org]

This is an automated email from the ASF dual-hosted git repository.

andrewzhaoluo pushed a commit to branch aluo/wombo-rebase-beforedyn-change-opaque-wuweiunetswitch
in repository https://gitbox.apache.org/repos/asf/tvm.git

commit 4977f91b30b5a31fbc840c94b922ff120b7ccbda
Author: AndrewZhaoLuo <an...@gmail.com>
AuthorDate: Wed Feb 1 21:44:18 2023 -0800

    [VM][DMLC] Lower memory usage when loading and dumping weights (#13877)
    
    * initial commit
    
    * update additional use cases
    
    * typo
    
    * asf header, summary
    
    * clean up
    
    * lint
    
    * move code to src/runtime/file_utils.h
    
    * file utils is cool
---
 src/runtime/file_utils.h     | 37 +++++++++++++++++++++++++++++++++++++
 src/runtime/vm/executable.cc | 22 ++++++----------------
 2 files changed, 43 insertions(+), 16 deletions(-)

diff --git a/src/runtime/file_utils.h b/src/runtime/file_utils.h
index bd3bec9031..20806f5ff1 100644
--- a/src/runtime/file_utils.h
+++ b/src/runtime/file_utils.h
@@ -129,6 +129,43 @@ std::string SaveParams(const Map<String, NDArray>& params);
  * \param params Parameters to save.
  */
 void SaveParams(dmlc::Stream* strm, const Map<String, NDArray>& params);
+
+/*!
+ * \brief A dmlc stream which wraps standard file operations.
+ */
+struct SimpleBinaryFileStream : public dmlc::Stream {
+ public:
+  SimpleBinaryFileStream(const std::string& path, std::string mode) {
+    const char* fname = path.c_str();
+
+    CHECK(mode == "wb" || mode == "rb") << "Only allowed modes are 'wb' and 'rb'";
+    read_ = mode == "rb";
+    fp_ = std::fopen(fname, mode.c_str());
+    CHECK(fp_ != nullptr) << "Unable to open file " << path;
+  }
+  virtual ~SimpleBinaryFileStream(void) { this->Close(); }
+  virtual size_t Read(void* ptr, size_t size) {
+    CHECK(read_) << "File opened in write-mode, cannot read.";
+    CHECK(fp_ != nullptr) << "File is closed";
+    return std::fread(ptr, 1, size, fp_);
+  }
+  virtual void Write(const void* ptr, size_t size) {
+    CHECK(!read_) << "File opened in read-mode, cannot write.";
+    CHECK(fp_ != nullptr) << "File is closed";
+    CHECK(std::fwrite(ptr, 1, size, fp_) == size) << "SimpleBinaryFileStream.Write incomplete";
+  }
+  inline void Close(void) {
+    if (fp_ != nullptr) {
+      std::fclose(fp_);
+      fp_ = nullptr;
+    }
+  }
+
+ private:
+  std::FILE* fp_ = nullptr;
+  bool read_;
+};  // class SimpleBinaryFileStream
+
 }  // namespace runtime
 }  // namespace tvm
 #endif  // TVM_RUNTIME_FILE_UTILS_H_
diff --git a/src/runtime/vm/executable.cc b/src/runtime/vm/executable.cc
index 082ff05565..5696bc5314 100644
--- a/src/runtime/vm/executable.cc
+++ b/src/runtime/vm/executable.cc
@@ -344,10 +344,8 @@ void Executable::MoveLateBoundConstantsToStream(dmlc::Stream* stream, size_t byt
 }
 
 void Executable::MoveLateBoundConstantsToFile(const std::string& path, size_t byte_limit) {
-  std::string bytes;
-  dmlc::MemoryStringStream stream(&bytes);
+  tvm::runtime::SimpleBinaryFileStream stream(path, "wb");
   MoveLateBoundConstantsToStream(&stream, byte_limit);
-  SaveBinaryToFile(path, bytes);
 }
 
 void Executable::LoadLateBoundConstantsFromStream(dmlc::Stream* stream) {
@@ -381,9 +379,7 @@ void Executable::LoadLateBoundConstantsFromMap(Map<String, NDArray> map) {
 }
 
 void Executable::LoadLateBoundConstantsFromFile(const std::string& path) {
-  std::string bytes;
-  LoadBinaryFromFile(path, &bytes);
-  dmlc::MemoryStringStream stream(&bytes);
+  tvm::runtime::SimpleBinaryFileStream stream(path, "rb");
   LoadLateBoundConstantsFromStream(&stream);
 }
 
@@ -1063,22 +1059,16 @@ Module ExecutableLoadBinary(void* strm) {
 }
 
 void Executable::SaveToFile(const std::string& path, const std::string& format) {
-  std::string data;
-  dmlc::MemoryStringStream writer(&data);
-  dmlc::SeekStream* strm = &writer;
-  SaveToBinary(strm);
-  SaveBinaryToFile(path, data);
+  tvm::runtime::SimpleBinaryFileStream stream(path, "wb");
+  SaveToBinary(&stream);
 }
 
 TVM_REGISTER_GLOBAL("runtime.module.loadbinary_VMExecutable").set_body_typed(ExecutableLoadBinary);
 
 // Load module from module.
 Module ExecutableLoadFile(const std::string& file_name, const std::string& format) {
-  std::string data;
-  LoadBinaryFromFile(file_name, &data);
-  dmlc::MemoryStringStream reader(&data);
-  dmlc::Stream* strm = &reader;
-  auto exec = ExecutableLoadBinary(reinterpret_cast<void*>(strm));
+  tvm::runtime::SimpleBinaryFileStream stream(file_name, "rb");
+  auto exec = ExecutableLoadBinary(reinterpret_cast<void*>(&stream));
   return exec;
 }