From 5a16aa7c2c13a67e7b5c29131ee5b0533c721a94 Mon Sep 17 00:00:00 2001 From: FantasyVR Date: Wed, 9 Nov 2022 15:43:07 +0800 Subject: [PATCH 01/10] replace cpu sparse matrix builder's array with ndarray --- misc/test_sm.py | 24 ++++++++++++++++ python/taichi/lang/kernel_impl.py | 2 +- python/taichi/linalg/sparse_matrix.py | 11 ++++++++ taichi/program/sparse_matrix.cpp | 28 ++++++++++++++++--- taichi/program/sparse_matrix.h | 8 +++++- taichi/python/export_lang.cpp | 4 ++- .../llvm/runtime_module/internal_functions.h | 9 +++++- 7 files changed, 78 insertions(+), 8 deletions(-) create mode 100644 misc/test_sm.py diff --git a/misc/test_sm.py b/misc/test_sm.py new file mode 100644 index 0000000000000..96b5e76ea0338 --- /dev/null +++ b/misc/test_sm.py @@ -0,0 +1,24 @@ +import taichi as ti + +ti.init(arch=ti.x64, debug=True) + +n = 2 + +K = ti.linalg.SparseMatrixBuilder(n, + n, + max_num_triplets=100, + dtype=ti.f32, + storage_format='col_major') + +# K.test_ndarray() + + +@ti.kernel +def fill(A: ti.types.sparse_matrix_builder()): + for i in range(n): + A[i, i] += 2 + + +fill(K) + +K.print_ndarray_data() diff --git a/python/taichi/lang/kernel_impl.py b/python/taichi/lang/kernel_impl.py index 4cc8ac3330c33..618d9855458b7 100644 --- a/python/taichi/lang/kernel_impl.py +++ b/python/taichi/lang/kernel_impl.py @@ -673,7 +673,7 @@ def func__(*args): elif isinstance(needed, sparse_matrix_builder): # Pass only the base pointer of the ti.types.sparse_matrix_builder() argument launch_ctx.set_arg_uint(actual_argument_slot, - v._get_addr()) + v._get_ndarray_addr()) elif isinstance(needed, ndarray_type.NdarrayType) and isinstance( v, taichi.lang._ndarray.Ndarray): diff --git a/python/taichi/linalg/sparse_matrix.py b/python/taichi/linalg/sparse_matrix.py index eac2896644af6..9a62c65334f93 100644 --- a/python/taichi/linalg/sparse_matrix.py +++ b/python/taichi/linalg/sparse_matrix.py @@ -277,10 +277,21 @@ def _get_addr(self): """Get the address of the sparse matrix""" return self.ptr.get_addr() + def _get_ndarray_addr(self): + """Get the address of the ndarray""" + return self.ptr.get_ndarray_data_ptr() + def print_triplets(self): """Print the triplets stored in the builder""" self.ptr.print_triplets() + def print_ndarray_data(self): + """Print the ndarray data stored in the builder""" + self.ptr.print_ndarray_data() + + def test_ndarray(self): + self.ptr.test_ndarray() + def build(self, dtype=f32, _format='CSR'): """Create a sparse matrix using the triplets""" sm = self.ptr.build() diff --git a/taichi/program/sparse_matrix.cpp b/taichi/program/sparse_matrix.cpp index befd15951cf40..447fb70fff08c 100644 --- a/taichi/program/sparse_matrix.cpp +++ b/taichi/program/sparse_matrix.cpp @@ -65,16 +65,21 @@ SparseMatrixBuilder::SparseMatrixBuilder(int rows, int cols, int max_num_triplets, DataType dtype, - const std::string &storage_format) + const std::string &storage_format, + Program *prog) : rows_(rows), cols_(cols), max_num_triplets_(max_num_triplets), dtype_(dtype), - storage_format_(storage_format) { + storage_format_(storage_format), + prog_(prog) { auto element_size = data_type_size(dtype); TI_ASSERT((element_size == 4 || element_size == 8)); - data_base_ptr_ = - std::make_unique(max_num_triplets_ * 3 * element_size); + data_base_ptr_ndarray_ = std::make_unique( + prog_, dtype_, std::vector{3 * (int)max_num_triplets_}); + for (auto i = 0; i < 3 * max_num_triplets_; i++) { + data_base_ptr_ndarray_->write_int(std::vector{i}, 0); + } } template @@ -104,6 +109,21 @@ void SparseMatrixBuilder::print_triplets() { } } +intptr_t SparseMatrixBuilder::get_ndarray_data_ptr() const { + return prog_->get_ndarray_data_ptr_as_int(data_base_ptr_ndarray_.get()); +} + +void SparseMatrixBuilder::print_ndarray_data() { + auto ptr = get_ndarray_data_ptr(); + int32 *data = reinterpret_cast(ptr); + int32 num_numbers = data[0]; + fmt::print("number of numebrs: {} \n", num_numbers); + data += 1; + for (int i = 0; i < 3 * num_numbers; i++) { + fmt::print("ndarray data:[{}] = {}\n", i, data[i]); + } +} + template void SparseMatrixBuilder::build_template(std::unique_ptr &m) { using V = Eigen::Triplet; diff --git a/taichi/program/sparse_matrix.h b/taichi/program/sparse_matrix.h index 91e58f983d82c..a73b82e6710b6 100644 --- a/taichi/program/sparse_matrix.h +++ b/taichi/program/sparse_matrix.h @@ -19,9 +19,13 @@ class SparseMatrixBuilder { int cols, int max_num_triplets, DataType dtype, - const std::string &storage_format); + const std::string &storage_format, + Program *prog); void print_triplets(); + void print_ndarray_data(); + + intptr_t get_ndarray_data_ptr() const; std::unique_ptr build(); @@ -37,12 +41,14 @@ class SparseMatrixBuilder { private: uint64 num_triplets_{0}; std::unique_ptr data_base_ptr_{nullptr}; + std::unique_ptr data_base_ptr_ndarray_{nullptr}; int rows_{0}; int cols_{0}; uint64 max_num_triplets_{0}; bool built_{false}; DataType dtype_{PrimitiveType::f32}; std::string storage_format_{"col_major"}; + Program *prog_{nullptr}; }; class SparseMatrix { diff --git a/taichi/python/export_lang.cpp b/taichi/python/export_lang.cpp index 834a9a16ebc0a..ab1ee53ab0518 100644 --- a/taichi/python/export_lang.cpp +++ b/taichi/python/export_lang.cpp @@ -401,7 +401,7 @@ void export_lang(py::module &m) { TI_ERROR_IF(!arch_is_cpu(program->this_thread_config().arch), "SparseMatrix Builder only supports CPU for now."); return SparseMatrixBuilder(n, m, max_num_entries, dtype, - storage_format); + storage_format, program); }) .def("create_sparse_matrix", [](Program *program, int n, int m, DataType dtype, @@ -1204,6 +1204,8 @@ void export_lang(py::module &m) { // Sparse Matrix py::class_(m, "SparseMatrixBuilder") .def("print_triplets", &SparseMatrixBuilder::print_triplets) + .def("print_ndarray_data", &SparseMatrixBuilder::print_ndarray_data) + .def("get_ndarray_data_ptr", &SparseMatrixBuilder::get_ndarray_data_ptr) .def("build", &SparseMatrixBuilder::build) .def("get_addr", [](SparseMatrixBuilder *mat) { return uint64(mat); }); diff --git a/taichi/runtime/llvm/runtime_module/internal_functions.h b/taichi/runtime/llvm/runtime_module/internal_functions.h index 89c23fd998877..03713e1541c57 100644 --- a/taichi/runtime/llvm/runtime_module/internal_functions.h +++ b/taichi/runtime/llvm/runtime_module/internal_functions.h @@ -36,7 +36,14 @@ i32 insert_triplet_f32(RuntimeContext *context, int i, int j, float value) { - ATOMIC_INSERT(int32); + auto runtime = context->runtime; + auto base_ptr = reinterpret_cast(base_ptr_); + int32 *num_triplets = base_ptr; + auto data_base_ptr = base_ptr + 1; + auto triplet_id = atomic_add_i32(num_triplets, 1); + data_base_ptr[triplet_id * 3] = i; + data_base_ptr[triplet_id * 3 + 1] = j; + data_base_ptr[triplet_id * 3 + 2] = taichi_union_cast(value); return 0; } From 52029da8b8cf236a1522a9d6f86f86b35411a834 Mon Sep 17 00:00:00 2001 From: FantasyVR Date: Wed, 9 Nov 2022 21:05:37 +0800 Subject: [PATCH 02/10] sparse matrix builder supports cuda arch --- misc/sparse_matrix.py | 6 ++++-- misc/test_sm.py | 2 +- taichi/program/sparse_matrix.cpp | 24 ++++++++++++++---------- taichi/python/export_lang.cpp | 5 +++-- 4 files changed, 22 insertions(+), 15 deletions(-) diff --git a/misc/sparse_matrix.py b/misc/sparse_matrix.py index f96748104e618..6e0b9af80df44 100644 --- a/misc/sparse_matrix.py +++ b/misc/sparse_matrix.py @@ -1,6 +1,6 @@ import taichi as ti -ti.init(arch=ti.x64) +ti.init(arch=ti.x64, debug=True) n = 8 @@ -34,7 +34,9 @@ def fill(A: ti.types.sparse_matrix_builder(), fill(K, f, 3) print(">>>> K.print_triplets()") -K.print_triplets() +# K.print_triplets() + +K.print_ndarray_data() A = K.build() diff --git a/misc/test_sm.py b/misc/test_sm.py index 96b5e76ea0338..e4cad0d487c40 100644 --- a/misc/test_sm.py +++ b/misc/test_sm.py @@ -1,6 +1,6 @@ import taichi as ti -ti.init(arch=ti.x64, debug=True) +ti.init(arch=ti.cuda, debug=True) n = 2 diff --git a/taichi/program/sparse_matrix.cpp b/taichi/program/sparse_matrix.cpp index 447fb70fff08c..9832348e39ab7 100644 --- a/taichi/program/sparse_matrix.cpp +++ b/taichi/program/sparse_matrix.cpp @@ -114,13 +114,15 @@ intptr_t SparseMatrixBuilder::get_ndarray_data_ptr() const { } void SparseMatrixBuilder::print_ndarray_data() { - auto ptr = get_ndarray_data_ptr(); - int32 *data = reinterpret_cast(ptr); - int32 num_numbers = data[0]; - fmt::print("number of numebrs: {} \n", num_numbers); - data += 1; - for (int i = 0; i < 3 * num_numbers; i++) { - fmt::print("ndarray data:[{}] = {}\n", i, data[i]); + num_triplets_ = data_base_ptr_ndarray_->read_int(std::vector{0}); + fmt::print("n={}, m={}, num_triplets={} (max={})\n", rows_, cols_, + num_triplets_, max_num_triplets_); + for (int i = 0; i < num_triplets_; i++) { + auto idx = 3 * i + 1; + auto row = data_base_ptr_ndarray_->read_int(std::vector{idx}); + auto col = data_base_ptr_ndarray_->read_int(std::vector{idx + 1}); + auto val = data_base_ptr_ndarray_->read_int(std::vector{idx + 2}); + fmt::print("[{}, {}] = {}\n", row, col, val); } } @@ -128,10 +130,12 @@ template void SparseMatrixBuilder::build_template(std::unique_ptr &m) { using V = Eigen::Triplet; std::vector triplets; - T *data = reinterpret_cast(data_base_ptr_.get()); + auto ptr = get_ndarray_data_ptr(); + G *data = reinterpret_cast(ptr); + num_triplets_ = data[0]; + data += 1; for (int i = 0; i < num_triplets_; i++) { - triplets.push_back(V(((G *)data)[i * 3], ((G *)data)[i * 3 + 1], - taichi_union_cast(data[i * 3 + 2]))); + triplets.push_back(V(data[i * 3], data[i * 3 + 1], data[i * 3 + 2])); } m->build_triplets(static_cast(&triplets)); clear(); diff --git a/taichi/python/export_lang.cpp b/taichi/python/export_lang.cpp index ab1ee53ab0518..53d6e7f453184 100644 --- a/taichi/python/export_lang.cpp +++ b/taichi/python/export_lang.cpp @@ -398,8 +398,9 @@ void export_lang(py::module &m) { .def("create_sparse_matrix_builder", [](Program *program, int n, int m, uint64 max_num_entries, DataType dtype, const std::string &storage_format) { - TI_ERROR_IF(!arch_is_cpu(program->this_thread_config().arch), - "SparseMatrix Builder only supports CPU for now."); + TI_ERROR_IF(!arch_is_cpu(program->this_thread_config().arch) && + !arch_is_cuda(program->this_thread_config().arch), + "SparseMatrix only supports CPU and CUDA for now."); return SparseMatrixBuilder(n, m, max_num_entries, dtype, storage_format, program); }) From dd7f38c3bdd8fd64c99f08ec7a1b3b06ab271650 Mon Sep 17 00:00:00 2001 From: FantasyVR Date: Wed, 9 Nov 2022 21:38:54 +0800 Subject: [PATCH 03/10] simplify --- misc/sparse_matrix.py | 4 ++-- .../llvm/runtime_module/internal_functions.h | 22 +++++++++++-------- 2 files changed, 15 insertions(+), 11 deletions(-) diff --git a/misc/sparse_matrix.py b/misc/sparse_matrix.py index 6e0b9af80df44..2967b9143f16e 100644 --- a/misc/sparse_matrix.py +++ b/misc/sparse_matrix.py @@ -7,12 +7,12 @@ K = ti.linalg.SparseMatrixBuilder(n, n, max_num_triplets=100, - dtype=ti.f32, + dtype=ti.f64, storage_format='col_major') f = ti.linalg.SparseMatrixBuilder(n, 1, max_num_triplets=100, - dtype=ti.f32, + dtype=ti.f64, storage_format='col_major') diff --git a/taichi/runtime/llvm/runtime_module/internal_functions.h b/taichi/runtime/llvm/runtime_module/internal_functions.h index 03713e1541c57..bf5abdca1323f 100644 --- a/taichi/runtime/llvm/runtime_module/internal_functions.h +++ b/taichi/runtime/llvm/runtime_module/internal_functions.h @@ -20,6 +20,17 @@ data_base_ptr[triplet_id * 3 + 2] = taichi_union_cast(value); \ } while (0); +#define ATOMIC_INSERT_V1(T) \ + do { \ + auto base_ptr = reinterpret_cast(base_ptr_); \ + int##T *num_triplets = base_ptr; \ + auto data_base_ptr = base_ptr + 1; \ + auto triplet_id = atomic_add_i##T(num_triplets, 1); \ + data_base_ptr[triplet_id * 3] = i; \ + data_base_ptr[triplet_id * 3 + 1] = j; \ + data_base_ptr[triplet_id * 3 + 2] = taichi_union_cast(value); \ + } while (0); + i32 do_nothing(RuntimeContext *context) { return 0; } @@ -36,14 +47,7 @@ i32 insert_triplet_f32(RuntimeContext *context, int i, int j, float value) { - auto runtime = context->runtime; - auto base_ptr = reinterpret_cast(base_ptr_); - int32 *num_triplets = base_ptr; - auto data_base_ptr = base_ptr + 1; - auto triplet_id = atomic_add_i32(num_triplets, 1); - data_base_ptr[triplet_id * 3] = i; - data_base_ptr[triplet_id * 3 + 1] = j; - data_base_ptr[triplet_id * 3 + 2] = taichi_union_cast(value); + ATOMIC_INSERT_V1(32); return 0; } @@ -52,7 +56,7 @@ i32 insert_triplet_f64(RuntimeContext *context, int i, int j, float64 value) { - ATOMIC_INSERT(int64); + ATOMIC_INSERT_V1(64); return 0; } From 2f1d60b8837161b41f77897ed7309aa203e0ff94 Mon Sep 17 00:00:00 2001 From: FantasyVR Date: Thu, 10 Nov 2022 10:45:14 +0800 Subject: [PATCH 04/10] fix float insertion bug and clean code --- misc/sparse_matrix.py | 10 ++--- misc/test_sm.py | 29 +++++++------ python/taichi/linalg/sparse_matrix.py | 4 -- taichi/program/sparse_matrix.cpp | 33 ++++----------- taichi/program/sparse_matrix.h | 1 - taichi/python/export_lang.cpp | 1 - .../llvm/runtime_module/internal_functions.h | 37 +++++++--------- tests/python/test_sparse_linear_solver.py | 42 ++++++------------- tests/python/test_sparse_matrix.py | 26 ++++++------ tests/python/test_spmv.py | 6 +-- 10 files changed, 70 insertions(+), 119 deletions(-) diff --git a/misc/sparse_matrix.py b/misc/sparse_matrix.py index 2967b9143f16e..c030d1f209734 100644 --- a/misc/sparse_matrix.py +++ b/misc/sparse_matrix.py @@ -1,18 +1,18 @@ import taichi as ti -ti.init(arch=ti.x64, debug=True) +ti.init(arch=ti.x64, debug=True, offline_cache=False) n = 8 K = ti.linalg.SparseMatrixBuilder(n, n, max_num_triplets=100, - dtype=ti.f64, + dtype=ti.f32, storage_format='col_major') f = ti.linalg.SparseMatrixBuilder(n, 1, max_num_triplets=100, - dtype=ti.f64, + dtype=ti.f32, storage_format='col_major') @@ -34,9 +34,7 @@ def fill(A: ti.types.sparse_matrix_builder(), fill(K, f, 3) print(">>>> K.print_triplets()") -# K.print_triplets() - -K.print_ndarray_data() +K.print_triplets() A = K.build() diff --git a/misc/test_sm.py b/misc/test_sm.py index e4cad0d487c40..d90a75a0411d1 100644 --- a/misc/test_sm.py +++ b/misc/test_sm.py @@ -1,24 +1,23 @@ import taichi as ti -ti.init(arch=ti.cuda, debug=True) +ti.init(arch=ti.x64, debug=True, offline_cache=False) -n = 2 - -K = ti.linalg.SparseMatrixBuilder(n, - n, - max_num_triplets=100, - dtype=ti.f32, - storage_format='col_major') - -# K.test_ndarray() +n = 8 +triplets = ti.Vector.ndarray(n=3, dtype=ti.f32, shape=n) +A = ti.linalg.SparseMatrix(n=10, + m=10, + dtype=ti.f32, + storage_format='col_major') @ti.kernel -def fill(A: ti.types.sparse_matrix_builder()): +def fill(triplets: ti.types.ndarray()): for i in range(n): - A[i, i] += 2 - + triplet = ti.Vector([i, i, i], dt=ti.f32) + triplets[i] = triplet -fill(K) -K.print_ndarray_data() +fill(triplets) +A.build_from_ndarray(triplets) +for i in range(n): + assert A[i, i] == i diff --git a/python/taichi/linalg/sparse_matrix.py b/python/taichi/linalg/sparse_matrix.py index 9a62c65334f93..a17ff69fcfb05 100644 --- a/python/taichi/linalg/sparse_matrix.py +++ b/python/taichi/linalg/sparse_matrix.py @@ -285,10 +285,6 @@ def print_triplets(self): """Print the triplets stored in the builder""" self.ptr.print_triplets() - def print_ndarray_data(self): - """Print the ndarray data stored in the builder""" - self.ptr.print_ndarray_data() - def test_ndarray(self): self.ptr.test_ndarray() diff --git a/taichi/program/sparse_matrix.cpp b/taichi/program/sparse_matrix.cpp index 9832348e39ab7..d0893d839cbe0 100644 --- a/taichi/program/sparse_matrix.cpp +++ b/taichi/program/sparse_matrix.cpp @@ -77,9 +77,6 @@ SparseMatrixBuilder::SparseMatrixBuilder(int rows, TI_ASSERT((element_size == 4 || element_size == 8)); data_base_ptr_ndarray_ = std::make_unique( prog_, dtype_, std::vector{3 * (int)max_num_triplets_}); - for (auto i = 0; i < 3 * max_num_triplets_; i++) { - data_base_ptr_ndarray_->write_int(std::vector{i}, 0); - } } template @@ -95,25 +92,6 @@ void SparseMatrixBuilder::print_template() { } void SparseMatrixBuilder::print_triplets() { - auto element_size = data_type_size(dtype_); - switch (element_size) { - case 4: - print_template(); - break; - case 8: - print_template(); - break; - default: - TI_ERROR("Unsupported sparse matrix data type!"); - break; - } -} - -intptr_t SparseMatrixBuilder::get_ndarray_data_ptr() const { - return prog_->get_ndarray_data_ptr_as_int(data_base_ptr_ndarray_.get()); -} - -void SparseMatrixBuilder::print_ndarray_data() { num_triplets_ = data_base_ptr_ndarray_->read_int(std::vector{0}); fmt::print("n={}, m={}, num_triplets={} (max={})\n", rows_, cols_, num_triplets_, max_num_triplets_); @@ -121,11 +99,15 @@ void SparseMatrixBuilder::print_ndarray_data() { auto idx = 3 * i + 1; auto row = data_base_ptr_ndarray_->read_int(std::vector{idx}); auto col = data_base_ptr_ndarray_->read_int(std::vector{idx + 1}); - auto val = data_base_ptr_ndarray_->read_int(std::vector{idx + 2}); + auto val = data_base_ptr_ndarray_->read_float(std::vector{idx + 2}); fmt::print("[{}, {}] = {}\n", row, col, val); } } +intptr_t SparseMatrixBuilder::get_ndarray_data_ptr() const { + return prog_->get_ndarray_data_ptr_as_int(data_base_ptr_ndarray_.get()); +} + template void SparseMatrixBuilder::build_template(std::unique_ptr &m) { using V = Eigen::Triplet; @@ -135,7 +117,10 @@ void SparseMatrixBuilder::build_template(std::unique_ptr &m) { num_triplets_ = data[0]; data += 1; for (int i = 0; i < num_triplets_; i++) { - triplets.push_back(V(data[i * 3], data[i * 3 + 1], data[i * 3 + 2])); + triplets.push_back( + V(data[i * 3], data[i * 3 + 1], taichi_union_cast(data[i * 3 + 2]))); + fmt::print("({}, {}) val={}\n", data[i * 3], data[i * 3 + 1], + taichi_union_cast(data[i * 3 + 2])); } m->build_triplets(static_cast(&triplets)); clear(); diff --git a/taichi/program/sparse_matrix.h b/taichi/program/sparse_matrix.h index a73b82e6710b6..7d5aea228eabc 100644 --- a/taichi/program/sparse_matrix.h +++ b/taichi/program/sparse_matrix.h @@ -23,7 +23,6 @@ class SparseMatrixBuilder { Program *prog); void print_triplets(); - void print_ndarray_data(); intptr_t get_ndarray_data_ptr() const; diff --git a/taichi/python/export_lang.cpp b/taichi/python/export_lang.cpp index 53d6e7f453184..cdef137813131 100644 --- a/taichi/python/export_lang.cpp +++ b/taichi/python/export_lang.cpp @@ -1205,7 +1205,6 @@ void export_lang(py::module &m) { // Sparse Matrix py::class_(m, "SparseMatrixBuilder") .def("print_triplets", &SparseMatrixBuilder::print_triplets) - .def("print_ndarray_data", &SparseMatrixBuilder::print_ndarray_data) .def("get_ndarray_data_ptr", &SparseMatrixBuilder::get_ndarray_data_ptr) .def("build", &SparseMatrixBuilder::build) .def("get_addr", [](SparseMatrixBuilder *mat) { return uint64(mat); }); diff --git a/taichi/runtime/llvm/runtime_module/internal_functions.h b/taichi/runtime/llvm/runtime_module/internal_functions.h index bf5abdca1323f..7a74a4311ee2c 100644 --- a/taichi/runtime/llvm/runtime_module/internal_functions.h +++ b/taichi/runtime/llvm/runtime_module/internal_functions.h @@ -9,26 +9,19 @@ } \ } while (0) -#define ATOMIC_INSERT(T) \ - do { \ - auto base_ptr = (int64 *)base_ptr_; \ - int64 *num_triplets = base_ptr; \ - auto data_base_ptr = *(T **)(base_ptr + 1); \ - auto triplet_id = atomic_add_i64(num_triplets, 1); \ - data_base_ptr[triplet_id * 3] = i; \ - data_base_ptr[triplet_id * 3 + 1] = j; \ - data_base_ptr[triplet_id * 3 + 2] = taichi_union_cast(value); \ - } while (0); - -#define ATOMIC_INSERT_V1(T) \ - do { \ - auto base_ptr = reinterpret_cast(base_ptr_); \ - int##T *num_triplets = base_ptr; \ - auto data_base_ptr = base_ptr + 1; \ - auto triplet_id = atomic_add_i##T(num_triplets, 1); \ - data_base_ptr[triplet_id * 3] = i; \ - data_base_ptr[triplet_id * 3 + 1] = j; \ - data_base_ptr[triplet_id * 3 + 2] = taichi_union_cast(value); \ +#define ATOMIC_INSERT(T) \ + do { \ + auto runtime = context->runtime; \ + auto base_ptr = reinterpret_cast(base_ptr_); \ + int##T *num_triplets = base_ptr; \ + auto data_base_ptr = base_ptr + 1; \ + auto triplet_id = atomic_add_i##T(num_triplets, 1); \ + data_base_ptr[triplet_id * 3] = i; \ + data_base_ptr[triplet_id * 3 + 1] = j; \ + taichi_printf(runtime, "atomic inserted %d %d %f\t\n", i, j, value); \ + data_base_ptr[triplet_id * 3 + 2] = taichi_union_cast(value); \ + taichi_printf(runtime, "atomic inserted %d %d %f\t\n", i, j, \ + data_base_ptr[triplet_id * 3 + 2]); \ } while (0); i32 do_nothing(RuntimeContext *context) { @@ -47,7 +40,7 @@ i32 insert_triplet_f32(RuntimeContext *context, int i, int j, float value) { - ATOMIC_INSERT_V1(32); + ATOMIC_INSERT(32); return 0; } @@ -56,7 +49,7 @@ i32 insert_triplet_f64(RuntimeContext *context, int i, int j, float64 value) { - ATOMIC_INSERT_V1(64); + ATOMIC_INSERT(64); return 0; } diff --git a/tests/python/test_sparse_linear_solver.py b/tests/python/test_sparse_linear_solver.py index f31e7bd7df988..ef9ad4cfcecc5 100644 --- a/tests/python/test_sparse_linear_solver.py +++ b/tests/python/test_sparse_linear_solver.py @@ -4,40 +4,19 @@ import taichi as ti from tests import test_utils -""" -The symmetric positive definite matrix is created in matlab using the following script: - A = diag([1,2,3,4]); - OrthM = [1 0 1 0; -1 -2 0 1; 0 1 -1 0; 0, 1, 0 1]; - U = orth(OrthM); - Aarray = U * A * U'; - b = [1,2,3,4]'; - res = inv(A) * b; -""" # yapf: disable - -Aarray = np.array([[ - 2.73999501130921, 0.518002544441220, 0.745119303009342, 0.0508907745638859 -], [0.518002544441220, 1.45111665837647, 0.757997555750432, 0.290885785873098], - [ - 0.745119303009342, 0.757997555750432, 2.96711176987733, - -0.518002544441220 - ], - [ - 0.0508907745638859, 0.290885785873098, - -0.518002544441220, 2.84177656043698 - ]]) - -res = np.array([ - -0.0754984396447588, 0.469972700892492, 1.18527357933586, 1.57686870529319 -]) - @pytest.mark.parametrize("dtype", [ti.f32]) @pytest.mark.parametrize("solver_type", ["LLT", "LDLT", "LU"]) @pytest.mark.parametrize("ordering", ["AMD", "COLAMD"]) @test_utils.test(arch=ti.cpu) def test_sparse_LLT_solver(dtype, solver_type, ordering): - n = 4 - Abuilder = ti.linalg.SparseMatrixBuilder(n, n, max_num_triplets=100) + n = 10 + A = np.random.rand(n, n) + A_psd = np.dot(A, A.transpose()) + dtype = ti.float32 + solver_type = "LLT" + ordering = "AMD" + Abuilder = ti.linalg.SparseMatrixBuilder(n, n, max_num_triplets=300) b = ti.field(ti.f32, shape=n) @ti.kernel @@ -48,16 +27,19 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), for i in range(n): b[i] = i + 1 - fill(Abuilder, Aarray, b) + fill(Abuilder, A_psd, b) A = Abuilder.build() + print(A) solver = ti.linalg.SparseSolver(dtype=dtype, solver_type=solver_type, ordering=ordering) solver.analyze_pattern(A) solver.factorize(A) x = solver.solve(b) + + res = np.linalg.solve(A_psd, b.to_numpy()) for i in range(n): - assert x[i] == test_utils.approx(res[i]) + assert x[i] == test_utils.approx(res[i], rel=1.0) @test_utils.test(arch=ti.cuda) diff --git a/tests/python/test_sparse_matrix.py b/tests/python/test_sparse_matrix.py index 10cf8451a25f0..2c4f597f24ad3 100644 --- a/tests/python/test_sparse_matrix.py +++ b/tests/python/test_sparse_matrix.py @@ -9,7 +9,7 @@ (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_builder_deprecated_anno(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -34,7 +34,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_builder(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -59,7 +59,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_build_sparse_matrix_frome_ndarray(dtype, storage_format): n = 8 triplets = ti.Vector.ndarray(n=3, dtype=ti.f32, shape=n) @@ -85,7 +85,7 @@ def fill(triplets: ti.types.ndarray()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_shape(dtype, storage_format): n, m = 8, 9 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -108,7 +108,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_element_access(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -132,7 +132,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_element_modify(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -156,7 +156,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_addition(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -190,7 +190,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_subtraction(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -224,7 +224,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_scalar_multiplication(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -250,7 +250,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_transpose(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -276,7 +276,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_elementwise_multiplication(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -310,7 +310,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_multiplication(dtype, storage_format): n = 2 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -345,7 +345,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_nonsymmetric_multiplication(dtype, storage_format): n, k, m = 2, 3, 4 Abuilder = ti.linalg.SparseMatrixBuilder(n, diff --git a/tests/python/test_spmv.py b/tests/python/test_spmv.py index 638c5e29563c3..d31dae74727a3 100644 --- a/tests/python/test_spmv.py +++ b/tests/python/test_spmv.py @@ -8,7 +8,7 @@ (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_vector_multiplication1(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -37,7 +37,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), b: ti.template()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_vector_multiplication2(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -69,7 +69,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), b: ti.template()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu) +@test_utils.test(arch=ti.cpu, offline_cache=False) def test_sparse_matrix_vector_multiplication3(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, From 80a670b2ef226b1dfeda1dc58bd555a92fa81ba9 Mon Sep 17 00:00:00 2001 From: FantasyVR Date: Thu, 10 Nov 2022 10:57:38 +0800 Subject: [PATCH 05/10] clean --- taichi/program/sparse_matrix.cpp | 26 +++++-------------- taichi/program/sparse_matrix.h | 6 +---- .../llvm/runtime_module/internal_functions.h | 4 --- 3 files changed, 7 insertions(+), 29 deletions(-) diff --git a/taichi/program/sparse_matrix.cpp b/taichi/program/sparse_matrix.cpp index d0893d839cbe0..73fb1cc070747 100644 --- a/taichi/program/sparse_matrix.cpp +++ b/taichi/program/sparse_matrix.cpp @@ -75,37 +75,25 @@ SparseMatrixBuilder::SparseMatrixBuilder(int rows, prog_(prog) { auto element_size = data_type_size(dtype); TI_ASSERT((element_size == 4 || element_size == 8)); - data_base_ptr_ndarray_ = std::make_unique( + ndarray_data_base_ptr_ = std::make_unique( prog_, dtype_, std::vector{3 * (int)max_num_triplets_}); } -template -void SparseMatrixBuilder::print_template() { - fmt::print("n={}, m={}, num_triplets={} (max={})\n", rows_, cols_, - num_triplets_, max_num_triplets_); - T *data = reinterpret_cast(data_base_ptr_.get()); - for (int64 i = 0; i < num_triplets_; i++) { - fmt::print("({}, {}) val={}\n", ((G *)data)[i * 3], ((G *)data)[i * 3 + 1], - taichi_union_cast(data[i * 3 + 2])); - } - fmt::print("\n"); -} - void SparseMatrixBuilder::print_triplets() { - num_triplets_ = data_base_ptr_ndarray_->read_int(std::vector{0}); + num_triplets_ = ndarray_data_base_ptr_->read_int(std::vector{0}); fmt::print("n={}, m={}, num_triplets={} (max={})\n", rows_, cols_, num_triplets_, max_num_triplets_); for (int i = 0; i < num_triplets_; i++) { auto idx = 3 * i + 1; - auto row = data_base_ptr_ndarray_->read_int(std::vector{idx}); - auto col = data_base_ptr_ndarray_->read_int(std::vector{idx + 1}); - auto val = data_base_ptr_ndarray_->read_float(std::vector{idx + 2}); + auto row = ndarray_data_base_ptr_->read_int(std::vector{idx}); + auto col = ndarray_data_base_ptr_->read_int(std::vector{idx + 1}); + auto val = ndarray_data_base_ptr_->read_float(std::vector{idx + 2}); fmt::print("[{}, {}] = {}\n", row, col, val); } } intptr_t SparseMatrixBuilder::get_ndarray_data_ptr() const { - return prog_->get_ndarray_data_ptr_as_int(data_base_ptr_ndarray_.get()); + return prog_->get_ndarray_data_ptr_as_int(ndarray_data_base_ptr_.get()); } template @@ -119,8 +107,6 @@ void SparseMatrixBuilder::build_template(std::unique_ptr &m) { for (int i = 0; i < num_triplets_; i++) { triplets.push_back( V(data[i * 3], data[i * 3 + 1], taichi_union_cast(data[i * 3 + 2]))); - fmt::print("({}, {}) val={}\n", data[i * 3], data[i * 3 + 1], - taichi_union_cast(data[i * 3 + 2])); } m->build_triplets(static_cast(&triplets)); clear(); diff --git a/taichi/program/sparse_matrix.h b/taichi/program/sparse_matrix.h index 7d5aea228eabc..24cd247c32ea8 100644 --- a/taichi/program/sparse_matrix.h +++ b/taichi/program/sparse_matrix.h @@ -31,16 +31,12 @@ class SparseMatrixBuilder { void clear(); private: - template - void print_template(); - template void build_template(std::unique_ptr &); private: uint64 num_triplets_{0}; - std::unique_ptr data_base_ptr_{nullptr}; - std::unique_ptr data_base_ptr_ndarray_{nullptr}; + std::unique_ptr ndarray_data_base_ptr_{nullptr}; int rows_{0}; int cols_{0}; uint64 max_num_triplets_{0}; diff --git a/taichi/runtime/llvm/runtime_module/internal_functions.h b/taichi/runtime/llvm/runtime_module/internal_functions.h index 7a74a4311ee2c..ad66fa7cba3e1 100644 --- a/taichi/runtime/llvm/runtime_module/internal_functions.h +++ b/taichi/runtime/llvm/runtime_module/internal_functions.h @@ -11,17 +11,13 @@ #define ATOMIC_INSERT(T) \ do { \ - auto runtime = context->runtime; \ auto base_ptr = reinterpret_cast(base_ptr_); \ int##T *num_triplets = base_ptr; \ auto data_base_ptr = base_ptr + 1; \ auto triplet_id = atomic_add_i##T(num_triplets, 1); \ data_base_ptr[triplet_id * 3] = i; \ data_base_ptr[triplet_id * 3 + 1] = j; \ - taichi_printf(runtime, "atomic inserted %d %d %f\t\n", i, j, value); \ data_base_ptr[triplet_id * 3 + 2] = taichi_union_cast(value); \ - taichi_printf(runtime, "atomic inserted %d %d %f\t\n", i, j, \ - data_base_ptr[triplet_id * 3 + 2]); \ } while (0); i32 do_nothing(RuntimeContext *context) { From a3528914630b559afd3ccc07d55a4332e3631765 Mon Sep 17 00:00:00 2001 From: FantasyVR Date: Thu, 10 Nov 2022 11:08:41 +0800 Subject: [PATCH 06/10] rm --- misc/test_sm.py | 23 ----------------------- 1 file changed, 23 deletions(-) delete mode 100644 misc/test_sm.py diff --git a/misc/test_sm.py b/misc/test_sm.py deleted file mode 100644 index d90a75a0411d1..0000000000000 --- a/misc/test_sm.py +++ /dev/null @@ -1,23 +0,0 @@ -import taichi as ti - -ti.init(arch=ti.x64, debug=True, offline_cache=False) - -n = 8 -triplets = ti.Vector.ndarray(n=3, dtype=ti.f32, shape=n) -A = ti.linalg.SparseMatrix(n=10, - m=10, - dtype=ti.f32, - storage_format='col_major') - - -@ti.kernel -def fill(triplets: ti.types.ndarray()): - for i in range(n): - triplet = ti.Vector([i, i, i], dt=ti.f32) - triplets[i] = triplet - - -fill(triplets) -A.build_from_ndarray(triplets) -for i in range(n): - assert A[i, i] == i From 98627175aefe6982c5a1e1db71f9c0f1faf89810 Mon Sep 17 00:00:00 2001 From: FantasyVR Date: Thu, 10 Nov 2022 11:28:14 +0800 Subject: [PATCH 07/10] clean --- python/taichi/linalg/sparse_matrix.py | 3 --- 1 file changed, 3 deletions(-) diff --git a/python/taichi/linalg/sparse_matrix.py b/python/taichi/linalg/sparse_matrix.py index a17ff69fcfb05..9e0e3a835b6d9 100644 --- a/python/taichi/linalg/sparse_matrix.py +++ b/python/taichi/linalg/sparse_matrix.py @@ -285,9 +285,6 @@ def print_triplets(self): """Print the triplets stored in the builder""" self.ptr.print_triplets() - def test_ndarray(self): - self.ptr.test_ndarray() - def build(self, dtype=f32, _format='CSR'): """Create a sparse matrix using the triplets""" sm = self.ptr.build() From 177d486aa2450ef4ad48a04b20903011b1f6cf77 Mon Sep 17 00:00:00 2001 From: pengyu <6712304+FantasyVR@users.noreply.github.com> Date: Thu, 10 Nov 2022 21:20:13 +0800 Subject: [PATCH 08/10] Update taichi/program/sparse_matrix.cpp Co-authored-by: Yi Xu --- taichi/program/sparse_matrix.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/taichi/program/sparse_matrix.cpp b/taichi/program/sparse_matrix.cpp index 73fb1cc070747..0f37d79ad2745 100644 --- a/taichi/program/sparse_matrix.cpp +++ b/taichi/program/sparse_matrix.cpp @@ -76,7 +76,7 @@ SparseMatrixBuilder::SparseMatrixBuilder(int rows, auto element_size = data_type_size(dtype); TI_ASSERT((element_size == 4 || element_size == 8)); ndarray_data_base_ptr_ = std::make_unique( - prog_, dtype_, std::vector{3 * (int)max_num_triplets_}); + prog_, dtype_, std::vector{3 * (int)max_num_triplets_ + 1}); } void SparseMatrixBuilder::print_triplets() { From 26698e0c0203be08f5273ddaab84adee097c0e97 Mon Sep 17 00:00:00 2001 From: pengyu <6712304+FantasyVR@users.noreply.github.com> Date: Fri, 11 Nov 2022 14:46:45 +0800 Subject: [PATCH 09/10] rm --- tests/python/test_sparse_linear_solver.py | 3 --- 1 file changed, 3 deletions(-) diff --git a/tests/python/test_sparse_linear_solver.py b/tests/python/test_sparse_linear_solver.py index ef9ad4cfcecc5..cd48f0581b3af 100644 --- a/tests/python/test_sparse_linear_solver.py +++ b/tests/python/test_sparse_linear_solver.py @@ -13,9 +13,6 @@ def test_sparse_LLT_solver(dtype, solver_type, ordering): n = 10 A = np.random.rand(n, n) A_psd = np.dot(A, A.transpose()) - dtype = ti.float32 - solver_type = "LLT" - ordering = "AMD" Abuilder = ti.linalg.SparseMatrixBuilder(n, n, max_num_triplets=300) b = ti.field(ti.f32, shape=n) From b6f938f66e8ac05a526dfdf7c523ac8a1d769713 Mon Sep 17 00:00:00 2001 From: FantasyVR Date: Fri, 11 Nov 2022 21:44:18 +0800 Subject: [PATCH 10/10] enable offlinecach --- misc/sparse_matrix.py | 2 +- tests/python/test_sparse_matrix.py | 26 +++++++++++++------------- tests/python/test_spmv.py | 6 +++--- 3 files changed, 17 insertions(+), 17 deletions(-) diff --git a/misc/sparse_matrix.py b/misc/sparse_matrix.py index c030d1f209734..f96748104e618 100644 --- a/misc/sparse_matrix.py +++ b/misc/sparse_matrix.py @@ -1,6 +1,6 @@ import taichi as ti -ti.init(arch=ti.x64, debug=True, offline_cache=False) +ti.init(arch=ti.x64) n = 8 diff --git a/tests/python/test_sparse_matrix.py b/tests/python/test_sparse_matrix.py index 2c4f597f24ad3..10cf8451a25f0 100644 --- a/tests/python/test_sparse_matrix.py +++ b/tests/python/test_sparse_matrix.py @@ -9,7 +9,7 @@ (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_builder_deprecated_anno(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -34,7 +34,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_builder(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -59,7 +59,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_build_sparse_matrix_frome_ndarray(dtype, storage_format): n = 8 triplets = ti.Vector.ndarray(n=3, dtype=ti.f32, shape=n) @@ -85,7 +85,7 @@ def fill(triplets: ti.types.ndarray()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_shape(dtype, storage_format): n, m = 8, 9 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -108,7 +108,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_element_access(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -132,7 +132,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_element_modify(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -156,7 +156,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_addition(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -190,7 +190,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_subtraction(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -224,7 +224,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_scalar_multiplication(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -250,7 +250,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_transpose(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -276,7 +276,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_elementwise_multiplication(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -310,7 +310,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_multiplication(dtype, storage_format): n = 2 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -345,7 +345,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_nonsymmetric_multiplication(dtype, storage_format): n, k, m = 2, 3, 4 Abuilder = ti.linalg.SparseMatrixBuilder(n, diff --git a/tests/python/test_spmv.py b/tests/python/test_spmv.py index d31dae74727a3..638c5e29563c3 100644 --- a/tests/python/test_spmv.py +++ b/tests/python/test_spmv.py @@ -8,7 +8,7 @@ (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_vector_multiplication1(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -37,7 +37,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), b: ti.template()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_vector_multiplication2(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n, @@ -69,7 +69,7 @@ def fill(Abuilder: ti.types.sparse_matrix_builder(), b: ti.template()): (ti.f32, 'row_major'), (ti.f64, 'col_major'), (ti.f64, 'row_major')]) -@test_utils.test(arch=ti.cpu, offline_cache=False) +@test_utils.test(arch=ti.cpu) def test_sparse_matrix_vector_multiplication3(dtype, storage_format): n = 8 Abuilder = ti.linalg.SparseMatrixBuilder(n,