diff --git a/CMakeLists.txt b/CMakeLists.txt index 93aaab654c02e..bc18b82fd9070 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -8,7 +8,7 @@ project(taichi) SET(TI_VERSION_MAJOR 0) SET(TI_VERSION_MINOR 5) -SET(TI_VERSION_PATCH 0) +SET(TI_VERSION_PATCH 1) execute_process( WORKING_DIRECTORY ${CMAKE_SOURCE_DIR} diff --git a/README.md b/README.md index 1b495cb0d3e58..192e15ccbe83b 100644 --- a/README.md +++ b/README.md @@ -20,23 +20,13 @@ python3 -m pip install taichi-nightly-cuda-10-0 python3 -m pip install taichi-nightly-cuda-10-1 ``` +## [Contribution Guidelines](https://taichi.readthedocs.io/en/latest/contributor_guide.html) + || **Linux (CUDA)** | **OS X (10.14+)** | **Windows** | |:------|:-----|:-----|:-----| |**Build**|[![Build Status](http://f11.csail.mit.edu:8080/job/taichi/badge/icon)](http://f11.csail.mit.edu:8080/job/taichi/)| [![Build Status](https://travis-ci.com/taichi-dev/taichi.svg?branch=master)](https://travis-ci.com/taichi-dev/taichi) | [![Build status](https://ci.appveyor.com/api/projects/status/yxm0uniin8xty4j7/branch/master?svg=true)](https://ci.appveyor.com/project/yuanming-hu/taichi/branch/master)| |**PyPI**|[![Build Status](https://travis-ci.com/yuanming-hu/taichi-wheels-test.svg?branch=master)](https://travis-ci.com/yuanming-hu/taichi-wheels-test)|[![Build Status](https://travis-ci.com/yuanming-hu/taichi-wheels-test.svg?branch=master)](https://travis-ci.com/yuanming-hu/taichi-wheels-test)|[![Build status](https://ci.appveyor.com/api/projects/status/39ar9wa8yd49je7o?svg=true)](https://ci.appveyor.com/project/IteratorAdvance/taichi-wheels-test)| -## [Contribution Guidelines](https://taichi.readthedocs.io/en/latest/contributor_guide.html) - -## Related papers -- [**(SIGGRAPH Asia 2019) High-Performance Computation on Sparse Data Structures**](http://taichi.graphics/wp-content/uploads/2019/09/taichi_lang.pdf) [[Video]](https://youtu.be/wKw8LMF3Djo) [[BibTex]](https://raw.githubusercontent.com/yuanming-hu/taichi/master/misc/taichi_bibtex.txt) - - by *Yuanming Hu, Tzu-Mao Li, Luke Anderson, Jonathan Ragan-Kelley, and Frédo Durand* -- [**(ICLR 2020) Differentiable Programming for Physical Simulation**](https://arxiv.org/abs/1910.00935) [[Video]](https://www.youtube.com/watch?v=Z1xvAZve9aE) [[BibTex]](https://raw.githubusercontent.com/yuanming-hu/taichi/master/misc/difftaichi_bibtex.txt) [[Code]](https://github.com/yuanming-hu/difftaichi) - - by *Yuanming Hu, Luke Anderson, Tzu-Mao Li, Qi Sun, Nathan Carr, Jonathan Ragan-Kelley, and Frédo Durand* - -
- -
- ## Short-term goals - (Done) Fully implement the LLVM backend to replace the legacy source-to-source C++/CUDA backends (By Dec 2019) - The only missing features compared to the old source-to-source backends: @@ -46,6 +36,14 @@ python3 -m pip install taichi-nightly-cuda-10-1 - (WIP) Tune the performance of the LLVM backend to match that of the legacy source-to-source backends (Hopefully by mid Feb, 2020. Current progress: setting up/tuning for final benchmarks) ## Updates +- (Feb 16, 2020) v0.5.1 released + - Keyboard and mouse events supported in the GUI system. Check out [mpm128.py](https://github.com/taichi-dev/taichi/blob/master/examples/mpm128.py) for a interactive demo! (by **Yubin Peng [archibate] and Ye Kuang [k-ye]**) + - Basic algebraic simplification passes (by **Mingkuan Xu [xumingkuan]**) + - (For developers) `ti` (`ti.exe`) command supported on Windows after setting `%PATH%` correctly (by **Mingkuan Xu [xumingkuan]**) + - General power operator `x ** y` now supported in Taichi kernels (by **Yubin Peng [archibate]**) + - `.dense(...).pointer()` now abbreviated as `.pointer(...)`. `pointer` now stands for a dense pointer array. This leads to cleaner code and better performance. (by **Kenneth Lozes [KLozes]**) + - (Advanced struct-fors only) `for i in X` now iterates all child instances of `X` instead of `X` itself. Skip this if you only use `X=leaf node` such as `ti.f32/i32/Vector/Matrix`. + - Fixed cuda random number generator racing conditions - (Feb 14, 2020) **v0.5.0 released with a new Apple Metal GPU backend for Mac OS X users!** (by **Ye Kuang [k-ye]**) - Just initialize your program with `ti.init(..., arch=ti.metal)` and run Taichi on your Mac GPUs! - A few takeaways if you do want to use the Metal backend: @@ -103,3 +101,14 @@ python3 -m pip install taichi-nightly-cuda-10-1 - Doc updated - [Full changelog](changelog.md) + + +## Related papers +- [**(SIGGRAPH Asia 2019) High-Performance Computation on Sparse Data Structures**](http://taichi.graphics/wp-content/uploads/2019/09/taichi_lang.pdf) [[Video]](https://youtu.be/wKw8LMF3Djo) [[BibTex]](https://raw.githubusercontent.com/yuanming-hu/taichi/master/misc/taichi_bibtex.txt) + - by *Yuanming Hu, Tzu-Mao Li, Luke Anderson, Jonathan Ragan-Kelley, and Frédo Durand* +- [**(ICLR 2020) Differentiable Programming for Physical Simulation**](https://arxiv.org/abs/1910.00935) [[Video]](https://www.youtube.com/watch?v=Z1xvAZve9aE) [[BibTex]](https://raw.githubusercontent.com/yuanming-hu/taichi/master/misc/difftaichi_bibtex.txt) [[Code]](https://github.com/yuanming-hu/difftaichi) + - by *Yuanming Hu, Luke Anderson, Tzu-Mao Li, Qi Sun, Nathan Carr, Jonathan Ragan-Kelley, and Frédo Durand* + +
+ +
diff --git a/appveyor.yml b/appveyor.yml index 7fdcd2d22d168..2813e3619df61 100644 --- a/appveyor.yml +++ b/appveyor.yml @@ -33,6 +33,8 @@ skip_commits: - '.*' build_script: + - set TAICHI_REPO_DIR=C:\taichi + # - 'if %PYTHON% %TAICHI_REPO_DIR%\misc\appveyor_filter.py; then exit 0; fi' - cd C:\ - curl --retry 10 --retry-delay 5 https://github.com/yuanming-hu/taichi_assets/releases/download/llvm8/taichi-llvm-8.0.1-msvc2017.zip -LO - 7z x taichi-llvm-8.0.1-msvc2017.zip -otaichi_llvm @@ -40,7 +42,6 @@ build_script: - 7z x clang-7.0.1-win.zip -otaichi_clang - set PATH=C:\taichi_llvm\bin;%PATH%; - set PATH=C:\taichi_clang\bin;%PATH% - - set TAICHI_REPO_DIR=C:\taichi - set PYTHONPATH=%TAICHI_REPO_DIR%/python - set PATH=%TAICHI_REPO_DIR%\bin;%PATH% - clang --version diff --git a/benchmarks/fill_sparse.py b/benchmarks/fill_sparse.py index 0e4d72533f362..0ba80f6f8d31a 100644 --- a/benchmarks/fill_sparse.py +++ b/benchmarks/fill_sparse.py @@ -6,7 +6,7 @@ def benchmark_nested_struct(): @ti.layout def place(): - ti.root.dense(ti.ij, [N, N]).pointer().dense(ti.ij, [8, 8]).place(a) + ti.root.pointer(ti.ij, [N, N]).dense(ti.ij, [8, 8]).place(a) @ti.kernel def fill(): @@ -24,7 +24,7 @@ def benchmark_nested_struct_fill_and_clear(): @ti.layout def place(): - ti.root.dense(ti.ij, [N, N]).pointer().dense(ti.ij, [8, 8]).place(a) + ti.root.pointer(ti.ij, [N, N]).dense(ti.ij, [8, 8]).place(a) @ti.kernel def fill(): diff --git a/ci_setup.py b/ci_setup.py index 5b91d49d6c34b..dbe0637cc8748 100644 --- a/ci_setup.py +++ b/ci_setup.py @@ -47,7 +47,7 @@ def get_shell_rc_name(): def get_username(): if build_type == 'ci': - os.environ['TC_CI'] = '1' + os.environ['TI_CI'] = '1' username = 'travis' else: assert get_os_name() != 'win' diff --git a/cmake/TaichiCXXFlags.cmake b/cmake/TaichiCXXFlags.cmake index 588ca9ba6db76..b03a90593404c 100644 --- a/cmake/TaichiCXXFlags.cmake +++ b/cmake/TaichiCXXFlags.cmake @@ -1,6 +1,6 @@ message("Using C++ compiler: " ${CMAKE_CXX_COMPILER}) -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTC_ISE_NONE") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTI_ISE_NONE") option(BUILD_WITH_ADDRESS_SANITIZER "Build with clang address sanitizer" OFF) @@ -49,19 +49,19 @@ if (USE_STDCPP) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -stdlib=libc++") endif() -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTC_PASS_EXCEPTION_TO_PYTHON") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTI_PASS_EXCEPTION_TO_PYTHON") -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTC_INCLUDED") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTI_INCLUDED") -if ($ENV{TC_USE_DOUBLE}) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTC_USE_DOUBLE") +if ($ENV{TI_USE_DOUBLE}) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTI_USE_DOUBLE") message("Using float64 (double) precision as real") else() message("Using float32 (single) precision as real") endif() -if (TC_USE_MPI) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTC_USE_MPI") +if (TI_USE_MPI) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTI_USE_MPI") message("Using MPI") endif () diff --git a/cmake/TaichiCore.cmake b/cmake/TaichiCore.cmake index b831df4405943..803b7c79fc601 100644 --- a/cmake/TaichiCore.cmake +++ b/cmake/TaichiCore.cmake @@ -40,7 +40,7 @@ if (TI_WITH_CUDA) message("Building with CUDA ${CUDA_VERSION}") set(CUDA_ARCH 61) message("Found CUDA. Arch = ${CUDA_ARCH}") - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DCUDA_FOUND -DTI_WITH_CUDA") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTI_WITH_CUDA") if (MSVC) include_directories(${CUDA_TOOLKIT_ROOT_DIR}/include) target_link_libraries(${LIBRARY_NAME} ${CUDA_TOOLKIT_ROOT_DIR}/lib/x64/cudart.lib ${CUDA_TOOLKIT_ROOT_DIR}/lib/x64/cuda.lib) @@ -83,7 +83,7 @@ llvm_map_components_to_libnames(llvm_libs ) target_link_libraries(${LIBRARY_NAME} ${llvm_libs}) -if (CUDA_FOUND) +if (TI_WITH_CUDA) llvm_map_components_to_libnames(llvm_ptx_libs NVPTX) target_link_libraries(${LIBRARY_NAME} ${llvm_ptx_libs}) endif() diff --git a/docs/cpp_style.rst b/docs/cpp_style.rst index b9c8295bac726..6104b5b861124 100644 --- a/docs/cpp_style.rst +++ b/docs/cpp_style.rst @@ -5,7 +5,7 @@ Naming -------------------------------------------------------------------------- - Variable names should consist of lowercase words connected by underscores, e.g. ``llvm_context``. - Class and struct names should consist of words with first letters capitalized, e.g. ``CodegenLLVM``. -- Macros should be capital start with ``TC``, such as ``TC_INFO``, ``TC_IMPLEMENTATION``. +- Macros should be capital start with ``TC``, such as ``TI_INFO``, ``TI_IMPLEMENTATION``. - We do not encourage the use of macro, although there are cases where macros are inevitable. diff --git a/docs/utilities.rst b/docs/utilities.rst index 55642ac85f719..3e14f5fe53e56 100644 --- a/docs/utilities.rst +++ b/docs/utilities.rst @@ -47,45 +47,45 @@ Serialization The serialization module of taichi allows you to serialize/deserialize objects into/from binary strings. -You can use ``TC_IO`` macros to explicit define fields necessary in Taichi. +You can use ``TI_IO`` macros to explicit define fields necessary in Taichi. .. code-block:: cpp - // TC_IO_DEF + // TI_IO_DEF struct Particle { Vector3f position, velocity; real mass; string name; - TC_IO_DEF(position, velocity, mass, name); + TI_IO_DEF(position, velocity, mass, name); } - // TC_IO_DECL + // TI_IO_DECL struct Particle { Vector3f position, velocity; real mass; bool has_name string name; - TC_IO_DECL() { - TC_IO(position); - TC_IO(velocity); - TC_IO(mass); - TC_IO(has_name); + TI_IO_DECL() { + TI_IO(position); + TI_IO(velocity); + TI_IO(mass); + TI_IO(has_name); // More flexibility: if (has_name) { - TC_IO(name); + TI_IO(name); } } } - // TC_IO_DEF_VIRT(); + // TI_IO_DEF_VIRT(); Progress Notification ---------------------------------- -The taichi messager can send an email to ``$TC_MONITOR_EMAIL`` when the task finished or crashed. +The taichi messager can send an email to ``$TI_MONITOR_EMAIL`` when the task finished or crashed. To enable: .. code-block:: python diff --git a/docs/version b/docs/version index 8f0916f768f04..4b9fcbec101a6 100644 --- a/docs/version +++ b/docs/version @@ -1 +1 @@ -0.5.0 +0.5.1 diff --git a/examples/cpp/cnn.cpp b/examples/cpp/cnn.cpp index 457fdbe0b7adb..04e2faeb12d5b 100644 --- a/examples/cpp/cnn.cpp +++ b/examples/cpp/cnn.cpp @@ -3,7 +3,7 @@ #include #include -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN using namespace Tlang; @@ -12,10 +12,10 @@ constexpr int n = 256; constexpr int num_ch1 = 16, num_ch2 = 16; auto cnn = [](std::vector cli_param) { - TC_WARN( + TI_WARN( "After refactoring the Texture class from Taichi is removed. Need to " "read from the raw bunny binary."); - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED #if (0) CoreState::set_trigger_gdb_when_crash(true); auto param = parse_param(cli_param); @@ -24,7 +24,7 @@ auto cnn = [](std::vector cli_param) { auto cache_l1 = param.get("cache_l1", true); auto use_dense = param.get("use_dense", false); auto write_input_voxel = param.get("write_input", true); - TC_P(use_dense); + TI_P(use_dense); Program prog(gpu ? Arch::gpu : Arch::x86_64); prog.config.simplify_before_lower_access = opt; @@ -228,6 +228,6 @@ auto cnn = [](std::vector cli_param) { #endif #endif }; -TC_REGISTER_TASK(cnn); +TI_REGISTER_TASK(cnn); -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/examples/cpp/diff_conv.cpp b/examples/cpp/diff_conv.cpp index 46466ef5756e5..e35cee5ef5a8b 100644 --- a/examples/cpp/diff_conv.cpp +++ b/examples/cpp/diff_conv.cpp @@ -1,23 +1,23 @@ #include -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN auto diff_conv = [](const std::vector args) { int grid_resolution = 254; - TC_ASSERT(args.size() == 3); + TI_ASSERT(args.size() == 3); float th = std::stof(args[2]); - TC_P(th); + TI_P(th); auto f = fopen(args[0].c_str(), "rb"); int n = pow<3>(grid_resolution); - TC_ASSERT(f); + TI_ASSERT(f); std::vector ret1(n); trash(std::fread(ret1.data(), sizeof(float32), ret1.size(), f)); std::fclose(f); f = fopen(args[1].c_str(), "rb"); - TC_ASSERT(f); + TI_ASSERT(f); std::vector ret2(n); trash(std::fread(ret2.data(), sizeof(float32), ret2.size(), f)); std::fclose(f); @@ -53,19 +53,19 @@ auto diff_conv = [](const std::vector args) { // fprintf(stderr, "ret1:%f, ret2:%f\n", ret1[i], ret2[i]); //} } - TC_INFO("same {} {}%", counter[0], 100.0f * counter[0] / n); - TC_INFO("non zero same {} {}%", counter[0], + TI_INFO("same {} {}%", counter[0], 100.0f * counter[0] / n); + TI_INFO("non zero same {} {}%", counter[0], 100.0f * counter[1] / total_non_zero); - TC_P(sum1 / n); - TC_P(sum2 / n); - TC_P(sum1 / total_non_zero); - TC_P(sum2 / total_non_zero); - TC_P(max1); - TC_P(max2); - TC_P(non_zero1); - TC_P(non_zero2); + TI_P(sum1 / n); + TI_P(sum2 / n); + TI_P(sum1 / total_non_zero); + TI_P(sum2 / total_non_zero); + TI_P(max1); + TI_P(max2); + TI_P(non_zero1); + TI_P(non_zero2); }; -TC_REGISTER_TASK(diff_conv); +TI_REGISTER_TASK(diff_conv); -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/examples/cpp/fem.cpp b/examples/cpp/fem.cpp index 45ddc351d8e29..df1a89b5b08c6 100644 --- a/examples/cpp/fem.cpp +++ b/examples/cpp/fem.cpp @@ -5,7 +5,7 @@ #include #include -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN #include "fem_coeff.h" @@ -30,30 +30,30 @@ auto fem = [](std::vector cli_param) { auto param = parse_param(cli_param); bool gpu = param.get("gpu", false); - TC_P(gpu); + TI_P(gpu); bool vectorization = param.get("vec", true); - TC_P(vectorization); + TI_P(vectorization); int threads = param.get("threads", 8); - TC_P(threads); + TI_P(threads); bool use_cache = param.get("cache", true); - TC_P(use_cache); + TI_P(use_cache); bool compute_gt = param.get("compute_gt", false); - TC_P(compute_gt); + TI_P(compute_gt); Program prog(gpu ? Arch::gpu : Arch::x86_64); prog.config.simplify_before_lower_access = param.get("simp1", true); - TC_P(prog.config.simplify_before_lower_access); + TI_P(prog.config.simplify_before_lower_access); prog.config.lower_access = param.get("lower_access", true); - TC_P(prog.config.lower_access); + TI_P(prog.config.lower_access); prog.config.print_ir = param.get("print_ir", false); - TC_P(prog.config.print_ir); + TI_P(prog.config.print_ir); prog.config.simplify_after_lower_access = param.get("simp2", true); - TC_P(prog.config.simplify_after_lower_access); + TI_P(prog.config.simplify_after_lower_access); prog.config.attempt_vectorized_load_cpu = param.get("vec_load_cpu", true); - TC_P(prog.config.attempt_vectorized_load_cpu); + TI_P(prog.config.attempt_vectorized_load_cpu); bool use_pointer = param.get("use_pointer", true); - TC_P(use_pointer); + TI_P(use_pointer); bool block_soa = param.get("block_soa", true); - TC_P(block_soa); + TI_P(block_soa); prog.config.lazy_compilation = false; Vector x(DataType::f32, dim), r(DataType::f32, dim), p(DataType::f32, dim), @@ -261,7 +261,7 @@ auto fem = [](std::vector cli_param) { for (int i = 0; i < n - 1; i++) { for (int j = 0; j < n - 1; j++) { for (int k = 0; k < n - 1; k++) { - TC_ASSERT(!active[i][j][k]); + TI_ASSERT(!active[i][j][k]); } } } @@ -275,16 +275,16 @@ auto fem = [](std::vector cli_param) { auto old_rTr = sum.val(); for (int i = 0; i < 1000; i++) { - TC_P(i); + TI_P(i); compute_Ap(); sum.val() = 0; reduce_pAp(); auto pAp = sum.val(); // alpha = rTr / pTAp alpha.val() = old_rTr / pAp; - TC_P(old_rTr); - // TC_P(pAp); - // TC_P(alpha.val()); + TI_P(old_rTr); + // TI_P(pAp); + // TI_P(alpha.val()); // x = x + alpha p update_x(); // r = r - alpha Ap @@ -293,12 +293,12 @@ auto fem = [](std::vector cli_param) { sum.val() = 0; reduce_r(); auto new_rTr = sum.val(); - // TC_P(new_rTr); + // TI_P(new_rTr); if (new_rTr < 1e-5f) break; // beta = new rTr / old rTr beta.val() = new_rTr / old_rTr; - // TC_P(beta.val()); + // TI_P(beta.val()); // p = r + beta p update_p(); old_rTr = new_rTr; @@ -326,7 +326,7 @@ auto fem = [](std::vector cli_param) { } } } - TC_P(residual); + TI_P(residual); auto difference = 0.0f; auto difference_max = 0.0f; for (int i = 0; i < n; i++) { @@ -341,8 +341,8 @@ auto fem = [](std::vector cli_param) { } } } - TC_P(difference); - TC_P(difference_max); + TI_P(difference); + TI_P(difference_max); int gui_res = 512; GUI gui("FEM", Vector2i(gui_res + 200, gui_res), false); @@ -371,6 +371,6 @@ auto fem = [](std::vector cli_param) { gui.update(); } }; -TC_REGISTER_TASK(fem); +TI_REGISTER_TASK(fem); -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/examples/cpp/mgpcg.cpp b/examples/cpp/mgpcg.cpp index 1276e9cd58d32..b1a5268d654eb 100644 --- a/examples/cpp/mgpcg.cpp +++ b/examples/cpp/mgpcg.cpp @@ -4,7 +4,7 @@ #include #include -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN using namespace Tlang; @@ -19,15 +19,15 @@ auto mgpcg_poisson = [](std::vector cli_param) { int block_size = 8; int threads = param.get("threads", 8); - TC_P(threads); + TI_P(threads); bool vec_option = param.get("vec", true); int vec = vec_option ? block_size : 1; - TC_ASSERT(vec == 1 || vec == block_size); - TC_P(vec); + TI_ASSERT(vec == 1 || vec == block_size); + TI_P(vec); bool load_gt = param.get("load_gt", false); - TC_P(load_gt) + TI_P(load_gt) bool gpu = param.get("gpu", false); - TC_P(gpu) + TI_P(gpu) CoreState::set_trigger_gdb_when_crash(true); @@ -35,15 +35,15 @@ auto mgpcg_poisson = [](std::vector cli_param) { // prog.config.print_ir = true; prog.config.simplify_before_lower_access = param.get("simp1", true); - TC_P(prog.config.simplify_before_lower_access); + TI_P(prog.config.simplify_before_lower_access); prog.config.lower_access = param.get("lower_access", true); - TC_P(prog.config.lower_access); + TI_P(prog.config.lower_access); prog.config.print_ir = param.get("print_ir", false); - TC_P(prog.config.print_ir); + TI_P(prog.config.print_ir); prog.config.simplify_after_lower_access = param.get("simp2", true); - TC_P(prog.config.simplify_after_lower_access); + TI_P(prog.config.simplify_after_lower_access); prog.config.attempt_vectorized_load_cpu = param.get("vec_load_cpu", true); - TC_P(prog.config.attempt_vectorized_load_cpu); + TI_P(prog.config.attempt_vectorized_load_cpu); prog.config.lazy_compilation = false; @@ -278,7 +278,7 @@ auto mgpcg_poisson = [](std::vector cli_param) { sum.val() = 0; reduce_r(); auto initial_rTr = sum.val(); - TC_P(initial_rTr); + TI_P(initial_rTr); // r = b - Ax = b since x = 0 // p = r = r + 0 p @@ -291,16 +291,16 @@ auto mgpcg_poisson = [](std::vector cli_param) { // CG auto t = Time::get_time(); for (int i = 0; i < 400; i++) { - TC_P(i); + TI_P(i); compute_Ap(); sum.val() = 0; reduce_pAp(); auto pAp = sum.val(); // alpha = rTr / pTAp alpha.val() = old_zTr / pAp; - // TC_P(old_zTr); - // TC_P(pAp); - // TC_P(alpha.val()); + // TI_P(old_zTr); + // TI_P(pAp); + // TI_P(alpha.val()); // x = x + alpha p update_x(); // r = r - alpha Ap @@ -311,21 +311,21 @@ auto mgpcg_poisson = [](std::vector cli_param) { sum.val() = 0; reduce_zTr(); auto new_zTr = sum.val(); - // TC_P(new_zTr); + // TI_P(new_zTr); sum.val() = 0; reduce_r(); auto rTr = sum.val(); - TC_P(rTr); + TI_P(rTr); if (rTr < initial_rTr * 1e-12f) break; // beta = new rTr / old rTr beta.val() = new_zTr / old_zTr; - // TC_P(beta.val()); + // TI_P(beta.val()); // p = z + beta p update_p(); old_zTr = new_zTr; } - TC_P(Time::get_time() - t); + TI_P(Time::get_time() - t); get_current_program().profiler_print(); compute_Ap(); @@ -339,8 +339,8 @@ auto mgpcg_poisson = [](std::vector cli_param) { } } } - TC_P(residual); - // TC_P(difference_max); + TI_P(residual); + // TI_P(difference_max); std::vector ref_input(pow<3>(n / 2)); if (load_gt) { @@ -351,7 +351,7 @@ auto mgpcg_poisson = [](std::vector cli_param) { for (auto r : ref_input) { absmax = std::max(std::abs(absmax), r); } - TC_P(absmax); + TI_P(absmax); int gui_res = 512; GUI gui("MGPCG Poisson", Vector2i(gui_res + 200, gui_res), false); @@ -376,6 +376,6 @@ auto mgpcg_poisson = [](std::vector cli_param) { gui.update(); } }; -TC_REGISTER_TASK(mgpcg_poisson); +TI_REGISTER_TASK(mgpcg_poisson); -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/examples/cpp/mpm_benchmark.cpp b/examples/cpp/mpm_benchmark.cpp index 618492974513f..74d34a4ce40b1 100644 --- a/examples/cpp/mpm_benchmark.cpp +++ b/examples/cpp/mpm_benchmark.cpp @@ -6,7 +6,7 @@ #include #include -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN using namespace Tlang; @@ -15,18 +15,18 @@ auto mpm_benchmark = [](std::vector cli_param) { auto param = parse_param(cli_param); bool particle_soa = param.get("particle_soa", false); - TC_P(particle_soa); + TI_P(particle_soa); bool block_soa = param.get("block_soa", true); - TC_P(block_soa); + TI_P(block_soa); bool use_cache = param.get("use_cache", true); - TC_P(use_cache); + TI_P(use_cache); bool initial_reorder = param.get("initial_reorder", true); - TC_P(initial_reorder); + TI_P(initial_reorder); bool initial_shuffle = param.get("initial_shuffle", false); - TC_P(initial_shuffle); + TI_P(initial_shuffle); prog.config.lower_access = param.get("lower_access", false); int stagger = param.get("stagger", true); - TC_P(stagger); + TI_P(stagger); constexpr int dim = 3, n = 256, grid_block_size = 4, n_particles = 775196; const real dt = 1e-5_f * 256 / n, dx = 1.0_f / n, inv_dx = 1.0_f / dx; @@ -45,7 +45,7 @@ auto mpm_benchmark = [](std::vector cli_param) { p_x.resize(n_particles); std::vector benchmark_particles; auto f = fopen("dragon_particles.bin", "rb"); - TC_ASSERT_INFO(f, "./dragon_particles.bin not found"); + TI_ASSERT_INFO(f, "./dragon_particles.bin not found"); benchmark_particles.resize(n_particles * 3); if (std::fread(benchmark_particles.data(), sizeof(float), n_particles * 3, f)) { @@ -81,7 +81,7 @@ auto mpm_benchmark = [](std::vector cli_param) { place(particle_x(i)); for (int i = 0; i < dim; i++) place(particle_v(i)); - TC_ASSERT(n % grid_block_size == 0); + TI_ASSERT(n % grid_block_size == 0); auto &block = root.dense({i, j, k}, n / grid_block_size).pointer(); if (block_soa) { block.dense({i, j, k}, grid_block_size).place(grid_v(0)); @@ -253,7 +253,7 @@ auto mpm_benchmark = [](std::vector cli_param) { } prog.profiler_print(); auto ms_per_substep = (Time::get_time() - t) / 200 * 1000; - TC_P(ms_per_substep); + TI_P(ms_per_substep); }; #if (0) @@ -323,6 +323,6 @@ auto mpm_benchmark = [](std::vector cli_param) { } #endif }; -TC_REGISTER_TASK(mpm_benchmark); +TI_REGISTER_TASK(mpm_benchmark); -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/examples/cpp/mpm_full.cpp b/examples/cpp/mpm_full.cpp index bcfa7aa495ca4..7272bf62d59b3 100644 --- a/examples/cpp/mpm_full.cpp +++ b/examples/cpp/mpm_full.cpp @@ -5,7 +5,7 @@ #include #include "volume_renderer.h" -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN using namespace Tlang; @@ -37,16 +37,16 @@ auto mpm_full = [](std::vector cli_param) { int seeding_frames = param.get("seeding_frames", 300); real G = -1000.0f; - TC_P(total_frames); - TC_P(seeding_frames); - TC_P(dt); + TI_P(total_frames); + TI_P(seeding_frames); + TI_P(dt); dt = frame_dt / std::ceil(frame_dt / dt - 1e-5f); - TC_P(dt); - TC_P(frame_dt); + TI_P(dt); + TI_P(frame_dt); int visualize_interval = param.get("visualize_interval", 5); - TC_P(visualize_interval); + TI_P(visualize_interval); real ground_friction = param.get("ground_friction", 0.2f); - TC_P(ground_friction); + TI_P(ground_friction); auto particle_mass = 1.0_f, vol = 1.0_f; auto E = param.get("E", 1e4), nu = 0.3f; real mu_0 = E / (2 * (1 + nu)), lambda_0 = E * nu / ((1 + nu) * (1 - 2 * nu)); @@ -61,7 +61,7 @@ auto mpm_full = [](std::vector cli_param) { bool bbox = param.get("bbox", false); int scene = param.get("scene", 0); if (scene == 0) { - TC_INFO( + TI_INFO( "Scene: [1] ball smash [2] one jet [3] two static jets [4] two " "moving jets [5] particle curtain"); exit(0); @@ -78,12 +78,12 @@ auto mpm_full = [](std::vector cli_param) { } else if (material_name == "jelly") { material = MPMMaterial::jelly; } else { - TC_ERROR("Unknown material {}", material_name); + TI_ERROR("Unknown material {}", material_name); } - TC_P(material_name); - TC_P(bbox); - TC_P(scene); + TI_P(material_name); + TI_P(bbox); + TI_P(scene); Vector particle_x("x", f32, dim), particle_v("v", f32, dim); Global(particle_color, i32); @@ -105,9 +105,9 @@ auto mpm_full = [](std::vector cli_param) { if (benchmark_dragon) { n_particles = 775196; p_x.resize(n_particles); - TC_ASSERT(n_particles <= max_n_particles); + TI_ASSERT(n_particles <= max_n_particles); auto f = fopen("dragon_particles.bin", "rb"); - TC_ASSERT_INFO(f, "./dragon_particles.bin not found"); + TI_ASSERT_INFO(f, "./dragon_particles.bin not found"); benchmark_particles.resize(n_particles * 3); if (std::fread(benchmark_particles.data(), sizeof(float), n_particles * 3, f)) { @@ -141,7 +141,7 @@ auto mpm_full = [](std::vector cli_param) { } } - TC_ASSERT(n_particles <= max_n_particles); + TI_ASSERT(n_particles <= max_n_particles); auto sample_unit_sphere = [&] { Vector3 offset = Vector3::rand() - Vector3(0.5_f); @@ -186,7 +186,7 @@ auto mpm_full = [](std::vector cli_param) { place(particle_J); place(particle_color); - TC_ASSERT(n % grid_block_size == 0); + TI_ASSERT(n % grid_block_size == 0); auto &block = root.dense({i, j, k}, grid_n / 4 / grid_block_size) .pointer() .dense({i, j, k}, 4) @@ -215,7 +215,7 @@ auto mpm_full = [](std::vector cli_param) { // prog.visualize_layout("layout.tex"); - TC_ASSERT(bit::is_power_of_two(n)); + TI_ASSERT(bit::is_power_of_two(n)); Kernel(summarize).def([&] { BlockDim(512); @@ -560,15 +560,15 @@ auto mpm_full = [](std::vector cli_param) { ->clear_data_and_deactivate(); auto t = Time::get_time(); for (int f = 0; f < std::round(frame_dt / dt); f++) { - TC_PROFILE("p2g", p2g()); - TC_PROFILE("grid_op", grid_op()); - TC_PROFILE("g2p", g2p()); + TI_PROFILE("p2g", p2g()); + TI_PROFILE("grid_op", grid_op()); + TI_PROFILE("g2p", g2p()); } prog.profiler_print(); auto ms_per_substep = (Time::get_time() - t) / 200 * 1000; - TC_P(ms_per_substep); + TI_P(ms_per_substep); auto sec_per_frame = Time::get_time() - t; - TC_P(sec_per_frame); + TI_P(sec_per_frame); }; Kernel(set_renderer_volume).def([&] { @@ -584,7 +584,7 @@ auto mpm_full = [](std::vector cli_param) { for (int frame = 0; frame < total_frames; frame++) { float32 current_t = frame_dt * frame; if (frame < seeding_frames) { - TC_P(scene); + TI_P(scene); if (scene == 2) { int N = 10000; if (n_particles + N <= max_n_particles) { @@ -625,7 +625,7 @@ auto mpm_full = [](std::vector cli_param) { n_particles += N; } else if (scene == 5) { int N = 10000; - TC_P(N); + TI_P(N); if (n_particles + N <= max_n_particles) { for (int i = 0; i < N; i++) { Vector3 color(0.5, 0.6, 0.4); @@ -640,7 +640,7 @@ auto mpm_full = [](std::vector cli_param) { } } } - TC_P(n_particles) + TI_P(n_particles) simulate_frame(); // auto res = canvas.img.get_res(); @@ -656,11 +656,11 @@ auto mpm_full = [](std::vector cli_param) { auto ot = Time::get_time(); std::vector particle_data(max_n_particles * 7); summarize(); -#if defined(CUDA_FOUND) +#if defined(TI_WITH_CUDA) cudaMemcpy(particle_data.data(), &particle_buffer.val(0), particle_data.size() * sizeof(float), cudaMemcpyDeviceToHost); #else - TC_ERROR("No CUDA."); + TI_ERROR("No CUDA."); #endif for (int i = 0; i < n_particles; i++) { @@ -672,7 +672,7 @@ auto mpm_full = [](std::vector cli_param) { std::max(renderer.parameters.box_max[k], v); } } - TC_P(Time::get_time() - ot); + TI_P(Time::get_time() - ot); create_directories(fmt::format("final_particles/{}", output)); /* write_to_binary_file( @@ -714,9 +714,9 @@ auto mpm_full = [](std::vector cli_param) { print_profile_info(); } }; -TC_REGISTER_TASK(mpm_full); +TI_REGISTER_TASK(mpm_full); -TC_NAMESPACE_END +TI_NAMESPACE_END // demos: // two sand jets: diff --git a/examples/cpp/ray_march.cpp b/examples/cpp/ray_march.cpp index 86fd8e9a52c2d..bf697c0c4637e 100644 --- a/examples/cpp/ray_march.cpp +++ b/examples/cpp/ray_march.cpp @@ -137,6 +137,6 @@ auto ray_march = [] { gui.update(); } }; -TC_REGISTER_TASK(ray_march); +TI_REGISTER_TASK(ray_march); TLANG_NAMESPACE_END diff --git a/examples/cpp/smoke_renderer.cpp b/examples/cpp/smoke_renderer.cpp index d3e61645fd6c1..74db4fdfb9cd6 100644 --- a/examples/cpp/smoke_renderer.cpp +++ b/examples/cpp/smoke_renderer.cpp @@ -9,10 +9,10 @@ extern bool use_gui; auto smoke_renderer = [](std::vector cli_param_) { auto cli_param = parse_param(cli_param_); bool gpu = cli_param.get("gpu", true); - TC_P(gpu); + TI_P(gpu); Program prog(gpu ? Arch::gpu : Arch::x86_64); bool benchmark = true; // benchmark the bunny cloud against tungsten? - TC_ASSERT(benchmark); + TI_ASSERT(benchmark); // CoreState::set_trigger_gdb_when_crash(true); // prog.config.print_ir = true; float32 target_max_density = 50.f; @@ -37,7 +37,7 @@ auto smoke_renderer = [](std::vector cli_param_) { if (benchmark) { auto f = fopen("bunny_cloud.bin", "rb"); - TC_ASSERT_INFO(f, "./bunny_cloud.bin not found"); + TI_ASSERT_INFO(f, "./bunny_cloud.bin not found"); int box_sizes[3]{584, 576, 440}; int total_voxels = box_sizes[0] * box_sizes[1] * box_sizes[2]; std::vector density_field(total_voxels); @@ -50,7 +50,7 @@ auto smoke_renderer = [](std::vector cli_param_) { max_density = std::max(max_density, density_field[i]); } - TC_P(max_density); + TI_P(max_density); for (int i = 0; i < total_voxels; i++) { density_field[i] /= max_density; // normalize to 1 first @@ -125,13 +125,13 @@ auto smoke_renderer = [](std::vector cli_param_) { } } }; -TC_REGISTER_TASK(smoke_renderer); +TI_REGISTER_TASK(smoke_renderer); auto smoke_renderer_gui = [](std::vector cli_param) { use_gui = true; smoke_renderer(cli_param); }; -TC_REGISTER_TASK(smoke_renderer_gui); +TI_REGISTER_TASK(smoke_renderer_gui); TLANG_NAMESPACE_END diff --git a/examples/cpp/smoke_renderer.h b/examples/cpp/smoke_renderer.h index be7399f9ac63f..4673eee5211f7 100644 --- a/examples/cpp/smoke_renderer.h +++ b/examples/cpp/smoke_renderer.h @@ -29,8 +29,8 @@ class SmokeRenderer { depth_limit = param.get("depth_limit", 128); output_res = param.get("output_res", Vector2i(1024, 512)); - TC_ASSERT(bit::is_power_of_two(output_res.x)); - TC_ASSERT(bit::is_power_of_two(output_res.y)); + TI_ASSERT(bit::is_power_of_two(output_res.x)); + TI_ASSERT(bit::is_power_of_two(output_res.y)); sky_map_size = Vector2i(512, 128); n_sky_samples = 1024; @@ -373,12 +373,12 @@ class SmokeRenderer { std::FILE *f; if (use_sky_map) { f = fopen("sky_map.bin", "rb"); - TC_ASSERT_INFO(f, "./sky_map.bin not found"); + TI_ASSERT_INFO(f, "./sky_map.bin not found"); std::vector sky_map_data(sky_map_size.prod() * 3); std::fread(sky_map_data.data(), sizeof(uint32), sky_map_data.size(), f); f = fopen("sky_samples.bin", "rb"); - TC_ASSERT_INFO(f, "./sky_samples.bin not found"); + TI_ASSERT_INFO(f, "./sky_samples.bin not found"); std::vector sky_sample_data(n_sky_samples * 5); trash(std::fread(sky_sample_data.data(), sizeof(uint32), sky_sample_data.size(), f)); diff --git a/examples/cpp/volume_renderer.cpp b/examples/cpp/volume_renderer.cpp index b4f1608a93001..ea5fe7e1e8f3a 100644 --- a/examples/cpp/volume_renderer.cpp +++ b/examples/cpp/volume_renderer.cpp @@ -4,7 +4,7 @@ TLANG_NAMESPACE_BEGIN bool use_gui = false; TLANG_NAMESPACE_END -#if defined(CUDA_FOUND) +#if defined(TI_WITH_CUDA) #include TLANG_NAMESPACE_BEGIN @@ -13,9 +13,9 @@ auto volume_renderer = [](std::vector cli_param) { auto param = parse_param(cli_param); bool gpu = param.get("gpu", true); - TC_P(gpu); + TI_P(gpu); std::string fn = param.get("fn", "snow_density_256.bin"); - TC_P(fn); + TI_P(fn); CoreState::set_trigger_gdb_when_crash(true); Program prog(gpu ? Arch::gpu : Arch::x86_64); TRenderer renderer((Dict())); @@ -71,7 +71,7 @@ auto volume_renderer = [](std::vector cli_param) { last_voxel_level = voxel_level; std::vector particles; auto f = fopen(fn.c_str(), "rb"); - TC_WARN_IF(!f, "{} not found", fn); + TI_WARN_IF(!f, "{} not found", fn); if (!f) return; @@ -91,7 +91,7 @@ auto volume_renderer = [](std::vector cli_param) { max_density = std::max(max_density, density_field[i]); } - TC_P(max_density); + TI_P(max_density); for (int i = 0; i < pow<3>(grid_resolution); i++) { density_field[i] /= max_density; // normalize to 1 first @@ -135,11 +135,11 @@ auto volume_renderer = [](std::vector cli_param) { } else if (voxel_level == 4) { coarsening = 64; } else { - TC_ASSERT(false); + TI_ASSERT(false); } renderer.parameters.grid_resolution = 256 / coarsening; renderer.check_param_update(); - TC_P(particles.size()); + TI_P(particles.size()); rasterize(); } for (int d = 0; d < 3; d++) { @@ -267,17 +267,17 @@ auto volume_renderer = [](std::vector cli_param) { frame += video_step - 1; gui->canvas->img.write_as_image(fmt::format("gui/{:05d}.png", frame)); Time::sleep(0.03); - TC_P(Time::get_time() - ft); + TI_P(Time::get_time() - ft); } }; -TC_REGISTER_TASK(volume_renderer); +TI_REGISTER_TASK(volume_renderer); auto volume_renderer_gui = [](std::vector cli_param) { use_gui = true; volume_renderer(cli_param); }; -TC_REGISTER_TASK(volume_renderer_gui); +TI_REGISTER_TASK(volume_renderer_gui); TLANG_NAMESPACE_END #endif diff --git a/examples/cpp/volume_renderer.h b/examples/cpp/volume_renderer.h index 853a68eadee32..a21c404137f6a 100644 --- a/examples/cpp/volume_renderer.h +++ b/examples/cpp/volume_renderer.h @@ -67,7 +67,7 @@ class TRenderer { } void place_data() { - TC_ASSERT(output_res == Vector2i(1280, 720)); + TI_ASSERT(output_res == Vector2i(1280, 720)); root.dense(Index(0), 1024 * 1024).place(buffer(0), buffer(1), buffer(2)); root.dense(Indices(0, 1, 2), 4) @@ -669,14 +669,14 @@ class TRenderer { std::FILE *f; if (use_sky_map) { f = fopen("sky_map.bin", "rb"); - TC_ASSERT_INFO(f, "./sky_map.bin not found"); + TI_ASSERT_INFO(f, "./sky_map.bin not found"); std::vector sky_map_data(sky_map_size.prod() * 3); if (std::fread(sky_map_data.data(), sizeof(uint32), sky_map_data.size(), f)) { } f = fopen("sky_samples.bin", "rb"); - TC_ASSERT_INFO(f, "./sky_samples.bin not found"); + TI_ASSERT_INFO(f, "./sky_samples.bin not found"); std::vector sky_sample_data(n_sky_samples * 5); if (std::fread(sky_sample_data.data(), sizeof(uint32), (int)sky_sample_data.size(), f)) { diff --git a/examples/cpp/voxel_renderer.cpp b/examples/cpp/voxel_renderer.cpp index 99ffb0e5b7054..bf0122d6042fc 100644 --- a/examples/cpp/voxel_renderer.cpp +++ b/examples/cpp/voxel_renderer.cpp @@ -9,7 +9,7 @@ auto voxel_renderer = [](const std::vector ¶ms) { int n = 512; if (params.size() < 2) { - TC_INFO("Usage: ti voxel renderer filename.bin resolution"); + TI_INFO("Usage: ti voxel renderer filename.bin resolution"); exit(-1); } @@ -183,7 +183,7 @@ auto voxel_renderer = [](const std::vector ¶ms) { fn = fmt::format(params[0], frame); } auto f = fopen(fn.c_str(), "rb"); - TC_ERROR_UNLESS(f, "File {} not found", params[0]); + TI_ERROR_UNLESS(f, "File {} not found", params[0]); std::vector density_field(pow<3>(grid_resolution), 0); trash(std::fread(density_field.data(), sizeof(char), density_field.size(), f)); @@ -214,6 +214,6 @@ auto voxel_renderer = [](const std::vector ¶ms) { gui.canvas->img.write_as_image(fn + ".png"); } }; -TC_REGISTER_TASK(voxel_renderer); +TI_REGISTER_TASK(voxel_renderer); TLANG_NAMESPACE_END diff --git a/examples/mpm128.py b/examples/mpm128.py index 2e7f207080b9e..ffdfdacf3da4a 100644 --- a/examples/mpm128.py +++ b/examples/mpm128.py @@ -98,7 +98,7 @@ def reset(): Jp.fill(1) C.fill(0) -print("[Hint] Use WSAD/arrow keys to control gravity. Use left/right mouse bottons to attract/repel. (OS X not yet supported)") +print("[Hint] Use WSAD/arrow keys to control gravity. Use left/right mouse bottons to attract/repel.") gui = ti.GUI("Taichi MLS-MPM-128", res=512, background_color=0x112F41) reset() diff --git a/examples/renderer.py b/examples/renderer.py index ab56aa2a5f975..4fd8862328872 100644 --- a/examples/renderer.py +++ b/examples/renderer.py @@ -60,8 +60,7 @@ def buffers(): ti.root.dense(ti.ijk, 2).dense(ti.ijk, particle_grid_res // 8).dense( ti.ijk, 8).place(voxel_has_particle) - ti.root.dense(ti.ijk, 4).dense( - ti.ijk, particle_grid_res // 8).pointer().dense(ti.ijk, 8).dynamic( + ti.root.dense(ti.ijk, 4).pointer(ti.ijk, particle_grid_res // 8).dense(ti.ijk, 8).dynamic( ti.l, max_num_particles_per_cell, 512).place(pid) ti.root.dense(ti.l, max_num_particles).place(particle_x, particle_v, diff --git a/examples/taichi_sparse.py b/examples/taichi_sparse.py index c4bd2a53d04ef..be8abd2c2c03b 100644 --- a/examples/taichi_sparse.py +++ b/examples/taichi_sparse.py @@ -7,9 +7,9 @@ res = n + n // 4 + n // 16 + n // 64 img = ti.var(ti.f32, shape=(res, res)) -block1 = ti.root.dense(ti.ij, n // 64).pointer() -block2 = block1.dense(ti.ij, 4).pointer() -block3 = block2.dense(ti.ij, 4).pointer() +block1 = ti.root.pointer(ti.ij, n // 64) +block2 = block1.pointer(ti.ij, 4) +block3 = block2.pointer(ti.ij, 4) block3.dense(ti.ij, 4).place(x) @ti.func diff --git a/external/include/miniz.h b/external/include/miniz.h index 2ef07e3cd77a5..68cd8717d3153 100644 --- a/external/include/miniz.h +++ b/external/include/miniz.h @@ -4297,7 +4297,7 @@ static FILE *mz_freopen(const char *pPath, const char *pMode, FILE *pStream) #define MZ_FFLUSH fflush #define MZ_FREOPEN(f, m, s) freopen(f, m, s) #define MZ_DELETE_FILE remove -#elif defined(__GNUC__) && _LARGEFILE64_SOURCE && !defined(TC_PLATFORM_OSX) +#elif defined(__GNUC__) && _LARGEFILE64_SOURCE && !defined(TI_PLATFORM_OSX) #ifndef MINIZ_NO_TIME #include #endif diff --git a/external/include/stb_image.h b/external/include/stb_image.h index de04259da53b9..aa053383b4b8e 100644 --- a/external/include/stb_image.h +++ b/external/include/stb_image.h @@ -1,4 +1,4 @@ -#if defined(TC_IMAGE_IO) +#if defined(TI_IMAGE_IO) /* stb_image - v2.16 - public domain image loader - http://nothings.org/stb_image.h no warranty implied; use at your own risk diff --git a/external/include/stb_image_write.h b/external/include/stb_image_write.h index b53ae306d57b7..7b4c1c3df6dc0 100644 --- a/external/include/stb_image_write.h +++ b/external/include/stb_image_write.h @@ -1,4 +1,4 @@ -#if defined(TC_IMAGE_IO) +#if defined(TI_IMAGE_IO) /* stb_image_write - v1.07 - public domain - http://nothings.org/stb/stb_image_write.h writes out PNG/BMP/TGA/JPEG/HDR images to C stdio - Sean Barrett 2010-2015 no warranty implied; use at your own risk diff --git a/misc/amalgamate.py b/misc/amalgamate.py index d6dfe56f571a4..893dd1e76c130 100644 --- a/misc/amalgamate.py +++ b/misc/amalgamate.py @@ -66,7 +66,7 @@ def include(self, fn): if l == '#endif': protected = False continue - if l == '#if !defined(TC_AMALGAMATED)': + if l == '#if !defined(TI_AMALGAMATED)': protected = True continue match = re.search(include_template, l) @@ -109,9 +109,9 @@ def run(self): print( "// DO NOT EDIT MANUALLY, unless you know that you are doing.", file=self.output_f) - print("#define TC_INCLUDED", file=self.output_f) - print("#define TC_AMALGAMATED", file=self.output_f) - print("#define TC_ISE_NONE", file=self.output_f) + print("#define TI_INCLUDED", file=self.output_f) + print("#define TI_AMALGAMATED", file=self.output_f) + print("#define TI_ISE_NONE", file=self.output_f) for f in self.files: self.include(f) print("Included files:") @@ -130,7 +130,7 @@ def test(self): Vector4 v(21); auto x = v + v; fmt::print("{}\\n", x.x); - TC_P(x); + TI_P(x); } ''') t = time.time() diff --git a/misc/appveyor_filter.py b/misc/appveyor_filter.py new file mode 100644 index 0000000000000..289daaf781f34 --- /dev/null +++ b/misc/appveyor_filter.py @@ -0,0 +1,8 @@ +import sys +import os + +msg = os.environ["APPVEYOR_REPO_COMMIT_MESSAGE"] +if msg.startswith('[release]') or sys.version_info[1] == 6: + exit(0) # Build for this configuration (starts with '[release]', or python version is 3.6) +else: + exit(1) # Do not build this configuration. See appveyor.yml diff --git a/misc/obsolete_cpp_tests/allocator.cpp b/misc/obsolete_cpp_tests/allocator.cpp index 0e9184ae2ec2a..7cac37e5191c6 100644 --- a/misc/obsolete_cpp_tests/allocator.cpp +++ b/misc/obsolete_cpp_tests/allocator.cpp @@ -4,7 +4,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("gpu_gc_basics") { +TI_TEST("gpu_gc_basics") { for (auto arch : {Arch::gpu}) { int n = 32; Program prog(arch); @@ -27,26 +27,26 @@ TC_TEST("gpu_gc_basics") { })(); auto stat = x.parent().parent().snode()->stat(); - TC_CHECK(stat.num_resident_blocks == n - 1); + TI_CHECK(stat.num_resident_blocks == n - 1); for (int i = 0; i < n; i++) { for (int j = 0; j < i; j++) { - TC_CHECK(x.val(i, j) == i + j); + TI_CHECK(x.val(i, j) == i + j); } } x.parent().parent().snode()->clear_data_and_deactivate(); stat = x.parent().parent().snode()->stat(); - TC_CHECK(stat.num_resident_blocks == 0); - TC_CHECK(stat.num_recycled_blocks == 0); + TI_CHECK(stat.num_resident_blocks == 0); + TI_CHECK(stat.num_recycled_blocks == 0); for (int i = 0; i < n; i++) { for (int j = 0; j < i; j++) { - TC_CHECK(x.val(i, j) == 0); + TI_CHECK(x.val(i, j) == 0); } } } }; -TC_TEST("parallel_particle_sort") { +TI_TEST("parallel_particle_sort") { Program prog(Arch::gpu); CoreState::set_trigger_gdb_when_crash(true); @@ -76,7 +76,7 @@ TC_TEST("parallel_particle_sort") { p_x[i] = Vector3(0.5_f) + offset * 0.7f; } - TC_ASSERT(n_particles <= max_n_particles); + TI_ASSERT(n_particles <= max_n_particles); auto i = Index(0), j = Index(1), k = Index(2); auto p = Index(3); @@ -89,14 +89,14 @@ TC_TEST("parallel_particle_sort") { root.dense(i, max_n_particles).place(flag); - TC_ASSERT(n % grid_block_size == 0); + TI_ASSERT(n % grid_block_size == 0); root.dense({i, j, k}, n / grid_block_size) .pointer() .dense({i, j, k}, grid_block_size) .place(grid_m); }); - TC_ASSERT(bit::is_power_of_two(n)); + TI_ASSERT(bit::is_power_of_two(n)); Kernel(sort).def([&] { BlockDim(256); @@ -121,7 +121,7 @@ TC_TEST("parallel_particle_sort") { sort(); prog.synchronize(); for (int k = 0; k < max_n_particles; k++) { - TC_CHECK(flag.val(k) == 1); + TI_CHECK(flag.val(k) == 1); } auto stat = grid_m.parent().parent().snode()->stat(); int nb = stat.num_resident_blocks; @@ -129,14 +129,14 @@ TC_TEST("parallel_particle_sort") { last_nb = nb; } else { if (last_nb != nb) { - TC_P(i); + TI_P(i); } - TC_CHECK(last_nb == nb); + TI_CHECK(last_nb == nb); } } }; -TC_TEST("struct_for") { +TI_TEST("struct_for") { Program prog(Arch::gpu); CoreState::set_trigger_gdb_when_crash(true); @@ -166,7 +166,7 @@ TC_TEST("struct_for") { p_x[i] = Vector3(0.5_f) + offset * 0.7f; } - TC_ASSERT(n_particles <= max_n_particles); + TI_ASSERT(n_particles <= max_n_particles); auto i = Index(0), j = Index(1), k = Index(2); auto p = Index(3); @@ -201,8 +201,8 @@ TC_TEST("struct_for") { big_count += 1; } } - TC_CHECK(zero_count == 0); - TC_CHECK(big_count == 0); + TI_CHECK(zero_count == 0); + TI_CHECK(big_count == 0); } }; diff --git a/misc/obsolete_cpp_tests/atomics.cpp b/misc/obsolete_cpp_tests/atomics.cpp index 055f8e606555d..5183caf3fe523 100644 --- a/misc/obsolete_cpp_tests/atomics.cpp +++ b/misc/obsolete_cpp_tests/atomics.cpp @@ -4,7 +4,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("atomics") { +TI_TEST("atomics") { CoreState::set_trigger_gdb_when_crash(true); int n = 10000000; Program prog(Arch::x86_64); @@ -23,11 +23,11 @@ TC_TEST("atomics") { func(); - TC_CHECK(sum.val() == n); - TC_CHECK(fsum.val() == 0); + TI_CHECK(sum.val() == n); + TI_CHECK(fsum.val() == 0); }; -TC_TEST("atomics2") { +TI_TEST("atomics2") { CoreState::set_trigger_gdb_when_crash(true); int n = 1000; Program prog(Arch::x86_64); @@ -42,10 +42,10 @@ TC_TEST("atomics2") { func(); - TC_CHECK(fsum.val() == 1000); + TI_CHECK(fsum.val() == 1000); }; -TC_TEST("parallel_reduce") { +TI_TEST("parallel_reduce") { CoreState::set_trigger_gdb_when_crash(true); int n = 1024 * 1024 * 32; Program prog(Arch::x86_64); @@ -70,7 +70,7 @@ TC_TEST("parallel_reduce") { reduce(); prog.profiler_print(); - TC_CHECK(fsum.val() == (n / 2) * (n - 1) * 10); + TI_CHECK(fsum.val() == (n / 2) * (n - 1) * 10); }; TLANG_NAMESPACE_END diff --git a/misc/obsolete_cpp_tests/compiler_basics.cpp b/misc/obsolete_cpp_tests/compiler_basics.cpp index 5c91df12fac10..349924776b8fe 100644 --- a/misc/obsolete_cpp_tests/compiler_basics.cpp +++ b/misc/obsolete_cpp_tests/compiler_basics.cpp @@ -5,7 +5,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("compiler_linalg") { +TI_TEST("compiler_linalg") { CoreState::set_trigger_gdb_when_crash(true); Program prog(Arch::x86_64); @@ -33,7 +33,7 @@ TC_TEST("compiler_linalg") { })(); }; -TC_TEST("select") { +TI_TEST("select") { CoreState::set_trigger_gdb_when_crash(true); int n = 128; Program prog(Arch::x86_64); @@ -53,11 +53,11 @@ TC_TEST("select") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == (2 - i % 2) * i); + TI_CHECK(a.val(i) == (2 - i % 2) * i); } }; -TC_TEST("compiler_basics") { +TI_TEST("compiler_basics") { CoreState::set_trigger_gdb_when_crash(true); int n = 128; Program prog(Arch::x86_64); @@ -79,11 +79,11 @@ TC_TEST("compiler_basics") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == (2 - i % 2) * i); + TI_CHECK(a.val(i) == (2 - i % 2) * i); } }; -TC_TEST("simplify_access") { +TI_TEST("simplify_access") { CoreState::set_trigger_gdb_when_crash(true); int n = 128; Program prog(Arch::x86_64); @@ -96,7 +96,7 @@ TC_TEST("simplify_access") { kernel([&]() { For(a, [&](Expr i) { a[i] = b[i] + 1; }); })(); }; -TC_TEST("fancy_for") { +TI_TEST("fancy_for") { CoreState::set_trigger_gdb_when_crash(true); int n = 128; Program prog(Arch::x86_64); @@ -118,11 +118,11 @@ TC_TEST("fancy_for") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == (2 - i % 2) * i); + TI_CHECK(a.val(i) == (2 - i % 2) * i); } }; -TC_TEST("simd_if") { +TI_TEST("simd_if") { CoreState::set_trigger_gdb_when_crash(true); int n = 128; Program prog(Arch::x86_64); @@ -144,11 +144,11 @@ TC_TEST("simd_if") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == (2 - i % 2) * i); + TI_CHECK(a.val(i) == (2 - i % 2) * i); } }; -TC_TEST("simd_if2") { +TI_TEST("simd_if2") { CoreState::set_trigger_gdb_when_crash(true); int n = 32; Program prog(Arch::x86_64); @@ -170,7 +170,7 @@ TC_TEST("simd_if2") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == (1 + i % 3) * i); + TI_CHECK(a.val(i) == (1 + i % 3) * i); } }; @@ -206,9 +206,9 @@ auto test_circle = [] { gui.update(); } }; -TC_REGISTER_TASK(test_circle); +TI_REGISTER_TASK(test_circle); -TC_TEST("vectorize") { +TI_TEST("vectorize") { CoreState::set_trigger_gdb_when_crash(true); int n = 128; Program prog(Arch::x86_64); @@ -223,11 +223,11 @@ TC_TEST("vectorize") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == i); + TI_CHECK(a.val(i) == i); } }; -TC_TEST("rand") { +TI_TEST("rand") { CoreState::set_trigger_gdb_when_crash(true); int n = 4; Program prog(Arch::x86_64); @@ -239,7 +239,7 @@ TC_TEST("rand") { kernel([&]() { For(0, n, [&](Expr i) { Print(Rand()); }); })(); }; -TC_TEST("while") { +TI_TEST("while") { CoreState::set_trigger_gdb_when_crash(true); int n = 4096; Program prog(Arch::x86_64); @@ -262,11 +262,11 @@ TC_TEST("while") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == (i - 1) * i / 2); + TI_CHECK(a.val(i) == (i - 1) * i / 2); } }; -TC_TEST("slp") { +TI_TEST("slp") { CoreState::set_trigger_gdb_when_crash(true); int n = 16; Program prog(Arch::x86_64); @@ -290,14 +290,14 @@ TC_TEST("slp") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == 1); - TC_CHECK(b.val(i) == 2); - TC_CHECK(c.val(i) == 3); - TC_CHECK(d.val(i) == 4); + TI_CHECK(a.val(i) == 1); + TI_CHECK(b.val(i) == 2); + TI_CHECK(c.val(i) == 3); + TI_CHECK(d.val(i) == 4); } }; -TC_TEST("slp1") { +TI_TEST("slp1") { CoreState::set_trigger_gdb_when_crash(true); int n = 16; for (auto slp1 : {true, false}) { @@ -322,13 +322,13 @@ TC_TEST("slp1") { for (int i = 0; i < n; i++) { for (int d = 0; d < 4; d++) { - TC_CHECK(grid(d).val(i) == d); + TI_CHECK(grid(d).val(i) == d); } } } }; -TC_TEST("slp2") { +TI_TEST("slp2") { CoreState::set_trigger_gdb_when_crash(true); int n = 16; Program prog(Arch::x86_64); @@ -348,12 +348,12 @@ TC_TEST("slp2") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == 1 + i * 7); - TC_CHECK(b.val(i) == 2 + i * 9); + TI_CHECK(a.val(i) == 1 + i * 7); + TI_CHECK(b.val(i) == 2 + i * 9); } }; -TC_TEST("slp3") { +TI_TEST("slp3") { CoreState::set_trigger_gdb_when_crash(true); int n = 16; Program prog(Arch::x86_64); @@ -375,12 +375,12 @@ TC_TEST("slp3") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == 1 + i * 7); - TC_CHECK(b.val(i) == 2 + i * 9); + TI_CHECK(a.val(i) == 1 + i * 7); + TI_CHECK(b.val(i) == 2 + i * 9); } }; -TC_TEST("slpmatvecmul") { +TI_TEST("slpmatvecmul") { CoreState::set_trigger_gdb_when_crash(true); int n = 16; Program prog(Arch::x86_64); @@ -414,14 +414,14 @@ TC_TEST("slpmatvecmul") { /* for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == 1 + i * 7); - TC_CHECK(b.val(i) == 2 + i * 9); + TI_CHECK(a.val(i) == 1 + i * 7); + TI_CHECK(b.val(i) == 2 + i * 9); } */ }; // scalar a * scalar b * vec c -TC_TEST("mixed_simd1") { +TI_TEST("mixed_simd1") { for (auto vec_size : {4, 8, 16}) { Program prog; @@ -461,14 +461,14 @@ TC_TEST("mixed_simd1") { for (int j = 0; j < vec_size; j++) { auto val = v(j).val(i); float32 gt = i * j * 2; - TC_CHECK_EQUAL(gt, val, 1e-3_f); + TI_CHECK_EQUAL(gt, val, 1e-3_f); } } } } // Vec reduction -TC_TEST("mixed_simd2") { +TI_TEST("mixed_simd2") { int n = 64; for (auto vec_size : {4, 8, 16}) { @@ -514,13 +514,13 @@ TC_TEST("mixed_simd2") { for (int i = 0; i < n; i++) { auto val = sum.val(i); float32 gt = vec_size * (vec_size - 1) / 2 + i * vec_size; - TC_CHECK_EQUAL(gt, val, 1e-5_f); + TI_CHECK_EQUAL(gt, val, 1e-5_f); } } } // reduce(vec_a ** 2 - vec_b ** 2) * vec_c<2n> -TC_TEST("mixed_simd3_slp") { +TI_TEST("mixed_simd3_slp") { for (auto vec_size : {16}) { // why vec_size = 16 fails?? Program prog; @@ -581,13 +581,13 @@ TC_TEST("mixed_simd3_slp") { for (int j = 0; j < vec_size * 2; j++) { auto val = c(j).val(i); auto gt = s * (i - 2 + j); - TC_CHECK_EQUAL(gt, val, 1e-3_f); + TI_CHECK_EQUAL(gt, val, 1e-3_f); } } } } -TC_TEST("vector_split1") { +TI_TEST("vector_split1") { CoreState::set_trigger_gdb_when_crash(true); int n = 32; Program prog(Arch::x86_64); @@ -603,11 +603,11 @@ TC_TEST("vector_split1") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == i); + TI_CHECK(a.val(i) == i); } }; -TC_TEST("vector_split_slp") { +TI_TEST("vector_split_slp") { CoreState::set_trigger_gdb_when_crash(true); int n = 256; Program prog(Arch::x86_64); @@ -632,14 +632,14 @@ TC_TEST("vector_split_slp") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == 1 + i); - TC_CHECK(b.val(i) == 2 + i); - TC_CHECK(c.val(i) == 3 + i); - TC_CHECK(d.val(i) == 4 + i); + TI_CHECK(a.val(i) == 1 + i); + TI_CHECK(b.val(i) == 2 + i); + TI_CHECK(c.val(i) == 3 + i); + TI_CHECK(d.val(i) == 4 + i); } }; -TC_TEST("union_cast") { +TI_TEST("union_cast") { CoreState::set_trigger_gdb_when_crash(true); for (auto arch : {Arch::x86_64, Arch::gpu}) { int n = 16; @@ -660,13 +660,13 @@ TC_TEST("union_cast") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == + TI_CHECK(a.val(i) == union_cast(union_cast(i * 1000) + 1234.0f)); } } }; -TC_TEST("logic_not") { +TI_TEST("logic_not") { CoreState::set_trigger_gdb_when_crash(true); for (auto arch : {Arch::x86_64, Arch::gpu}) { int n = 16; @@ -685,14 +685,14 @@ TC_TEST("logic_not") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val() == 0); - TC_CHECK(b.val() != 0); - TC_CHECK(c.val() == 0); + TI_CHECK(a.val() == 0); + TI_CHECK(b.val() != 0); + TI_CHECK(c.val() == 0); } } }; -TC_TEST("simd_if_5") { +TI_TEST("simd_if_5") { CoreState::set_trigger_gdb_when_crash(true); for (auto arch : {Arch::x86_64}) { for (auto vec : {1, 4, 8}) { @@ -713,13 +713,13 @@ TC_TEST("simd_if_5") { }); })(); for (int i = 0; i < n; i++) { - TC_CHECK(c.val(i) == 1); + TI_CHECK(c.val(i) == 1); } } } }; -TC_TEST("point_inside_box") { +TI_TEST("point_inside_box") { CoreState::set_trigger_gdb_when_crash(true); for (auto arch : {Arch::x86_64}) { for (auto vec : {1, 4, 8}) { @@ -746,13 +746,13 @@ TC_TEST("point_inside_box") { }); })(); for (int i = 0; i < n; i++) { - TC_CHECK(bool(c.val(i)) == true); + TI_CHECK(bool(c.val(i)) == true); } } } }; -TC_TEST("while_in_while") { +TI_TEST("while_in_while") { CoreState::set_trigger_gdb_when_crash(true); for (auto arch : {Arch::x86_64}) { for (auto vec : {1, 4, 8}) { @@ -782,13 +782,13 @@ TC_TEST("while_in_while") { }); })(); for (int i = 0; i < n; i++) { - TC_CHECK(c.val(i) == ((i + 1) % 2) * 100); + TI_CHECK(c.val(i) == ((i + 1) % 2) * 100); } } } }; -TC_TEST("cmp") { +TI_TEST("cmp") { CoreState::set_trigger_gdb_when_crash(true); for (auto arch : {Arch::x86_64}) { for (auto vec : {1, 4, 8}) { @@ -812,7 +812,7 @@ TC_TEST("cmp") { For(0, n, [&](Expr i) { c[i] = a[i] OP b[i]; }); \ })(); \ for (int i = 0; i < n; i++) { \ - TC_CHECK(bool(c.val(i)) == \ + TI_CHECK(bool(c.val(i)) == \ bool(a.val(i) OP b.val(i))); \ } diff --git a/misc/obsolete_cpp_tests/dynamic.cpp b/misc/obsolete_cpp_tests/dynamic.cpp index 2d7c1e832c9a2..6d866ff4f6e75 100644 --- a/misc/obsolete_cpp_tests/dynamic.cpp +++ b/misc/obsolete_cpp_tests/dynamic.cpp @@ -4,7 +4,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("append_and_probe") { +TI_TEST("append_and_probe") { CoreState::set_trigger_gdb_when_crash(true); for (auto arch : {Arch::x86_64, Arch::gpu}) { int n = 32; @@ -31,12 +31,12 @@ TC_TEST("append_and_probe") { })(); for (int i = 0; i < n; i++) - TC_CHECK(x.val(i) == i); - TC_CHECK(len.val() == n); + TI_CHECK(x.val(i) == i); + TI_CHECK(len.val() == n); } }; -TC_TEST("activate") { +TI_TEST("activate") { for (auto arch : {Arch::x86_64, Arch::gpu}) { int n = 32; Program prog(arch); @@ -66,13 +66,13 @@ TC_TEST("activate") { for (int i = 0; i < n; i++) { for (int j = 0; j < i; j++) { - TC_CHECK(x.val(i, j) == i + j); + TI_CHECK(x.val(i, j) == i + j); } } } }; -TC_TEST("task_list") { +TI_TEST("task_list") { for (auto arch : {Arch::gpu}) { int n = 262144; int m = 64; @@ -105,13 +105,13 @@ TC_TEST("task_list") { for (int i = 0; i < n; i++) { if (i % 5 == 4) for (int j = 0; j < m; j++) { - TC_CHECK(x.val(i, j) == i + j + P); + TI_CHECK(x.val(i, j) == i + j + P); } } } }; -TC_TEST("task_list_dynamic") { +TI_TEST("task_list_dynamic") { for (auto arch : {Arch::gpu}) { int n = 262144; int m = 64; @@ -142,13 +142,13 @@ TC_TEST("task_list_dynamic") { for (int i = 0; i < n; i++) { if (i % 5 == 4) for (int j = 0; j < 1; j++) { - TC_CHECK(x.val(i, j) == i + j + P); + TI_CHECK(x.val(i, j) == i + j + P); } } } }; -TC_TEST("parallel_append") { +TI_TEST("parallel_append") { for (auto arch : {Arch::gpu}) { int n = 32; Program prog(arch); @@ -178,18 +178,18 @@ TC_TEST("parallel_append") { else append(); auto stat = x.parent().parent().snode()->stat(); - TC_CHECK(stat.num_resident_blocks == n); + TI_CHECK(stat.num_resident_blocks == n); if (i % 2) for (int i = 0; i < n; i++) { for (int j = 0; j < n; j++) { - TC_CHECK(x.val(i, j) == 0); + TI_CHECK(x.val(i, j) == 0); } } } } }; -TC_TEST("append_2d") { +TI_TEST("append_2d") { for (auto arch : {Arch::x86_64, Arch::gpu}) { int n = 32; Program prog(arch); @@ -211,13 +211,13 @@ TC_TEST("append_2d") { for (int i = 0; i < n; i++) { for (int j = 0; j < i; j++) { - TC_CHECK(x.val(i, j) == i + j); + TI_CHECK(x.val(i, j) == i + j); } } } }; -TC_TEST("clear") { +TI_TEST("clear") { CoreState::set_trigger_gdb_when_crash(true); for (auto arch : {Arch::x86_64, Arch::gpu}) { int n = 32; @@ -249,15 +249,15 @@ TC_TEST("clear") { for (int i = 0; i < n; i++) { for (int j = 0; j < n; j++) { if (j < i) - TC_CHECK(x.val(i, j) == i + j); + TI_CHECK(x.val(i, j) == i + j); else - TC_CHECK(x.val(i, j) == 0); + TI_CHECK(x.val(i, j) == 0); } } } }; -TC_TEST("sort") { +TI_TEST("sort") { for (auto arch : {Arch::x86_64, Arch::gpu}) { int n = 4; Program prog(arch); @@ -273,7 +273,7 @@ TC_TEST("sort") { Global(c, i32); Global(coord, i32); Global(p, i32); - TC_P(particles); + TI_P(particles); layout([&]() { auto i = Index(0); auto j = Index(1); @@ -309,12 +309,12 @@ TC_TEST("sort") { })(); for (int i = 0; i < n * n; i++) { - TC_CHECK(c.val(i) == count[i]); + TI_CHECK(c.val(i) == count[i]); } } }; -TC_TEST("dilate") { +TI_TEST("dilate") { for (auto arch : {Arch::x86_64, Arch::gpu}) { for (auto ds : {1}) { int n = 32; @@ -350,16 +350,16 @@ TC_TEST("dilate") { for (int i = 0; i < n; i++) { int bid = i / bs; - TC_CHECK(x.val(i) == (1 <= bid && bid < 4)); + TI_CHECK(x.val(i) == (1 <= bid && bid < 4)); } for (int i = 0; i < n / bs; i++) { - TC_CHECK(y.val(i) == (1 <= i && i < 4)); + TI_CHECK(y.val(i) == (1 <= i && i < 4)); } } } }; -TC_TEST("dynamic_sort") { +TI_TEST("dynamic_sort") { for (auto arch : {Arch::x86_64, Arch::gpu}) { int n = 4; Program prog(arch); @@ -375,7 +375,7 @@ TC_TEST("dynamic_sort") { Global(c, i32); Global(coord, i32); Global(p, i32); - TC_P(particles); + TI_P(particles); layout([&]() { auto i = Index(0); auto j = Index(1); @@ -414,7 +414,7 @@ TC_TEST("dynamic_sort") { })(); for (int i = 0; i < n * n; i++) { - TC_CHECK(c.val(i) == count[i]); + TI_CHECK(c.val(i) == count[i]); } } }; @@ -434,7 +434,7 @@ auto reset_grid_benchmark = []() { auto i = Index(0), j = Index(1), k = Index(2); layout([&]() { - TC_ASSERT(n % grid_block_size == 0); + TI_ASSERT(n % grid_block_size == 0); auto &block = root.dense({i, j, k}, n / grid_block_size); constexpr bool block_soa = false; if (block_soa) { @@ -449,7 +449,7 @@ auto reset_grid_benchmark = []() { } }); - TC_ASSERT(bit::is_power_of_two(n)); + TI_ASSERT(bit::is_power_of_two(n)); auto &reset_grid = kernel([&]() { Declare(i); @@ -464,8 +464,8 @@ auto reset_grid_benchmark = []() { }); while (1) - TC_TIME(reset_grid()); + TI_TIME(reset_grid()); }; -TC_REGISTER_TASK(reset_grid_benchmark); +TI_REGISTER_TASK(reset_grid_benchmark); TLANG_NAMESPACE_END diff --git a/misc/obsolete_cpp_tests/gpu.cpp b/misc/obsolete_cpp_tests/gpu.cpp index 3b34b3cdab72b..ea7d6a46e4222 100644 --- a/misc/obsolete_cpp_tests/gpu.cpp +++ b/misc/obsolete_cpp_tests/gpu.cpp @@ -4,7 +4,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("compiler_basics_gpu") { +TI_TEST("compiler_basics_gpu") { CoreState::set_trigger_gdb_when_crash(true); int n = 128; Program prog(Arch::gpu); @@ -24,18 +24,18 @@ TC_TEST("compiler_basics_gpu") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == (2 - i % 2) * i); + TI_CHECK(a.val(i) == (2 - i % 2) * i); } }; -#if defined(CUDA_FOUND) -TC_TEST("cuda_malloc_managed") { +#if defined(TI_WITH_CUDA) +TI_TEST("cuda_malloc_managed") { void *ptr; cudaMallocManaged(&ptr, 1LL << 40); int *data = (int *)ptr; for (int i = 0; i < 100000; i++) { - TC_CHECK(data[i * 749] == 0); + TI_CHECK(data[i * 749] == 0); } cudaFree(ptr); } diff --git a/misc/obsolete_cpp_tests/hash.cpp b/misc/obsolete_cpp_tests/hash.cpp index 6c8bf4e16bbb0..a8c1aaec439ee 100644 --- a/misc/obsolete_cpp_tests/hash.cpp +++ b/misc/obsolete_cpp_tests/hash.cpp @@ -3,7 +3,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("hash") { +TI_TEST("hash") { for (auto arch : {Arch::gpu, Arch::x86_64}) { Program prog(arch); CoreState::set_trigger_gdb_when_crash(true); @@ -22,10 +22,10 @@ TC_TEST("hash") { })(); for (int i = 0; i < n * n / 2; i++) { - TC_CHECK(u.val(i) == i * 2); + TI_CHECK(u.val(i) == i * 2); } for (int i = n * n / 2; i < n * n; i++) { - TC_CHECK(u.val(i) == 0); + TI_CHECK(u.val(i) == 0); } } } diff --git a/misc/obsolete_cpp_tests/micro_access_ops.cpp b/misc/obsolete_cpp_tests/micro_access_ops.cpp index 031c033895cee..8f38699acd1b1 100644 --- a/misc/obsolete_cpp_tests/micro_access_ops.cpp +++ b/misc/obsolete_cpp_tests/micro_access_ops.cpp @@ -3,7 +3,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("lower_access_basics") { +TI_TEST("lower_access_basics") { for (auto arch : {Arch::x86_64, Arch::gpu}) { for (auto vec : {1, 4, 8}) { CoreState::set_trigger_gdb_when_crash(true); @@ -21,7 +21,7 @@ TC_TEST("lower_access_basics") { func(); for (int i = 0; i < n; i++) { - TC_CHECK(a.val(i) == i); + TI_CHECK(a.val(i) == i); } } } diff --git a/misc/obsolete_cpp_tests/opt.cpp b/misc/obsolete_cpp_tests/opt.cpp index 39d539321f0f6..0fe75557e3779 100644 --- a/misc/obsolete_cpp_tests/opt.cpp +++ b/misc/obsolete_cpp_tests/opt.cpp @@ -4,7 +4,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("access_simp") { +TI_TEST("access_simp") { CoreState::set_trigger_gdb_when_crash(true); int n = 16; Program prog(Arch::x86_64); @@ -29,10 +29,10 @@ TC_TEST("access_simp") { }); })(); - TC_CHECK(sum.val() == 16); + TI_CHECK(sum.val() == 16); }; -TC_TEST("root_leaf_path_weakening") { +TI_TEST("root_leaf_path_weakening") { CoreState::set_trigger_gdb_when_crash(true); int n = 16; Program prog(Arch::x86_64); diff --git a/misc/obsolete_cpp_tests/scalar_svd.h b/misc/obsolete_cpp_tests/scalar_svd.h index 6dc470a7de91f..925f75c4e42b8 100644 --- a/misc/obsolete_cpp_tests/scalar_svd.h +++ b/misc/obsolete_cpp_tests/scalar_svd.h @@ -37,7 +37,7 @@ A. McAdams, A. Selle, R. Tamstorf, J. Teran and E. Sifakis // POSSIBILITY OF SUCH DAMAGE. //##################################################################### -TC_FORCE_INLINE float rsqrt(const float f) { +TI_FORCE_INLINE float rsqrt(const float f) { return 1.0f / std::sqrt(f); } @@ -48,7 +48,7 @@ constexpr float Cosine_Pi_Over_Eight = 0.9238795325112867f; //.5 * sqrt(2. + sqrt(2.)); template -TC_FORCE_INLINE void svd(const float a11, +TI_FORCE_INLINE void svd(const float a11, const float a12, const float a13, const float a21, diff --git a/misc/obsolete_cpp_tests/scratch_pad.cpp b/misc/obsolete_cpp_tests/scratch_pad.cpp index 5a7e2e1a17d65..99ef7a3596c64 100644 --- a/misc/obsolete_cpp_tests/scratch_pad.cpp +++ b/misc/obsolete_cpp_tests/scratch_pad.cpp @@ -4,7 +4,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("scratch_pad_bounds") { +TI_TEST("scratch_pad_bounds") { Program prog; int N = 8; @@ -22,26 +22,26 @@ TC_TEST("scratch_pad_bounds") { pad.access({1, 2, -3}, ScratchPad::AccessFlag::read); - TC_CHECK(pad.bounds[0][0] == 1); - TC_CHECK(pad.bounds[0][1] == 2); - TC_CHECK(pad.bounds[0][2] == -3); + TI_CHECK(pad.bounds[0][0] == 1); + TI_CHECK(pad.bounds[0][1] == 2); + TI_CHECK(pad.bounds[0][2] == -3); - TC_CHECK(pad.bounds[1][0] == 2); - TC_CHECK(pad.bounds[1][1] == 3); - TC_CHECK(pad.bounds[1][2] == -2); + TI_CHECK(pad.bounds[1][0] == 2); + TI_CHECK(pad.bounds[1][1] == 3); + TI_CHECK(pad.bounds[1][2] == -2); pad.access({4, -2, 5}, ScratchPad::AccessFlag::read); - TC_CHECK(pad.bounds[0][0] == 1); - TC_CHECK(pad.bounds[0][1] == -2); - TC_CHECK(pad.bounds[0][2] == -3); + TI_CHECK(pad.bounds[0][0] == 1); + TI_CHECK(pad.bounds[0][1] == -2); + TI_CHECK(pad.bounds[0][2] == -3); - TC_CHECK(pad.bounds[1][0] == 5); - TC_CHECK(pad.bounds[1][1] == 3); - TC_CHECK(pad.bounds[1][2] == 6); + TI_CHECK(pad.bounds[1][0] == 5); + TI_CHECK(pad.bounds[1][1] == 3); + TI_CHECK(pad.bounds[1][2] == 6); } -TC_TEST("range_assumption") { +TI_TEST("range_assumption") { CoreState::set_trigger_gdb_when_crash(true); Program prog(Arch::gpu); @@ -70,7 +70,7 @@ TC_TEST("range_assumption") { })(); }; -TC_TEST("scratch_pad_3d") { +TI_TEST("scratch_pad_3d") { CoreState::set_trigger_gdb_when_crash(true); int n = 10000000; Program prog(Arch::gpu); @@ -157,14 +157,14 @@ TC_TEST("scratch_pad_3d") { x_val(i, j + 1, k) - x_val(i - 1, j, k) - x_val(i + 1, j, k); if (std::abs(gt - y.val(i, j, k)) > 1) { - TC_P(d); - TC_P(gt); - TC_P(y.val(i, j, k)); - TC_P(i); - TC_P(j); - TC_P(k); + TI_P(d); + TI_P(gt); + TI_P(y.val(i, j, k)); + TI_P(i); + TI_P(j); + TI_P(k); } - TC_CHECK_EQUAL(gt, y.val(i, j, k), + TI_CHECK_EQUAL(gt, y.val(i, j, k), 1e-1f / domain_size / domain_size); } } diff --git a/misc/obsolete_cpp_tests/stencil1d.cpp b/misc/obsolete_cpp_tests/stencil1d.cpp index e71fefcb38355..82a071527fac5 100644 --- a/misc/obsolete_cpp_tests/stencil1d.cpp +++ b/misc/obsolete_cpp_tests/stencil1d.cpp @@ -197,7 +197,7 @@ void benchmark_layers() { } }; - TC_P(measure_cpe(test_hash_table, n)); + TI_P(measure_cpe(test_hash_table, n)); auto &tile = data.begin()->second; @@ -210,13 +210,13 @@ void benchmark_layers() { } }; - TC_P(measure_cpe(test_block, n)); - TC_P(cnt); + TI_P(measure_cpe(test_block, n)); + TI_P(cnt); } TLANG_NAMESPACE_BEGIN -TC_TEST("stencil1d") { +TI_TEST("stencil1d") { CoreState::set_trigger_gdb_when_crash(true); Program prog; @@ -277,24 +277,24 @@ TC_TEST("stencil1d") { } } } - TC_P(total_tiles); - TC_P(total_blocks); - TC_P(total_nodes); + TI_P(total_tiles); + TI_P(total_blocks); + TI_P(total_nodes); } // benchmark_layers(); - // TC_P(measure_cpe(stencil, total_nodes)); + // TI_P(measure_cpe(stencil, total_nodes)); for (int i = 0; i < 10; i++) - TC_TIME(copy_ref()); + TI_TIME(copy_ref()); for (int i = 0; i < 10; i++) - TC_TIME(copy_optimized()); + TI_TIME(copy_optimized()); for (int i = 0; i < 10; i++) - TC_TIME(copy()); + TI_TIME(copy()); for (int i = 0; i < 10; i++) - TC_TIME(copy_parallelized()); + TI_TIME(copy_parallelized()); // test copy to x for (auto &it : data) { @@ -305,25 +305,25 @@ TC_TEST("stencil1d") { continue; for (int n = 0; n < ::Block::size; n++) { int i = it.first * dim0 + b * dim1 + n; - TC_CHECK(block->nodes[n].y == y.val(i)); + TI_CHECK(block->nodes[n].y == y.val(i)); } } } for (int i = 0; i < 10; i++) - TC_TIME(stencil_ref()); + TI_TIME(stencil_ref()); for (int i = 0; i < 10; i++) - TC_TIME(stencil_optimized()); + TI_TIME(stencil_optimized()); for (int i = 0; i < 10; i++) - TC_TIME(stencil_optimized2()); + TI_TIME(stencil_optimized2()); for (int i = 0; i < 10; i++) - TC_TIME(stencil_optimized3()); + TI_TIME(stencil_optimized3()); for (int i = 0; i < 10; i++) - TC_TIME(stencil()); + TI_TIME(stencil()); // test stencil to x for (auto &it : data) { @@ -334,7 +334,7 @@ TC_TEST("stencil1d") { continue; for (int n = 0; n < ::Block::size; n++) { int i = it.first * dim0 + b * dim1 + n; - TC_CHECK_EQUAL(block->nodes[n].y, y.val(i), 1e-5f); + TI_CHECK_EQUAL(block->nodes[n].y, y.val(i), 1e-5f); } } } diff --git a/misc/obsolete_cpp_tests/struct.cpp b/misc/obsolete_cpp_tests/struct.cpp index 139572d07c24d..48432e732d77c 100644 --- a/misc/obsolete_cpp_tests/struct.cpp +++ b/misc/obsolete_cpp_tests/struct.cpp @@ -3,7 +3,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("snode") { +TI_TEST("snode") { Program prog(Arch::x86_64); auto i = Index(0); @@ -19,11 +19,11 @@ TC_TEST("snode") { } for (int i = 0; i < n; i++) { - TC_CHECK_EQUAL(u.val(i), i + 1, 0); + TI_CHECK_EQUAL(u.val(i), i + 1, 0); } } -TC_TEST("snode_loop") { +TI_TEST("snode_loop") { for (auto arch : {Arch::x86_64, Arch::gpu}) { Program prog(arch); CoreState::set_trigger_gdb_when_crash(true); @@ -43,12 +43,12 @@ TC_TEST("snode_loop") { })(); for (int i = 0; i < n; i++) { - TC_CHECK_EQUAL(u.val(i), i * 2, 0); + TI_CHECK_EQUAL(u.val(i), i * 2, 0); } } } -TC_TEST("snode_loop2") { +TI_TEST("snode_loop2") { for (auto arch : {Arch::x86_64, Arch::gpu}) { Program prog(arch); CoreState::set_trigger_gdb_when_crash(true); @@ -65,10 +65,10 @@ TC_TEST("snode_loop2") { root.dense(j, n).place(v); }); - TC_ASSERT( + TI_ASSERT( u.cast()->snode->physical_index_position[0] == 0); - TC_ASSERT( + TI_ASSERT( v.cast()->snode->physical_index_position[0] == 1); @@ -83,13 +83,13 @@ TC_TEST("snode_loop2") { })(); for (int i = 0; i < n; i++) { - TC_CHECK_EQUAL(u.val(i), i * 2, 0); - TC_CHECK_EQUAL(v.val(i), i * 3, 0); + TI_CHECK_EQUAL(u.val(i), i * 2, 0); + TI_CHECK_EQUAL(v.val(i), i * 3, 0); } } } -TC_TEST("2d_blocked_array") { +TI_TEST("2d_blocked_array") { int n = 8, block_size = 4; for (auto arch : {Arch::x86_64, Arch::gpu}) @@ -103,7 +103,7 @@ TC_TEST("2d_blocked_array") { auto i = Index(0); auto j = Index(1); if (blocked) { - TC_ASSERT(n % block_size == 0); + TI_ASSERT(n % block_size == 0); root.dense({i, j}, {n / block_size, n * 2 / block_size}) .dense({i, j}, {block_size, block_size}) .place(a, b); @@ -124,14 +124,14 @@ TC_TEST("2d_blocked_array") { for (int i = 0; i < n; i++) { for (int j = 0; j < n * 2; j++) { - TC_CHECK(a.val(i, j) == i + j * 3); - TC_CHECK(b.val(i, j) == i * 2 + j * 3); + TI_CHECK(a.val(i, j) == i + j * 3); + TI_CHECK(b.val(i, j) == i * 2 + j * 3); } } } } -TC_TEST("2d_blocked_array_morton") { +TI_TEST("2d_blocked_array_morton") { int n = 16, block_size = 4; for (auto arch : {Arch::x86_64}) { @@ -143,7 +143,7 @@ TC_TEST("2d_blocked_array_morton") { layout([&] { auto i = Index(0); auto j = Index(1); - TC_ASSERT(n % block_size == 0); + TI_ASSERT(n % block_size == 0); root.dense({i, j}, {n / block_size, n / block_size}) .morton() .dense({i, j}, {block_size, block_size}) @@ -160,14 +160,14 @@ TC_TEST("2d_blocked_array_morton") { for (int i = 0; i < n; i++) { for (int j = 0; j < n; j++) { - TC_CHECK(a.val(i, j) == i + j * 3); - TC_CHECK(b.val(i, j) == i * 2 + j * 3); + TI_CHECK(a.val(i, j) == i + j * 3); + TI_CHECK(b.val(i, j) == i * 2 + j * 3); } } } } -TC_TEST("bitmask_clear") { +TI_TEST("bitmask_clear") { CoreState::set_trigger_gdb_when_crash(true); int n = 256, block_size = 16; @@ -180,7 +180,7 @@ TC_TEST("bitmask_clear") { layout([&] { auto i = Index(0); auto j = Index(1); - TC_ASSERT(n % block_size == 0); + TI_ASSERT(n % block_size == 0); root.dense({i, j}, {n / block_size, n / block_size}) .bitmasked() .dense({i, j}, {block_size, block_size}) @@ -197,11 +197,11 @@ TC_TEST("bitmask_clear") { for (int i = 0; i < n / 2; i++) { for (int j = n / 2; j < n; j++) { - TC_CHECK(a.val(i, j) == i + j * 3); - TC_CHECK(b.val(i, j) == i + j * 3 + 1); + TI_CHECK(a.val(i, j) == i + j * 3); + TI_CHECK(b.val(i, j) == i + j * 3 + 1); } } - TC_CHECK(b.val(1, 1) == 0); + TI_CHECK(b.val(1, 1) == 0); a.parent().snode()->clear_data(); @@ -209,11 +209,11 @@ TC_TEST("bitmask_clear") { for (int i = 0; i < n / 2; i++) { for (int j = n / 2; j < n; j++) { - TC_CHECK(a.val(i, j) == 0); - TC_CHECK(b.val(i, j) == i); + TI_CHECK(a.val(i, j) == 0); + TI_CHECK(b.val(i, j) == i); } } - TC_CHECK(b.val(block_size + 1, 1) == 0); + TI_CHECK(b.val(block_size + 1, 1) == 0); a.parent().snode()->clear_data_and_deactivate(); @@ -221,15 +221,15 @@ TC_TEST("bitmask_clear") { for (int i = 0; i < n / 2; i++) { for (int j = n / 2; j < n; j++) { - TC_CHECK(a.val(i, j) == 0); - TC_CHECK(b.val(i, j) == 0); + TI_CHECK(a.val(i, j) == 0); + TI_CHECK(b.val(i, j) == 0); } } - TC_CHECK(b.val(block_size + 1, 1) == 0); + TI_CHECK(b.val(block_size + 1, 1) == 0); } } -TC_TEST("2d_blocked_array_bitmasked") { +TI_TEST("2d_blocked_array_bitmasked") { int n = 16, block_size = 4; for (auto arch : {Arch::x86_64, Arch::gpu}) { @@ -245,7 +245,7 @@ TC_TEST("2d_blocked_array_bitmasked") { layout([&] { auto i = Index(0); auto j = Index(1); - TC_ASSERT(n % block_size == 0); + TI_ASSERT(n % block_size == 0); root.dense({i, j}, {n / block_size, n / block_size}) .morton(morton) .bitmasked() @@ -264,16 +264,16 @@ TC_TEST("2d_blocked_array_bitmasked") { for (int i = 0; i < n / 2; i++) { for (int j = n / 2; j < n; j++) { - TC_CHECK(a.val(i, j) == i + j * 3); - TC_CHECK(b.val(i, j) == i * 2 + j * 3); + TI_CHECK(a.val(i, j) == i + j * 3); + TI_CHECK(b.val(i, j) == i * 2 + j * 3); } } - TC_CHECK(b.val(1, 1) == 0); + TI_CHECK(b.val(1, 1) == 0); } } } -TC_TEST("2d_blocked_array_vec") { +TI_TEST("2d_blocked_array_vec") { int n = 8, block_size = 4; for (auto arch : {Arch::x86_64}) @@ -287,7 +287,7 @@ TC_TEST("2d_blocked_array_vec") { auto i = Index(0); auto j = Index(1); if (blocked) { - TC_ASSERT(n % block_size == 0); + TI_ASSERT(n % block_size == 0); root.dense({i, j}, {n / block_size, n * 2 / block_size}) .dense({i, j}, {block_size, block_size}) .place(a, b); @@ -310,14 +310,14 @@ TC_TEST("2d_blocked_array_vec") { for (int i = 0; i < n; i++) { for (int j = 0; j < n * 2; j++) { - TC_CHECK(a.val(i, j) == i + j * 3); - TC_CHECK(b.val(i, j) == i * 2 + j * 3); + TI_CHECK(a.val(i, j) == i + j * 3); + TI_CHECK(b.val(i, j) == i * 2 + j * 3); } } } } -TC_TEST("loop_over_blocks") { +TI_TEST("loop_over_blocks") { int n = 64, block_size = 4; for (auto arch : {Arch::x86_64, Arch::gpu}) { @@ -329,7 +329,7 @@ TC_TEST("loop_over_blocks") { layout([&] { auto ij = Indices(0, 1); - TC_ASSERT(n % block_size == 0); + TI_ASSERT(n % block_size == 0); root.dense(ij, n / block_size) .dense(ij, {block_size, block_size * 2}) .place(a); @@ -351,18 +351,18 @@ TC_TEST("loop_over_blocks") { sum_j_gt += j; } } - TC_CHECK(sum_i.val() == sum_i_gt); - TC_CHECK(sum_j.val() == sum_j_gt); + TI_CHECK(sum_i.val() == sum_i_gt); + TI_CHECK(sum_j.val() == sum_j_gt); } } #if (0) -TC_TEST("spmv") { +TI_TEST("spmv") { initialize_benchmark(); int n = 8192; int band = 256; int k = 128; - TC_ASSERT(k <= band); + TI_ASSERT(k <= band); int m = n * k; Eigen::SparseMatrix M(n, n); @@ -430,31 +430,31 @@ TC_TEST("spmv") { vec_val.val(i) = val; } - TC_TIME(M.setFromTriplets(entries.begin(), entries.end())); + TI_TIME(M.setFromTriplets(entries.begin(), entries.end())); - TC_TIME(populate()); + TI_TIME(populate()); int T = 1; for (int i = 0; i < T; i++) { - TC_TIME(matvecmul()); + TI_TIME(matvecmul()); } - TC_P(n = Eigen::nbThreads()); + TI_P(n = Eigen::nbThreads()); for (int i = 0; i < T; i++) { - TC_TIME(Vret = M * V); + TI_TIME(Vret = M * V); } for (int i = 0; i < n; i++) { - TC_CHECK_EQUAL(Vret(i), result.val(i) / T, 1e-3_f); + TI_CHECK_EQUAL(Vret(i), result.val(i) / T, 1e-3_f); } } -TC_TEST("spmv_dynamic") { +TI_TEST("spmv_dynamic") { initialize_benchmark(); int n = 8192; int band = 256; int k = 128; - TC_ASSERT(k <= band); + TI_ASSERT(k <= band); int m = n * k; Eigen::SparseMatrix M(n, n); @@ -524,27 +524,27 @@ TC_TEST("spmv_dynamic") { vec_val.val(i) = val; } - TC_TIME(M.setFromTriplets(entries.begin(), entries.end())); + TI_TIME(M.setFromTriplets(entries.begin(), entries.end())); - TC_TIME(populate()); + TI_TIME(populate()); int T = 30; for (int i = 0; i < T; i++) { - TC_TIME(matvecmul()); + TI_TIME(matvecmul()); } - TC_P(n = Eigen::nbThreads()); + TI_P(n = Eigen::nbThreads()); for (int i = 0; i < T; i++) { - TC_TIME(Vret = M * V); + TI_TIME(Vret = M * V); } for (int i = 0; i < n; i++) { - TC_CHECK_EQUAL(Vret(i), result.val(i) / T, 1e-3_f); + TI_CHECK_EQUAL(Vret(i), result.val(i) / T, 1e-3_f); } } // array of linked list -TC_TEST("indirect") { +TI_TEST("indirect") { Program prog; int n = 4; @@ -586,12 +586,12 @@ TC_TEST("indirect") { for (int i = 0; i < n; i++) { auto reduced = sum.val(i); - TC_CHECK(reduced == (i * k + (i + 1) * k + 1) * k / 2); + TI_CHECK(reduced == (i * k + (i + 1) * k + 1) * k / 2); } } #endif -TC_TEST("leaf_context") { +TI_TEST("leaf_context") { Program prog; int n = 64; @@ -620,10 +620,10 @@ TC_TEST("leaf_context") { For(i, a, [&] { sum[Expr(0)] += a[i]; }); })(); - TC_CHECK(sum.val() == sum_gt); + TI_CHECK(sum.val() == sum_gt); } -TC_TEST("pointer") { +TI_TEST("pointer") { for (auto arch : {Arch::x86_64, Arch::gpu}) { Program prog(arch); @@ -651,11 +651,11 @@ TC_TEST("pointer") { kernel([&]() { For(a, [&](Expr i) { Atomic(sum[Expr(0)]) += a[i]; }); })(); auto reduced = sum.val(); - TC_CHECK(reduced == sum_gt); + TI_CHECK(reduced == sum_gt); } } -TC_TEST("gpu_listgen") { +TI_TEST("gpu_listgen") { for (auto arch : {Arch::gpu}) { for (auto level : {1, 2, 3}) { Program prog(arch); @@ -691,12 +691,12 @@ TC_TEST("gpu_listgen") { [&]() { For(a, [&](Expr i) { Atomic(sum[Expr(0)]) += a[i]; }); })(); auto reduced = sum.val(); - TC_CHECK(reduced == sum_gt); + TI_CHECK(reduced == sum_gt); } } } -TC_TEST("misaligned") { +TI_TEST("misaligned") { // On the same tree, x has indices i while y has indices i & j Program prog; @@ -730,11 +730,11 @@ TC_TEST("misaligned") { })(); for (int i = 0; i < n; i++) { - TC_CHECK(x_gt[i] == x.val(i)); + TI_CHECK(x_gt[i] == x.val(i)); } } -TC_TEST("hashed") { +TI_TEST("hashed") { Program prog; int n = 64; @@ -766,10 +766,10 @@ TC_TEST("hashed") { })(); auto reduced = sum.val(); - TC_CHECK(reduced == sum_gt); + TI_CHECK(reduced == sum_gt); } -TC_TEST("mpm_layout") { +TI_TEST("mpm_layout") { Program prog(Arch::gpu); // Program prog(Arch::x86_64); // prog.config.print_ir = true; @@ -823,7 +823,7 @@ TC_TEST("mpm_layout") { place(particle_v(i)); place(particle_J); - TC_ASSERT(n % grid_block_size == 0); + TI_ASSERT(n % grid_block_size == 0); auto &block = root.dense({i, j, k}, grid_n / 4 / grid_block_size) .pointer() .dense({i, j, k}, 4) diff --git a/misc/obsolete_cpp_tests/svd.cpp b/misc/obsolete_cpp_tests/svd.cpp index c763568619efa..f54dbabf387fb 100644 --- a/misc/obsolete_cpp_tests/svd.cpp +++ b/misc/obsolete_cpp_tests/svd.cpp @@ -43,7 +43,7 @@ void sifakis_svd_gt(Matrix3 &a, Matrix3 &u, Matrix3 &v, Vector3 &sig) { // clang-format on } -TC_TEST("svd_scalar") { +TI_TEST("svd_scalar") { using Matrix = TMatrix; using Vector = TVector; float32 tolerance = 2e-3_f32; @@ -54,30 +54,30 @@ TC_TEST("svd_scalar") { sifakis_svd_gt<6>(m, U, V, sig_vec); sig = Matrix(sig_vec); - TC_CHECK_EQUAL(m, U * sig * transposed(V), tolerance); - TC_CHECK_EQUAL(Matrix(1), U * transposed(U), tolerance); - TC_CHECK_EQUAL(Matrix(1), V * transposed(V), tolerance); - TC_CHECK_EQUAL(sig, Matrix(sig.diag()), tolerance); + TI_CHECK_EQUAL(m, U * sig * transposed(V), tolerance); + TI_CHECK_EQUAL(Matrix(1), U * transposed(U), tolerance); + TI_CHECK_EQUAL(Matrix(1), V * transposed(V), tolerance); + TI_CHECK_EQUAL(sig, Matrix(sig.diag()), tolerance); /* if (dim == 2) { qr_decomp(m, Q, R); - TC_CHECK_EQUAL(m, Q * R, tolerance); - TC_CHECK_EQUAL(Q * transposed(Q), Matrix(1), tolerance); + TI_CHECK_EQUAL(m, Q * R, tolerance); + TI_CHECK_EQUAL(Q * transposed(Q), Matrix(1), tolerance); CHECK(abs(R[0][1]) < 1e-6_f); CHECK(R[0][0] > -1e-6_f); CHECK(R[1][1] > -1e-6_f); } polar_decomp(m, R, S); - TC_CHECK_EQUAL(m, R * S, tolerance); - TC_CHECK_EQUAL(Matrix(1), R * transposed(R), tolerance); - TC_CHECK_EQUAL(S, transposed(S), tolerance); + TI_CHECK_EQUAL(m, R * S, tolerance); + TI_CHECK_EQUAL(Matrix(1), R * transposed(R), tolerance); + TI_CHECK_EQUAL(S, transposed(S), tolerance); */ } } -TC_TEST("svd_dsl") { +TI_TEST("svd_dsl") { for (auto vec : {1}) { CoreState::set_trigger_gdb_when_crash(true); using TMat = TMatrix; @@ -136,15 +136,15 @@ TC_TEST("svd_dsl") { sig(p, p) = gSigma(p).val(i); } - TC_CHECK_EQUAL(m, U * sig * transposed(V), tolerance); - TC_CHECK_EQUAL(TMat(1), U * transposed(U), tolerance); - TC_CHECK_EQUAL(TMat(1), V * transposed(V), tolerance); - TC_CHECK_EQUAL(sig, TMat(sig.diag()), tolerance); + TI_CHECK_EQUAL(m, U * sig * transposed(V), tolerance); + TI_CHECK_EQUAL(TMat(1), U * transposed(U), tolerance); + TI_CHECK_EQUAL(TMat(1), V * transposed(V), tolerance); + TI_CHECK_EQUAL(sig, TMat(sig.diag()), tolerance); } } } -TC_TEST("svd_dsl_float64") { +TI_TEST("svd_dsl_float64") { for (auto vec : {1}) { CoreState::set_trigger_gdb_when_crash(true); using TMat = TMatrix; @@ -204,10 +204,10 @@ TC_TEST("svd_dsl_float64") { sig(p, p) = gSigma(p).val(i); } - TC_CHECK_EQUAL(m, U * sig * transposed(V), tolerance); - TC_CHECK_EQUAL(TMat(1), U * transposed(U), tolerance); - TC_CHECK_EQUAL(TMat(1), V * transposed(V), tolerance); - TC_CHECK_EQUAL(sig, TMat(sig.diag()), tolerance); + TI_CHECK_EQUAL(m, U * sig * transposed(V), tolerance); + TI_CHECK_EQUAL(TMat(1), U * transposed(U), tolerance); + TI_CHECK_EQUAL(TMat(1), V * transposed(V), tolerance); + TI_CHECK_EQUAL(sig, TMat(sig.diag()), tolerance); } } } diff --git a/misc/obsolete_cpp_tests/test_linalg.cpp b/misc/obsolete_cpp_tests/test_linalg.cpp index f211e3b6bd845..a24c94f74b658 100644 --- a/misc/obsolete_cpp_tests/test_linalg.cpp +++ b/misc/obsolete_cpp_tests/test_linalg.cpp @@ -6,7 +6,7 @@ #include #include -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN template void test_matrix() { @@ -16,16 +16,16 @@ void test_matrix() { Matrix m = Matrix::rand(); if (determinant(m) > tolerance * 1e3_f) { if (!math::equal(m * inversed(m), Matrix(1), tolerance)) { - TC_P(m * inversed(m) - Matrix(1)); - TC_P(math::abs(m * inversed(m) - Matrix(1))); - TC_P(math::maximum(math::abs(m * inversed(m) - Matrix(1)))); + TI_P(m * inversed(m) - Matrix(1)); + TI_P(math::abs(m * inversed(m) - Matrix(1))); + TI_P(math::maximum(math::abs(m * inversed(m) - Matrix(1)))); } - TC_CHECK_EQUAL(m * inversed(m), Matrix(1), tolerance); + TI_CHECK_EQUAL(m * inversed(m), Matrix(1), tolerance); } } } -TC_TEST("vector arith") { +TI_TEST("vector arith") { Vector3 a(1, 2, 3), b(4, 2, 5); CHECK(a + b == Vector3(5, 4, 8)); CHECK(b - a == Vector3(3, 0, 2)); @@ -46,7 +46,7 @@ TC_TEST("vector arith") { CHECK(c + d == Vector2(3, 7)); CHECK(Vector4(1, 2, 3, 1).length2() == 15.0_f); -#if !defined(TC_USE_DOUBLE) && !defined(TC_ISE_NONE) +#if !defined(TI_USE_DOUBLE) && !defined(TI_ISE_NONE) CHECK(Vector3(1, 2, 3, 1).length2() == 14.0_f); #endif CHECK(dot(Vector2(1, 2), Vector2(3, 2)) == 7.0_f); @@ -88,4 +88,4 @@ TC_TEST("vector arith") { "std::string is not VectorND."); } -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/misc/obsolete_cpp_tests/test_system.cpp b/misc/obsolete_cpp_tests/test_system.cpp index 4e8f5d1a429dd..a00958c211d32 100644 --- a/misc/obsolete_cpp_tests/test_system.cpp +++ b/misc/obsolete_cpp_tests/test_system.cpp @@ -7,9 +7,9 @@ #include #include -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN -TC_TEST("Virtual Memory") { +TI_TEST("Virtual Memory") { for (int i = 0; i < 3; i++) { // Allocate 1 TB of virtual memory std::size_t size = 1LL << 40; @@ -23,4 +23,4 @@ TC_TEST("Virtual Memory") { } } -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/misc/obsolete_cpp_tests/types.cpp b/misc/obsolete_cpp_tests/types.cpp index 78b7eb83bcbb3..5dde418c4d2b7 100644 --- a/misc/obsolete_cpp_tests/types.cpp +++ b/misc/obsolete_cpp_tests/types.cpp @@ -4,7 +4,7 @@ TLANG_NAMESPACE_BEGIN -TC_TEST("float64") { +TI_TEST("float64") { CoreState::set_trigger_gdb_when_crash(true); int n = 32; Program prog(Arch::gpu); @@ -20,12 +20,12 @@ TC_TEST("float64") { func(); - // TC_CHECK(sum.val() == n); + // TI_CHECK(sum.val() == n); for (int i = 0; i < n; i++) - TC_CHECK(a.val(i) == i * 2); + TI_CHECK(a.val(i) == i * 2); }; -TC_TEST("llvm_exception") { +TI_TEST("llvm_exception") { CoreState::set_trigger_gdb_when_crash(true); int n = 1; for (int i = 0; i < 2; i++) { @@ -41,7 +41,7 @@ TC_TEST("llvm_exception") { try { throw IRModified(); } catch (IRModified) { - TC_TAG; + TI_TAG; } }; diff --git a/python/taichi/core/util.py b/python/taichi/core/util.py index 7a3f31090803d..64a56e83d2638 100644 --- a/python/taichi/core/util.py +++ b/python/taichi/core/util.py @@ -50,7 +50,7 @@ def locale_encode(s): def is_ci(): - return os.environ.get('TC_CI', '') == '1' + return os.environ.get('TI_CI', '') == '1' def package_root(): diff --git a/python/taichi/lang/__init__.py b/python/taichi/lang/__init__.py index 5040d1c268820..62592a39f0cae 100644 --- a/python/taichi/lang/__init__.py +++ b/python/taichi/lang/__init__.py @@ -45,7 +45,7 @@ def reset(): global runtime runtime = get_runtime() -def init(default_fp=None, default_ip=None, print_processed=None, debug=None, **kwargs): +def init(default_fp=None, default_ip=None, print_preprocessed=None, debug=None, **kwargs): if debug is None: debug = bool(int(os.environ.get('TI_DEBUG', '0'))) @@ -61,8 +61,8 @@ def init(default_fp=None, default_ip=None, print_processed=None, debug=None, **k ti.get_runtime().set_default_fp(default_fp) if default_ip is not None: ti.get_runtime().set_default_ip(default_ip) - if print_processed is not None: - ti.get_runtime().print_preprocessed = print_processed + if print_preprocessed is not None: + ti.get_runtime().print_preprocessed = print_preprocessed if debug: ti.set_logging_level(ti.DEBUG) ti.cfg.debug = debug diff --git a/python/taichi/lang/snode.py b/python/taichi/lang/snode.py index 0f12bd9413c07..348732f2e6023 100644 --- a/python/taichi/lang/snode.py +++ b/python/taichi/lang/snode.py @@ -8,15 +8,22 @@ def dense(self, indices, dimensions): dimensions = [dimensions] * len(indices) return SNode(self.ptr.dense(indices, dimensions)) + def pointer(self, indices, dimensions): + if isinstance(dimensions, int): + dimensions = [dimensions] * len(indices) + return SNode(self.ptr.pointer(indices, dimensions)) + + def hash(self, indices, dimensions): + if isinstance(dimensions, int): + dimensions = [dimensions] * len(indices) + return SNode(self.ptr.hash(indices, dimensions)) + def dynamic(self, index, dimension, chunk_size=None): assert len(index) == 1 if chunk_size is None: chunk_size = dimension return SNode(self.ptr.dynamic(index[0], dimension, chunk_size)) - def pointer(self): - return SNode(self.ptr.pointer()) - def bitmasked(self, val=True): self.ptr.bitmasked(val) return self @@ -51,10 +58,10 @@ def get_shape(self, i): def loop_range(self): import taichi as ti return ti.Expr(ti.core.global_var_expr_from_snode(self.ptr)) - + def snode(self): return self - + def get_children(self): children = [] for i in range(self.ptr.get_num_ch()): diff --git a/python/taichi/main.py b/python/taichi/main.py index af98e5ec7162b..aab2ac686d372 100644 --- a/python/taichi/main.py +++ b/python/taichi/main.py @@ -6,16 +6,24 @@ from taichi.tools.video import make_video, interpolate_frames, mp4_to_gif, scale_video, crop_video, accelerate_video -def test_python(verbose=False): +def test_python(test_files=None, verbose=False): print("\nRunning python tests...\n") import taichi as ti import pytest + test_dir = None if ti.is_release(): - test_dir = os.path.join(ti.package_root(), 'tests') + test_dir = ti.package_root() else: - test_dir = os.path.join(ti.get_repo_directory(), 'tests') - - args = [test_dir] + test_dir = ti.get_repo_directory() + test_dir = os.path.join(test_dir, 'tests', 'python') + args = [] + if test_files: + # run individual tests + for f in test_files: + args.append(os.path.join(test_dir, f)) + else: + # run all the tests + args = [test_dir] if verbose: args += ['-s'] @@ -92,7 +100,7 @@ def main(debug=False): script = script.read() exec(script, {'__name__': '__main__'}) elif mode == "test_python": - return test_python() + return test_python(test_files=sys.argv[2:]) elif mode == "test_cpp": return test_cpp() elif mode == "test": @@ -100,7 +108,7 @@ def main(debug=False): return -1 return test_cpp() elif mode == "test_verbose": - if test_python(True) != 0: + if test_python(verbose=True) != 0: return -1 return test_cpp() elif mode == "build": diff --git a/python/taichi/tools/messager.py b/python/taichi/tools/messager.py index eebbd533e1100..069ae610e0c44 100644 --- a/python/taichi/tools/messager.py +++ b/python/taichi/tools/messager.py @@ -16,9 +16,9 @@ def send_crash_report(message, receiver=None): return emailed = True if receiver is None: - receiver = os.environ.get('TC_MONITOR_EMAIL', None) + receiver = os.environ.get('TI_MONITOR_EMAIL', None) if receiver is None: - tc.warning('No receiver in $TC_MONITOR_EMAIL') + tc.warning('No receiver in $TI_MONITOR_EMAIL') return tc.warning('Emailing {}'.format(receiver)) TO = receiver diff --git a/taichi/analysis/value_diff.cpp b/taichi/analysis/value_diff.cpp index 672d210a4dd1a..6b1a071ed23e4 100644 --- a/taichi/analysis/value_diff.cpp +++ b/taichi/analysis/value_diff.cpp @@ -48,7 +48,7 @@ class ValueDiff : public IRVisitor { void visit(ElementShuffleStmt *stmt) override { int old_lane = lane; - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); auto src = stmt->elements[lane].stmt; lane = stmt->elements[lane].index; src->accept(this); diff --git a/taichi/arch.h b/taichi/arch.h index aa0e12c09f5c0..24e941b6afb46 100644 --- a/taichi/arch.h +++ b/taichi/arch.h @@ -21,7 +21,7 @@ inline std::string arch_name(Arch arch) { #include "inc/archs.inc.h" #undef PER_ARCH default: - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } } diff --git a/taichi/backends/base.cpp b/taichi/backends/base.cpp index 701e863dcab7e..26bd63900eb09 100644 --- a/taichi/backends/base.cpp +++ b/taichi/backends/base.cpp @@ -1,7 +1,7 @@ // The code generator base class #include "base.h" -#if !defined(TC_PLATFORM_WINDOWS) +#if !defined(TI_PLATFORM_WINDOWS) #include #endif #include @@ -14,7 +14,7 @@ std::string CodeGenBase::get_source_path() { } std::string CodeGenBase::get_library_path() { -#if defined(TC_PLATFORM_OSX) +#if defined(TI_PLATFORM_OSX) // Note: use .so here will lead to wired behavior... return fmt::format("{}/tmp{:04d}.dylib", folder, id); #else @@ -28,7 +28,7 @@ void CodeGenBase::write_source() { std::string firstline; std::getline(ifs, firstline); if (firstline.find("debug") != firstline.npos) { - TC_WARN("Debugging file {}. Code overridden.", get_source_path()); + TI_WARN("Debugging file {}. Code overridden.", get_source_path()); return; } } @@ -42,24 +42,24 @@ void CodeGenBase::write_source() { } std::string CodeGenBase::get_source() { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED return ""; } void CodeGenBase::load_dll() { -#if defined(TC_PLATFORM_UNIX) +#if defined(TI_PLATFORM_UNIX) dll = dlopen((get_library_path()).c_str(), RTLD_LAZY); if (dll == nullptr) { - TC_ERROR("{}", dlerror()); + TI_ERROR("{}", dlerror()); } - TC_ASSERT(dll != nullptr); + TI_ASSERT(dll != nullptr); #else - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED #endif } void CodeGenBase::disassemble() { -#if defined(TC_PLATFORM_LINUX) +#if defined(TI_PLATFORM_LINUX) auto objdump_ret = system(fmt::format("objdump {} -d > {}.s", get_library_path(), get_library_path()) .c_str()); @@ -91,16 +91,16 @@ void CodeGenBase::generate_binary(std::string extra_flags) { get_current_program() .config.preprocess_cmd(get_source_path(), pp_fn, extra_flags, true) .c_str())); - TC_ERROR("Preprocessing failed."); + TI_ERROR("Preprocessing failed."); } std::ifstream ifs(pp_fn); - TC_ASSERT(ifs); + TI_ASSERT(ifs); auto hash_input = preprocess_cmd + std::string(std::istreambuf_iterator(ifs), std::istreambuf_iterator()); - // TC_P(preprocess_cmd); + // TI_P(preprocess_cmd); auto hash = XXH64(hash_input.data(), hash_input.size(), 0); - // TC_P(hash); + // TI_P(hash); std::string cached_binary_fn = db_folder() + fmt::format("/{}.so", hash); std::ifstream key_file(cached_binary_fn); @@ -117,11 +117,11 @@ void CodeGenBase::generate_binary(std::string extra_flags) { get_source_path(), get_library_path(), extra_flags); auto compile_ret = std::system(cmd.c_str()); if (compile_ret != 0) { - TC_WARN("Compilation cmd: {}", cmd); + TI_WARN("Compilation cmd: {}", cmd); auto cmd = get_current_program().config.compile_cmd( get_source_path(), get_library_path(), extra_flags, true); trash(std::system(cmd.c_str())); - TC_ERROR("Source {} compilation failed.", get_source_path()); + TI_ERROR("Source {} compilation failed.", get_source_path()); } else { trash(std::system( fmt::format("cp {} {}", get_library_path(), cached_binary_fn) @@ -129,9 +129,9 @@ void CodeGenBase::generate_binary(std::string extra_flags) { } } trash(std::system(fmt::format("rm {}", pp_fn).c_str())); - TC_INFO("Compilation time: {:.1f} ms", 1000 * (Time::get_time() - t)); + TI_INFO("Compilation time: {:.1f} ms", 1000 * (Time::get_time() - t)); #else - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED #endif } diff --git a/taichi/backends/base.h b/taichi/backends/base.h index b83dfee9323ac..7d57eb2ec0dae 100644 --- a/taichi/backends/base.h +++ b/taichi/backends/base.h @@ -5,7 +5,7 @@ #include "../snode.h" #include "../ir.h" #include "../program.h" -#if defined(TC_PLATFORM_UNIX) +#if defined(TI_PLATFORM_UNIX) #include #endif @@ -68,7 +68,7 @@ class CodeGenBase { static int get_kernel_id() { static int id = 0; - TC_ASSERT(id < 10000); + TI_ASSERT(id < 10000); return id++; } @@ -112,16 +112,16 @@ class CodeGenBase { template T load_function(std::string name) { -#if !defined(TC_PLATFORM_WINDOWS) +#if !defined(TI_PLATFORM_WINDOWS) using FP = decltype(function_pointer_helper(std::declval())); if (dll == nullptr) { load_dll(); } auto ret = dlsym(dll, name.c_str()); - TC_ASSERT(ret != nullptr); + TI_ASSERT(ret != nullptr); return T((FP)ret); #else - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED #endif } diff --git a/taichi/backends/codegen_cuda.cpp b/taichi/backends/codegen_cuda.cpp index 45e4010a4b103..289aad2ed1023 100644 --- a/taichi/backends/codegen_cuda.cpp +++ b/taichi/backends/codegen_cuda.cpp @@ -41,7 +41,7 @@ class GPUIRCodeGen : public IRVisitor { } void struct_for(Stmt *for_stmt_) { - TC_ASSERT_INFO(current_struct_for == nullptr, + TI_ASSERT_INFO(current_struct_for == nullptr, "Struct for cannot be nested."); auto for_stmt = for_stmt_->as(); current_struct_for = for_stmt; @@ -52,7 +52,7 @@ class GPUIRCodeGen : public IRVisitor { std::min(1 << leaf->total_num_bits, max_gpu_block_dim); } - TC_ASSERT((1 << leaf->total_num_bits) % for_stmt->block_dim == 0); + TI_ASSERT((1 << leaf->total_num_bits) % for_stmt->block_dim == 0); int block_division = (1 << leaf->total_num_bits) / for_stmt->block_dim; std::vector path; @@ -148,7 +148,7 @@ class GPUIRCodeGen : public IRVisitor { for (auto &pad : scratch_pads->pads) { emit("__shared__ {}::val_type {}[{}];", pad.first->node_type_name, pad.second.name(), pad.second.linear_size()); - TC_ASSERT(pad.second.is_pure()); + TI_ASSERT(pad.second.is_pure()); if (pad.second.total_flags == AccessFlag::read || pad.second.total_flags == AccessFlag::accumulate) { // read & accumulate case @@ -302,7 +302,7 @@ class GPUIRCodeGen : public IRVisitor { for (auto &o : opt) { if (o.first == 1) { ldg.insert(o.second); - TC_INFO("Caching to L1: {}", o.second->node_type_name); + TI_INFO("Caching to L1: {}", o.second->node_type_name); } else { new_opt.push_back(o); } @@ -321,7 +321,7 @@ class GPUIRCodeGen : public IRVisitor { for_stmt_counter += 1; if (for_stmt_->is()) { auto range_for = for_stmt_->as(); - TC_ASSERT(range_for->vectorize == 1); + TI_ASSERT(range_for->vectorize == 1); int begin = range_for->begin->as()->val[0].val_int32(); int end = range_for->end->as()->val[0].val_int32(); @@ -347,7 +347,7 @@ class GPUIRCodeGen : public IRVisitor { emit("{{"); int block_dim = range_for->block_dim; if (block_dim == 0) { - TC_WARN("Using default block size = 256"); + TI_WARN("Using default block size = 256"); block_dim = 256; } emit("gpu_runtime_init();"); @@ -361,7 +361,7 @@ class GPUIRCodeGen : public IRVisitor { emit("}}"); } else { auto for_stmt = for_stmt_->as(); - TC_ASSERT(for_stmt->vectorize == 1); + TI_ASSERT(for_stmt->vectorize == 1); extract_ldg(for_stmt->scratch_opt); struct_for(for_stmt_); } @@ -451,7 +451,7 @@ class GPUIRCodeGen : public IRVisitor { } void visit(RandStmt *stmt) override { - TC_ASSERT(stmt->ret_type.data_type == DataType::f32); + TI_ASSERT(stmt->ret_type.data_type == DataType::f32); emit("const auto {} = randf();", stmt->raw_name(), stmt->ret_data_type_name()); } @@ -483,7 +483,7 @@ class GPUIRCodeGen : public IRVisitor { } void visit(TernaryOpStmt *tri) override { - TC_ASSERT(tri->op_type == TernaryOpType::select); + TI_ASSERT(tri->op_type == TernaryOpType::select); emit("const {} {} = {} ? {} : {};", tri->ret_data_type_name(), tri->raw_name(), tri->op1->raw_name(), tri->op2->raw_name(), tri->op3->raw_name()); @@ -511,7 +511,7 @@ class GPUIRCodeGen : public IRVisitor { emit("printf(\"[debug] {}\" \" = %f\\n\", {});", print_stmt->str, print_stmt->stmt->raw_name()); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } } @@ -533,7 +533,7 @@ class GPUIRCodeGen : public IRVisitor { void visit(StructForStmt *for_stmt) override { // generate_loop_header(for_stmt->snode, for_stmt, true); - TC_ASSERT_INFO(current_struct_for == nullptr, + TI_ASSERT_INFO(current_struct_for == nullptr, "StructFor cannot be nested."); current_struct_for = for_stmt; for_stmt->body->accept(this); @@ -560,7 +560,7 @@ class GPUIRCodeGen : public IRVisitor { emit("{} = {}_;", loop_var->raw_name(), loop_var->raw_name()); } } else { - TC_ASSERT(!for_stmt->reversed); + TI_ASSERT(!for_stmt->reversed); emit("for ({} {} = {}; {} < {}; {} = {} + {}({})) {{", loop_var->ret_data_type_name(), loop_var->raw_name(), for_stmt->begin->raw_name(), loop_var->raw_name(), @@ -594,8 +594,8 @@ class GPUIRCodeGen : public IRVisitor { } void visit(ExternalPtrStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); - TC_ASSERT(stmt->indices.size() == 1); + TI_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->indices.size() == 1); auto dt = stmt->ret_type.data_type; emit("const {} *{}[1] = {{&{}[{}]}};", data_type_name(dt), stmt->raw_name(), stmt->base_ptrs[0]->raw_name(), stmt->indices[0]->raw_name()); @@ -607,7 +607,7 @@ class GPUIRCodeGen : public IRVisitor { return; } - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); emit("{} *{}[{}];", data_type_name(stmt->ret_type.data_type), stmt->raw_name(), stmt->ret_type.width); for (int l = 0; l < stmt->ret_type.width; l++) { @@ -618,7 +618,7 @@ class GPUIRCodeGen : public IRVisitor { std::vector indices(max_num_indices, "0"); // = "(root, "; for (int i = 0; i < (int)stmt->indices.size(); i++) { if (snode->physical_index_position[i] != -1) { - // TC_ASSERT(snode->physical_index_position[i] != -1); + // TI_ASSERT(snode->physical_index_position[i] != -1); indices[snode->physical_index_position[i]] = stmt->indices[i]->raw_name(); } @@ -634,7 +634,7 @@ class GPUIRCodeGen : public IRVisitor { void visit(SNodeOpStmt *stmt) override { /* - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); auto snode = stmt->snodes[0]; auto indices = indices_str(snode, -1, stmt->indices); @@ -650,7 +650,7 @@ class GPUIRCodeGen : public IRVisitor { make_list(indices, "")); if (stmt->op_type == SNodeOpType::append) { - TC_ASSERT(stmt->val->width() == 1); + TI_ASSERT(stmt->val->width() == 1); emit("{}_tmp->append({}({}));", snode->node_type_name, snode->ch[0]->node_type_name, stmt->val->raw_name()); } else if (stmt->op_type == SNodeOpType::clear) { @@ -670,7 +670,7 @@ class GPUIRCodeGen : public IRVisitor { emit("activate_{}(root, {});", snode->node_type_name, make_list(indices, "")); } else { - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; } emit("}}"); */ @@ -699,7 +699,7 @@ class GPUIRCodeGen : public IRVisitor { } void visit(GlobalLoadStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); if (stmt->ptr->is()) { auto ptr = stmt->ptr->as(); auto snode = ptr->snodes[0]; @@ -737,11 +737,11 @@ class GPUIRCodeGen : public IRVisitor { } void visit(AtomicOpStmt *stmt) override { - TC_ASSERT(stmt->val->ret_type.data_type == DataType::f32 || + TI_ASSERT(stmt->val->ret_type.data_type == DataType::f32 || stmt->val->ret_type.data_type == DataType::i32 || stmt->val->ret_type.data_type == DataType::f64 || stmt->val->ret_type.data_type == DataType::i64); - TC_ASSERT(stmt->op_type == AtomicOpType::add); + TI_ASSERT(stmt->op_type == AtomicOpType::add); auto ptr = stmt->dest->as(); auto snode = ptr->snodes[0]; if (current_scratch_pads && current_scratch_pads->has(snode)) { @@ -759,7 +759,7 @@ class GPUIRCodeGen : public IRVisitor { void visit(ElementShuffleStmt *stmt) override { auto init = stmt->elements.serialize( [&](const VectorElement &elem) { - TC_ASSERT(elem.index == 0); + TI_ASSERT(elem.index == 0); if (stmt->pointer) { return fmt::format("{}[0]", elem.stmt->raw_name(), elem.index); } else { @@ -780,9 +780,9 @@ class GPUIRCodeGen : public IRVisitor { // this does not necessarily hold since any index within the leaf block can // be the base /* - emit("TC_ASSERT({} + {} <= {});", stmt->base->raw_name(), stmt->low, + emit("TI_ASSERT({} + {} <= {});", stmt->base->raw_name(), stmt->low, stmt->input->raw_name()); - emit("TC_ASSERT({} < {} + {});", stmt->input->raw_name(), + emit("TI_ASSERT({} < {} + {});", stmt->input->raw_name(), stmt->base->raw_name(), stmt->high); */ emit("const auto {} = {};", stmt->raw_name(), stmt->input->raw_name()); @@ -790,7 +790,7 @@ class GPUIRCodeGen : public IRVisitor { void visit(AssertStmt *stmt) override { emit("#if defined(TL_DEBUG)"); - emit(R"(TC_ASSERT_INFO({}, "{}");)", stmt->val->raw_name(), stmt->text); + emit(R"(TI_ASSERT_INFO({}, "{}");)", stmt->val->raw_name(), stmt->text); emit("#endif"); } @@ -978,7 +978,7 @@ class GPUIRCodeGen : public IRVisitor { }; void GPUCodeGen::lower_cuda() { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } void GPUCodeGen::lower_llvm() { @@ -995,7 +995,7 @@ void GPUCodeGen::lower_llvm() { if (kernel->grad) { irpass::reverse_segments(ir); if (print_ir) { - TC_TRACE("Segment reversed (for autodiff):"); + TI_TRACE("Segment reversed (for autodiff):"); irpass::re_id(ir); irpass::print(ir); } @@ -1014,7 +1014,7 @@ void GPUCodeGen::lower_llvm() { irpass::demote_dense_struct_fors(ir); irpass::typecheck(ir); if (print_ir) { - TC_TRACE("Dense Struct-for demoted:"); + TI_TRACE("Dense Struct-for demoted:"); irpass::print(ir); } } @@ -1023,85 +1023,85 @@ void GPUCodeGen::lower_llvm() { irpass::simplify(ir); irpass::re_id(ir); if (print_ir) { - TC_TRACE("Simplified I:"); + TI_TRACE("Simplified I:"); irpass::print(ir); } } if (kernel->grad) { // irpass::re_id(ir); - // TC_TRACE("Primal:"); + // TI_TRACE("Primal:"); // irpass::print(ir); irpass::demote_atomics(ir); irpass::full_simplify(ir); irpass::typecheck(ir); if (print_ir) { - TC_TRACE("Before make_adjoint:"); + TI_TRACE("Before make_adjoint:"); irpass::print(ir); } irpass::make_adjoint(ir); if (print_ir) { - TC_TRACE("After make_adjoint:"); + TI_TRACE("After make_adjoint:"); irpass::print(ir); } irpass::typecheck(ir); // irpass::re_id(ir); - // TC_TRACE("Adjoint:"); + // TI_TRACE("Adjoint:"); // irpass::print(ir); } irpass::lower_access(ir, prog->config.use_llvm); if (print_ir) { - TC_TRACE("Access Lowered:"); + TI_TRACE("Access Lowered:"); irpass::re_id(ir); irpass::print(ir); } if (prog->config.simplify_after_lower_access) { irpass::die(ir); if (print_ir) { - TC_TRACE("DIEd:"); + TI_TRACE("DIEd:"); irpass::re_id(ir); irpass::print(ir); } irpass::simplify(ir); if (print_ir) { - TC_TRACE("Simplified II:"); + TI_TRACE("Simplified II:"); irpass::re_id(ir); irpass::print(ir); } } irpass::die(ir); if (print_ir) { - TC_TRACE("DIEd:"); + TI_TRACE("DIEd:"); irpass::re_id(ir); irpass::print(ir); } irpass::flag_access(ir); if (print_ir) { - TC_TRACE("Access Flagged:"); + TI_TRACE("Access Flagged:"); irpass::re_id(ir); irpass::print(ir); } irpass::offload(ir); if (print_ir) { - TC_TRACE("Offloaded:"); + TI_TRACE("Offloaded:"); irpass::re_id(ir); irpass::print(ir); } irpass::demote_atomics(ir); if (print_ir) { - TC_TRACE("Atomics Demoted:"); + TI_TRACE("Atomics Demoted:"); irpass::re_id(ir); irpass::print(ir); } irpass::full_simplify(ir); if (print_ir) { - TC_TRACE("Simplified III:"); + TI_TRACE("Simplified III:"); irpass::re_id(ir); irpass::print(ir); } } void GPUCodeGen::lower() { - TC_PROFILER(__FUNCTION__) + TI_PROFILER(__FUNCTION__) if (prog->config.use_llvm) { lower_llvm(); } else { @@ -1110,7 +1110,7 @@ void GPUCodeGen::lower() { } void GPUCodeGen::codegen() { - emit("#define TC_GPU"); + emit("#define TI_GPU"); generate_header(); // Body diff --git a/taichi/backends/codegen_llvm.h b/taichi/backends/codegen_llvm.h index 8dfc1da42af2f..358394f2c617d 100644 --- a/taichi/backends/codegen_llvm.h +++ b/taichi/backends/codegen_llvm.h @@ -77,15 +77,15 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } void compile() { - TC_ASSERT(!func); + TI_ASSERT(!func); auto kernel_symbol = codegen->jit->lookup(name); - TC_ASSERT_INFO(kernel_symbol, "Function not found"); + TI_ASSERT_INFO(kernel_symbol, "Function not found"); func = (task_fp_type)(void *)(llvm::cantFail(kernel_symbol.getAddress())); } void operator()(Context *context) { - TC_ASSERT(func); + TI_ASSERT(func); func(context); } }; @@ -199,8 +199,8 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { meta->call("set_bitmasked", tlctx->get_constant(snode->_bitmasked)); meta->call("set_morton_dim", tlctx->get_constant((int)snode->_morton)); } else if (snode->type == SNodeType::pointer) { - meta = std::make_unique("PointerMeta", this, builder); - emit_struct_meta_base("Pointer", meta->ptr, snode); + meta = std::make_unique("pointerMeta", this, builder); + emit_struct_meta_base("pointer", meta->ptr, snode); } else if (snode->type == SNodeType::root) { meta = std::make_unique("RootMeta", this, builder); emit_struct_meta_base("Root", meta->ptr, snode); @@ -209,8 +209,8 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { emit_struct_meta_base("Dynamic", meta->ptr, snode); meta->call("set_chunk_size", tlctx->get_constant(snode->chunk_size)); } else { - TC_P(snode_type_name(snode->type)); - TC_NOT_IMPLEMENTED; + TI_P(snode_type_name(snode->type)); + TI_NOT_IMPLEMENTED; } if (false) { // auto ptr_type = llvm::Type::getInt8PtrTy(*llvm_context, 0); @@ -225,7 +225,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { llvm::Value *emit_struct_meta(SNode *snode) { auto obj = emit_struct_meta_object(snode); - TC_ASSERT(obj != nullptr); + TI_ASSERT(obj != nullptr); return obj->ptr; } @@ -254,7 +254,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { template void emit(std::string f, Args &&... args) { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED codegen->emit(f, std::forward(args)...); } @@ -265,7 +265,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } void visit(AllocaStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); stmt->value = create_entry_block_alloca(stmt->ret_type.data_type); // initialize as zero builder->CreateStore(tlctx->get_constant(stmt->ret_type.data_type, 0), @@ -295,7 +295,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { stmt->value = \ builder->CreateCall(get_runtime_function(#x "_i32"), input); \ } else { \ - TC_NOT_IMPLEMENTED \ + TI_NOT_IMPLEMENTED \ } \ } if (false) { @@ -316,8 +316,8 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { {input_type}, {input}); } else { - TC_P(unary_op_type_name(op)); - TC_NOT_IMPLEMENTED + TI_P(unary_op_type_name(op)); + TI_NOT_IMPLEMENTED } #undef UNARY_STD } @@ -359,16 +359,16 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { llvm::CastInst::CastOps cast_op; auto from = stmt->operand->ret_type.data_type; auto to = stmt->cast_type; - TC_ASSERT(from != to); + TI_ASSERT(from != to); if (is_real(from) != is_real(to)) { if (is_real(from) && is_integral(to)) { cast_op = llvm::Instruction::CastOps::FPToSI; } else if (is_integral(from) && is_real(to)) { cast_op = llvm::Instruction::CastOps::SIToFP; } else { - TC_P(data_type_name(from)); - TC_P(data_type_name(to)); - TC_NOT_IMPLEMENTED; + TI_P(data_type_name(from)); + TI_P(data_type_name(to)); + TI_NOT_IMPLEMENTED; } stmt->value = builder->CreateCast(cast_op, stmt->operand->value, @@ -391,7 +391,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } } } else { - TC_ASSERT(data_type_size(stmt->ret_type.data_type) == + TI_ASSERT(data_type_size(stmt->ret_type.data_type) == data_type_size(stmt->cast_type)); stmt->value = builder->CreateBitCast( stmt->operand->value, tlctx->get_data_type(stmt->cast_type)); @@ -409,7 +409,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } else if (dt == DataType::f64) { return llvm::Type::getDoubleTy(*llvm_context); } else { - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; } return nullptr; } @@ -466,8 +466,8 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { stmt->value = create_call("max_i32", {stmt->lhs->value, stmt->rhs->value}); } else { - TC_P(data_type_name(ret_type)); - TC_NOT_IMPLEMENTED + TI_P(data_type_name(ret_type)); + TI_NOT_IMPLEMENTED } } else if (op == BinaryOpType::atan2) { if (current_arch() == Arch::x86_64) { @@ -478,8 +478,8 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { stmt->value = create_call("atan2_f64", {stmt->lhs->value, stmt->rhs->value}); } else { - TC_P(data_type_name(ret_type)); - TC_NOT_IMPLEMENTED + TI_P(data_type_name(ret_type)); + TI_NOT_IMPLEMENTED } } else if (current_arch() == Arch::cuda) { if (ret_type == DataType::f32) { @@ -489,11 +489,11 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { stmt->value = create_call("__nv_atan2", {stmt->lhs->value, stmt->rhs->value}); } else { - TC_P(data_type_name(ret_type)); - TC_NOT_IMPLEMENTED + TI_P(data_type_name(ret_type)); + TI_NOT_IMPLEMENTED } } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } } else if (op == BinaryOpType::pow) { if (current_arch() == Arch::x86_64) { @@ -510,8 +510,8 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { stmt->value = create_call("pow_i64", {stmt->lhs->value, stmt->rhs->value}); } else { - TC_P(data_type_name(ret_type)); - TC_NOT_IMPLEMENTED + TI_P(data_type_name(ret_type)); + TI_NOT_IMPLEMENTED } } else if (current_arch() == Arch::cuda) { if (ret_type == DataType::f32) { @@ -527,11 +527,11 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { stmt->value = create_call("pow_i64", {stmt->lhs->value, stmt->rhs->value}); } else { - TC_P(data_type_name(ret_type)); - TC_NOT_IMPLEMENTED + TI_P(data_type_name(ret_type)); + TI_NOT_IMPLEMENTED } } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } } else if (op == BinaryOpType::min) { if (is_real(ret_type)) { @@ -540,8 +540,8 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { stmt->value = create_call("min_i32", {stmt->lhs->value, stmt->rhs->value}); } else { - TC_P(data_type_name(ret_type)); - TC_NOT_IMPLEMENTED + TI_P(data_type_name(ret_type)); + TI_NOT_IMPLEMENTED } } else if (is_comparison(op)) { llvm::Value *cmp = nullptr; @@ -599,17 +599,17 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { cmp = builder->CreateICmpNE(stmt->lhs->value, stmt->rhs->value); } } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } stmt->value = builder->CreateSExt(cmp, llvm_type(DataType::i32)); } else { - TC_P(binary_op_type_name(op)); - TC_NOT_IMPLEMENTED + TI_P(binary_op_type_name(op)); + TI_NOT_IMPLEMENTED } } void visit(TernaryOpStmt *stmt) override { - TC_ASSERT(stmt->op_type == TernaryOpType::select); + TI_ASSERT(stmt->op_type == TernaryOpType::select); stmt->value = builder->CreateSelect( builder->CreateTrunc(stmt->op1->value, llvm_type(DataType::i1)), stmt->op2->value, stmt->op3->value); @@ -644,7 +644,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { if (dt == DataType::i32) { format = "%d"; } else if (dt == DataType::i64) { -#if defined(TC_PLATFORM_UNIX) +#if defined(TI_PLATFORM_UNIX) format = "%lld"; #else format = "%I64d"; @@ -655,7 +655,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } else if (dt == DataType::f64) { format = "%.12f"; } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } args.push_back(builder->CreateGlobalStringPtr( ("[llvm codegen debug] " + tag + " = " + format + "\n").c_str(), @@ -666,7 +666,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } void visit(PrintStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); std::vector args; std::string format; auto value = stmt->stmt->value; @@ -674,7 +674,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { if (dt == DataType::i32) { format = "%d"; } else if (dt == DataType::i64) { -#if defined(TC_PLATFORM_UNIX) +#if defined(TI_PLATFORM_UNIX) format = "%lld"; #else format = "%I64d"; @@ -685,7 +685,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } else if (dt == DataType::f64) { format = "%.12f"; } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } args.push_back(builder->CreateGlobalStringPtr( ("[debug] " + stmt->str + " = " + format + "\n").c_str(), @@ -697,7 +697,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } void visit(ConstStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); auto val = stmt->val[0]; if (val.dt == DataType::f32) { stmt->value = llvm::ConstantFP::get(*llvm_context, @@ -712,14 +712,14 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { stmt->value = llvm::ConstantInt::get( *llvm_context, llvm::APInt(64, val.val_int64(), true)); } else { - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; } } void visit(WhileControlStmt *stmt) override { BasicBlock *after_break = BasicBlock::Create(*llvm_context, "after_break", func); - TC_ASSERT(while_after_loop); + TI_ASSERT(while_after_loop); auto cond = builder->CreateICmpEQ(stmt->cond->value, tlctx->get_constant(0)); builder->CreateCondBr(cond, while_after_loop, after_break); @@ -865,7 +865,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { void visit(ArgStoreStmt *stmt) override { if (stmt->is_ptr) { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } else { auto intermediate_bits = tlctx->get_data_type(stmt->val->ret_type.data_type) @@ -882,14 +882,14 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } void visit(LocalLoadStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); stmt->value = builder->CreateLoad(stmt->ptr[0].var->value); } void visit(LocalStoreStmt *stmt) override { auto mask = stmt->parent->mask(); if (mask && stmt->width() != 1) { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } else { builder->CreateStore(stmt->data->value, stmt->ptr->value); } @@ -903,30 +903,33 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { void visit(SNodeOpStmt *stmt) override { auto snode = stmt->snode; if (stmt->op_type == SNodeOpType::append) { - TC_ASSERT(snode->type == SNodeType::dynamic); - TC_ASSERT(stmt->ret_type.data_type == DataType::i32); + TI_ASSERT(snode->type == SNodeType::dynamic); + TI_ASSERT(stmt->ret_type.data_type == DataType::i32); stmt->value = call(snode, stmt->ptr->value, "append", {stmt->val->value}); } else if (stmt->op_type == SNodeOpType::length) { - TC_ASSERT(snode->type == SNodeType::dynamic); + TI_ASSERT(snode->type == SNodeType::dynamic); stmt->value = call(snode, stmt->ptr->value, "get_num_elements", {}); } else if (stmt->op_type == SNodeOpType::is_active) { stmt->value = call(snode, stmt->ptr->value, "is_active", {stmt->val->value}); } else if (stmt->op_type == SNodeOpType::deactivate) { - TC_ASSERT(snode->type == SNodeType::pointer || - snode->type == SNodeType::dynamic); - stmt->value = call(snode, stmt->ptr->value, "deactivate", {}); + if (snode->type == SNodeType::pointer || snode->type == SNodeType::hash) { + stmt->value = + call(snode, stmt->ptr->value, "deactivate", {stmt->val->value}); + } else if (snode->type == SNodeType::dynamic) { + stmt->value = call(snode, stmt->ptr->value, "deactivate", {}); + } } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } } void visit(AtomicOpStmt *stmt) override { // auto mask = stmt->parent->mask(); // TODO: deal with mask when vectorized - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); for (int l = 0; l < stmt->width(); l++) { - TC_ASSERT(stmt->op_type == AtomicOpType::add); + TI_ASSERT(stmt->op_type == AtomicOpType::add); llvm::Value *old_value; if (stmt->val->ret_type.data_type == DataType::i32) old_value = builder->CreateAtomicRMW( @@ -939,26 +942,26 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { old_value = builder->CreateCall(get_runtime_function("atomic_add_f64"), {stmt->dest->value, stmt->val->value}); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } stmt->value = old_value; } } void visit(GlobalPtrStmt *stmt) override { - TC_ERROR("Global Ptrs should have been lowered."); + TI_ERROR("Global Ptrs should have been lowered."); } void visit(GlobalStoreStmt *stmt) override { - TC_ASSERT(!stmt->parent->mask() || stmt->width() == 1); - TC_ASSERT(stmt->data->value); - TC_ASSERT(stmt->ptr->value); + TI_ASSERT(!stmt->parent->mask() || stmt->width() == 1); + TI_ASSERT(stmt->data->value); + TI_ASSERT(stmt->ptr->value); builder->CreateStore(stmt->data->value, stmt->ptr->value); } void visit(GlobalLoadStmt *stmt) override { int width = stmt->width(); - TC_ASSERT(width == 1); + TI_ASSERT(width == 1); stmt->value = builder->CreateLoad( tlctx->get_data_type(stmt->ret_type.data_type), stmt->ptr->value); } @@ -986,12 +989,12 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } else if (snode->type == SNodeType::dynamic) { return "Dynamic"; } else if (snode->type == SNodeType::pointer) { - return "Pointer"; + return "pointer"; } else if (snode->type == SNodeType::hash) { return "Hash"; } else { - TC_P(snode_type_name(snode->type)); - TC_NOT_IMPLEMENTED + TI_P(snode_type_name(snode->type)); + TI_NOT_IMPLEMENTED } } @@ -1008,6 +1011,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { llvm::Type::getInt8PtrTy(*llvm_context)); std::vector func_arguments{s_ptr, node_ptr}; + func_arguments.insert(func_arguments.end(), arguments.begin(), arguments.end()); @@ -1040,7 +1044,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } void visit(IntegerOffsetStmt *stmt) override { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED if (stmt->input->is() && stmt->input->as()->output_snode->type == SNodeType::place) { auto input = stmt->input->as(); @@ -1056,7 +1060,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { void visit(SNodeLookupStmt *stmt) override { llvm::Value *parent = nullptr; parent = stmt->input_snode->value; - TC_ASSERT(parent); + TI_ASSERT(parent); auto snode = stmt->snode; if (snode->type == SNodeType::root) { stmt->value = builder->CreateGEP(parent, stmt->input_index->value); @@ -1070,8 +1074,8 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { stmt->value = call(snode, stmt->input_snode->value, "lookup_element", {stmt->input_index->value}); } else { - TC_INFO(snode_type_name(snode->type)); - TC_NOT_IMPLEMENTED + TI_INFO(snode_type_name(snode->type)); + TI_NOT_IMPLEMENTED } } @@ -1085,7 +1089,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } void visit(ExternalPtrStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); auto argload = stmt->base_ptrs[0]->as(); auto arg_id = argload->arg_id; @@ -1162,11 +1166,11 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { builder->CreateBr(func_body_bb); if (prog->config.print_kernel_llvm_ir) { - TC_INFO("Kernel Module IR"); + TI_INFO("Kernel Module IR"); module->print(errs(), nullptr); } - TC_ASSERT(!llvm::verifyFunction(*func, &errs())); - // TC_INFO("Kernel function verified."); + TI_ASSERT(!llvm::verifyFunction(*func, &errs())); + // TI_INFO("Kernel function verified."); } class FunctionCreationGuard { @@ -1275,7 +1279,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { void create_offload_struct_for(OffloadedStmt *stmt, bool spmd = false) { llvm::Function *body; - auto leaf_block = stmt->snode->parent; + auto leaf_block = stmt->snode; { // Create the loop body function auto guard = get_function_creation_gurad({ @@ -1408,7 +1412,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { auto buffer = call("Runtime_get_temporary_pointer", runtime, tlctx->get_constant((int64)stmt->offset)); - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); auto ptr_type = llvm::PointerType::get( tlctx->get_data_type(stmt->ret_type.data_type), 0); stmt->value = builder->CreatePointerCast(buffer, ptr_type); @@ -1442,7 +1446,7 @@ class CodeGenLLVM : public IRVisitor, public ModuleBuilder { } else if (stmt->task_type == Type::gc) { emit_gc(stmt); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } if (prog->config.enable_profiler) { call(builder, "Runtime_profiler_stop", {get_runtime()}); diff --git a/taichi/backends/codegen_llvm_ptx.cpp b/taichi/backends/codegen_llvm_ptx.cpp index edb4037301d41..e5921fc198fe1 100644 --- a/taichi/backends/codegen_llvm_ptx.cpp +++ b/taichi/backends/codegen_llvm_ptx.cpp @@ -72,17 +72,17 @@ class CodeGenLLVMGPU : public CodeGenLLVM { auto offloaded_local = offloaded_tasks; for (auto &task : offloaded_local) { llvm::Function *func = module->getFunction(task.name); - TC_ASSERT(func); + TI_ASSERT(func); mark_function_as_cuda_kernel(func); } if (prog->config.print_kernel_llvm_ir) { - TC_INFO("IR before global optimization"); + TI_INFO("IR before global optimization"); module->print(errs(), nullptr); } auto ptx = compile_module_to_ptx(module); if (prog->config.print_kernel_llvm_ir_optimized) { - TC_P(ptx); + TI_P(ptx); } auto cuda_module = cuda_context->compile(ptx); @@ -93,7 +93,7 @@ class CodeGenLLVMGPU : public CodeGenLLVM { auto prog = this->prog; return [offloaded_local, prog](Context context) { for (auto task : offloaded_local) { - TC_DEBUG("Launching kernel {}<<<{}, {}>>>", task.name, task.grid_dim, + TI_DEBUG("Launching kernel {}<<<{}, {}>>>", task.name, task.grid_dim, task.block_dim); ProfilerBase *profiler = nullptr; @@ -105,13 +105,13 @@ class CodeGenLLVMGPU : public CodeGenLLVM { } }; #else - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; return nullptr; #endif } void visit(PrintStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); auto value_type = tlctx->get_data_type(stmt->stmt->ret_type.data_type); @@ -130,7 +130,7 @@ class CodeGenLLVMGPU : public CodeGenLLVM { } else if (stmt->stmt->ret_type.data_type == DataType::f64) { format = "%.12f"; } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } std::vector types{value_type}; @@ -166,7 +166,7 @@ class CodeGenLLVMGPU : public CodeGenLLVM { } else if (input_taichi_type == DataType::i32) { \ stmt->value = builder->CreateCall(get_runtime_function(#x), input); \ } else { \ - TC_NOT_IMPLEMENTED \ + TI_NOT_IMPLEMENTED \ } \ } if (op == UnaryOpType::abs) { @@ -180,7 +180,7 @@ class CodeGenLLVMGPU : public CodeGenLLVM { stmt->value = builder->CreateCall(get_runtime_function("__nv_abs"), input); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } } else if (op == UnaryOpType::sqrt) { if (input_taichi_type == DataType::f32) { @@ -190,14 +190,14 @@ class CodeGenLLVMGPU : public CodeGenLLVM { stmt->value = builder->CreateCall(get_runtime_function("__nv_sqrt"), input); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } } else if (op == UnaryOpType::logic_not) { if (input_taichi_type == DataType::i32) { stmt->value = builder->CreateCall(get_runtime_function("logic_not_i32"), input); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } } UNARY_STD(exp) @@ -210,21 +210,21 @@ class CodeGenLLVMGPU : public CodeGenLLVM { UNARY_STD(cos) UNARY_STD(sin) else { - TC_P(unary_op_type_name(op)); - TC_NOT_IMPLEMENTED + TI_P(unary_op_type_name(op)); + TI_NOT_IMPLEMENTED } #undef UNARY_STD } void visit(AtomicOpStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); // https://llvm.org/docs/NVPTXUsage.html#address-spaces bool is_local = stmt->dest->is(); if (is_local) { - TC_ERROR("Local atomics should have been demoted."); + TI_ERROR("Local atomics should have been demoted."); } else { for (int l = 0; l < stmt->width(); l++) { - TC_ASSERT(stmt->op_type == AtomicOpType::add); + TI_ASSERT(stmt->op_type == AtomicOpType::add); llvm::Value *old_value; if (is_integral(stmt->val->ret_type.data_type)) { old_value = builder->CreateAtomicRMW( @@ -243,7 +243,7 @@ class CodeGenLLVMGPU : public CodeGenLLVM { {llvm::PointerType::get(dt, 0)}, {stmt->dest->value, stmt->val->value}); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } stmt->value = old_value; } @@ -344,7 +344,7 @@ class CodeGenLLVMGPU : public CodeGenLLVM { if (kernel_block_dim == 0) kernel_block_dim = prog->config.default_gpu_block_dim; kernel_block_dim = - std::min(stmt->snode->parent->max_num_elements(), kernel_block_dim); + std::min(stmt->snode->max_num_elements(), kernel_block_dim); stmt->block_dim = kernel_block_dim; create_offload_struct_for(stmt, true); } else if (stmt->task_type == Type::clear_list) { @@ -355,7 +355,7 @@ class CodeGenLLVMGPU : public CodeGenLLVM { kernel_block_dim = std::min(branching, 64); emit_list_gen(stmt); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } finalize_offloaded_task_function(); current_task->grid_dim = kernel_grid_dim; @@ -364,13 +364,13 @@ class CodeGenLLVMGPU : public CodeGenLLVM { current_task = nullptr; } #else - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED #endif } }; FunctionType GPUCodeGen::codegen_llvm() { - TC_PROFILER("cuda codegen"); + TI_PROFILER("cuda codegen"); return CodeGenLLVMGPU(this, kernel).gen(); } diff --git a/taichi/backends/codegen_llvm_x86.cpp b/taichi/backends/codegen_llvm_x86.cpp index 84ccca6053e06..a70831bc49884 100644 --- a/taichi/backends/codegen_llvm_x86.cpp +++ b/taichi/backends/codegen_llvm_x86.cpp @@ -24,7 +24,7 @@ class CodeGenLLVMCPU : public CodeGenLLVM { }; FunctionType CPUCodeGen::codegen_llvm() { - TC_PROFILER("cpu codegen"); + TI_PROFILER("cpu codegen"); return CodeGenLLVMCPU(this, kernel).gen(); } @@ -32,7 +32,7 @@ void global_optimize_module_x86_64(std::unique_ptr &module) { TI_AUTO_PROF auto JTMB = JITTargetMachineBuilder::detectHost(); if (!JTMB) { - TC_ERROR("Target machine creation failed."); + TI_ERROR("Target machine creation failed."); } module->setTargetTriple(JTMB->getTargetTriple().str()); llvm::Triple triple(module->getTargetTriple()); @@ -40,7 +40,7 @@ void global_optimize_module_x86_64(std::unique_ptr &module) { std::string err_str; const llvm::Target *target = TargetRegistry::lookupTarget(triple.str(), err_str); - TC_ERROR_UNLESS(target, err_str); + TI_ERROR_UNLESS(target, err_str); TargetOptions options; options.PrintMachineCode = false; @@ -69,7 +69,7 @@ void global_optimize_module_x86_64(std::unique_ptr &module) { triple.str(), mcpu.str(), "", options, llvm::Reloc::PIC_, llvm::CodeModel::Small, CodeGenOpt::Aggressive)); - TC_ERROR_UNLESS(target_machine.get(), "Could not allocate target machine!"); + TI_ERROR_UNLESS(target_machine.get(), "Could not allocate target machine!"); module->setDataLayout(target_machine->createDataLayout()); @@ -98,9 +98,9 @@ void global_optimize_module_x86_64(std::unique_ptr &module) { auto t = Time::get_time(); module_pass_manager.run(*module); t = Time::get_time() - t; - // TC_INFO("Global optimization time: {} ms", t * 1000); + // TI_INFO("Global optimization time: {} ms", t * 1000); if (get_current_program().config.print_kernel_llvm_ir_optimized) { - TC_INFO("Global optimized IR:"); + TI_INFO("Global optimized IR:"); module->print(llvm::errs(), nullptr); } } diff --git a/taichi/backends/codegen_metal.cpp b/taichi/backends/codegen_metal.cpp index ebe7af5bca16b..a747fff03f298 100644 --- a/taichi/backends/codegen_metal.cpp +++ b/taichi/backends/codegen_metal.cpp @@ -50,7 +50,7 @@ class MetalKernelCodegen : public IRVisitor { } void visit(ConstStmt *const_stmt) override { - TC_ASSERT(const_stmt->width() == 1); + TI_ASSERT(const_stmt->width() == 1); emit("const {} {} = {};", metal_data_type_name(const_stmt->element_type()), const_stmt->raw_name(), const_stmt->val[0].stringify()); } @@ -69,7 +69,7 @@ class MetalKernelCodegen : public IRVisitor { emit("const {} {}({});", metal_data_type_name(stmt->element_type()), stmt->raw_name(), ptr->raw_name()); } else { - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; } } @@ -114,7 +114,7 @@ class MetalKernelCodegen : public IRVisitor { if (stmt->input_snode) { parent = stmt->input_snode->raw_name(); } else { - TC_ASSERT(root_stmt_ != nullptr); + TI_ASSERT(root_stmt_ != nullptr); parent = root_stmt_->raw_name(); } @@ -123,12 +123,12 @@ class MetalKernelCodegen : public IRVisitor { } void visit(GlobalStoreStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); emit(R"(*{} = {};)", stmt->ptr->raw_name(), stmt->data->raw_name()); } void visit(GlobalLoadStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); emit(R"({} {} = *{};)", metal_data_type_name(stmt->element_type()), stmt->raw_name(), stmt->ptr->raw_name()); } @@ -146,7 +146,7 @@ class MetalKernelCodegen : public IRVisitor { void visit(ArgStoreStmt *stmt) override { const auto dt = metal_data_type_name(stmt->element_type()); - TC_ASSERT(!stmt->is_ptr); + TI_ASSERT(!stmt->is_ptr); emit("*{}.arg{}() = {};", kArgsContextName, stmt->arg_id, stmt->val->raw_name()); } @@ -154,7 +154,7 @@ class MetalKernelCodegen : public IRVisitor { void visit(ExternalPtrStmt *stmt) override { // Used mostly for transferring data between host (e.g. numpy array) and // Metal. - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); const auto linear_index_name = fmt::format("{}_linear_index_", stmt->raw_name()); emit("int {} = 0;", linear_index_name); @@ -184,16 +184,16 @@ class MetalKernelCodegen : public IRVisitor { } void visit(GlobalTemporaryStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); const auto dt = metal_data_type_name(stmt->element_type()); emit("device {}* {} = reinterpret_cast({} + {});", dt, stmt->raw_name(), dt, kGlobalTmpsBufferName, stmt->offset); } void visit(LoopIndexStmt *stmt) override { - TC_ASSERT(current_kernel_attribs_->task_type == + TI_ASSERT(current_kernel_attribs_->task_type == OffloadedStmt::TaskType::range_for); - TC_ASSERT(!stmt->is_struct_for && stmt->index == 0); + TI_ASSERT(!stmt->is_struct_for && stmt->index == 0); if (current_kernel_attribs_->range_for_attribs.const_begin) { emit("const int {} = (static_cast({}) + {});", stmt->raw_name(), kKernelThreadIdName, @@ -221,7 +221,7 @@ class MetalKernelCodegen : public IRVisitor { // reinterpret the bit pattern const auto to_type = to_metal_type(stmt->cast_type); const auto to_type_name = metal_data_type_name(to_type); - TC_ASSERT(metal_data_type_bytes( + TI_ASSERT(metal_data_type_bytes( to_metal_type(stmt->operand->element_type())) == metal_data_type_bytes(to_type)); emit("const {} {} = union_cast<{}>({});", to_type_name, @@ -257,15 +257,15 @@ class MetalKernelCodegen : public IRVisitor { } void visit(TernaryOpStmt *tri) override { - TC_ASSERT(tri->op_type == TernaryOpType::select); + TI_ASSERT(tri->op_type == TernaryOpType::select); emit("const {} {} = ({}) ? ({}) : ({});", metal_data_type_name(tri->element_type()), tri->raw_name(), tri->op1->raw_name(), tri->op2->raw_name(), tri->op3->raw_name()); } void visit(AtomicOpStmt *stmt) override { - TC_ASSERT(stmt->width() == 1); - TC_ASSERT(stmt->op_type == AtomicOpType::add); + TI_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->op_type == AtomicOpType::add); const auto dt = stmt->val->element_type(); if (dt == DataType::i32) { emit( @@ -277,7 +277,7 @@ class MetalKernelCodegen : public IRVisitor { emit("const float {} = fatomic_fetch_add({}, {});", stmt->raw_name(), stmt->dest->raw_name(), stmt->val->raw_name()); } else { - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; } } @@ -294,7 +294,7 @@ class MetalKernelCodegen : public IRVisitor { } void visit(RangeForStmt *for_stmt) override { - TC_ASSERT(for_stmt->width() == 1); + TI_ASSERT(for_stmt->width() == 1); auto *loop_var = for_stmt->loop_var; if (loop_var->ret_type.data_type == DataType::i32) { if (!for_stmt->reversed) { @@ -312,7 +312,7 @@ class MetalKernelCodegen : public IRVisitor { emit(" int {} = {}_;", loop_var->raw_name(), loop_var->raw_name()); } } else { - TC_ASSERT(!for_stmt->reversed); + TI_ASSERT(!for_stmt->reversed); const auto type_name = metal_data_type_name(loop_var->element_type()); emit("for ({} {} = {}; {} < {}; {} = {} + ({})1) {{", type_name, loop_var->raw_name(), for_stmt->begin->raw_name(), @@ -324,11 +324,11 @@ class MetalKernelCodegen : public IRVisitor { } void visit(StructForStmt *) override { - TC_ERROR("Struct for cannot be nested."); + TI_ERROR("Struct for cannot be nested."); } void visit(OffloadedStmt *stmt) override { - TC_ASSERT(is_top_level_); + TI_ASSERT(is_top_level_); is_top_level_ = false; using Type = OffloadedStmt::TaskType; if (stmt->task_type == Type::serial) { @@ -338,7 +338,7 @@ class MetalKernelCodegen : public IRVisitor { } else { // struct_for is automatically lowered to ranged_for for dense snodes // (#378). So we only need to support serial and range_for tasks. - TC_ERROR("Unsupported offload type={} on Metal arch", stmt->task_name()); + TI_ERROR("Unsupported offload type={} on Metal arch", stmt->task_name()); } is_top_level_ = true; } @@ -354,12 +354,12 @@ class MetalKernelCodegen : public IRVisitor { } void visit(RandStmt *stmt) override { - TC_ERROR("Metal arch doesn't support ti.random() yet"); + TI_ERROR("Metal arch doesn't support ti.random() yet"); } void visit(PrintStmt *stmt) override { // TODO: Add a flag to control whether ignoring print() stmt is allowed. - TC_WARN("Cannot print inside Metal kernel, ignored"); + TI_WARN("Cannot print inside Metal kernel, ignored"); } private: @@ -376,10 +376,10 @@ class MetalKernelCodegen : public IRVisitor { } void generate_common_functions() { -#define TC_INSIDE_METAL_CODEGEN +#define TI_INSIDE_METAL_CODEGEN #include kernel_src_code_ += kMetalHelpersSourceCode; -#undef TC_INSIDE_METAL_CODEGEN +#undef TI_INSIDE_METAL_CODEGEN emit("\n"); } @@ -425,7 +425,7 @@ class MetalKernelCodegen : public IRVisitor { } void generate_serial_kernel(OffloadedStmt *stmt) { - TC_ASSERT(stmt->task_type == OffloadedStmt::TaskType::serial); + TI_ASSERT(stmt->task_type == OffloadedStmt::TaskType::serial); const std::string mtl_kernel_name = make_kernel_name(); emit_mtl_kernel_func_sig(mtl_kernel_name); emit(" // serial"); @@ -445,7 +445,7 @@ class MetalKernelCodegen : public IRVisitor { } void generate_range_for_kernel(OffloadedStmt *stmt) { - TC_ASSERT(stmt->task_type == OffloadedStmt::TaskType::range_for); + TI_ASSERT(stmt->task_type == OffloadedStmt::TaskType::range_for); const std::string mtl_kernel_name = make_kernel_name(); emit_mtl_kernel_func_sig(mtl_kernel_name); @@ -570,7 +570,7 @@ void MetalCodeGen::lower() { auto ir = kernel_->ir; const bool print_ir = prog_->config.print_ir; if (print_ir) { - TC_TRACE("Initial IR:"); + TI_TRACE("Initial IR:"); irpass::print(ir); } @@ -578,7 +578,7 @@ void MetalCodeGen::lower() { irpass::reverse_segments(ir); irpass::re_id(ir); if (print_ir) { - TC_TRACE("Segment reversed (for autodiff):"); + TI_TRACE("Segment reversed (for autodiff):"); irpass::print(ir); } } @@ -586,21 +586,21 @@ void MetalCodeGen::lower() { irpass::lower(ir); irpass::re_id(ir); if (print_ir) { - TC_TRACE("Lowered:"); + TI_TRACE("Lowered:"); irpass::print(ir); } irpass::typecheck(ir); irpass::re_id(ir); if (print_ir) { - TC_TRACE("Typechecked:"); + TI_TRACE("Typechecked:"); irpass::print(ir); } irpass::demote_dense_struct_fors(ir); irpass::typecheck(ir); if (print_ir) { - TC_TRACE("Dense Struct-for demoted:"); + TI_TRACE("Dense Struct-for demoted:"); irpass::print(ir); } @@ -609,7 +609,7 @@ void MetalCodeGen::lower() { irpass::simplify(ir); irpass::re_id(ir); if (print_ir) { - TC_TRACE("Simplified I:"); + TI_TRACE("Simplified I:"); irpass::print(ir); } } @@ -619,12 +619,12 @@ void MetalCodeGen::lower() { irpass::full_simplify(ir); irpass::typecheck(ir); if (print_ir) { - TC_TRACE("Before make_adjoint:"); + TI_TRACE("Before make_adjoint:"); irpass::print(ir); } irpass::make_adjoint(ir); if (print_ir) { - TC_TRACE("After make_adjoint:"); + TI_TRACE("After make_adjoint:"); irpass::print(ir); } irpass::typecheck(ir); @@ -633,27 +633,27 @@ void MetalCodeGen::lower() { irpass::lower_access(ir, prog_->config.use_llvm); irpass::re_id(ir); if (print_ir) { - TC_TRACE("Access Lowered:"); + TI_TRACE("Access Lowered:"); irpass::print(ir); } irpass::die(ir); irpass::re_id(ir); if (print_ir) { - TC_TRACE("DIEd:"); + TI_TRACE("DIEd:"); irpass::print(ir); } irpass::flag_access(ir); irpass::re_id(ir); if (print_ir) { - TC_TRACE("Access Flagged:"); + TI_TRACE("Access Flagged:"); irpass::print(ir); } irpass::constant_fold(ir); if (print_ir) { - TC_TRACE("Constant folded:"); + TI_TRACE("Constant folded:"); irpass::re_id(ir); irpass::print(ir); } @@ -661,21 +661,21 @@ void MetalCodeGen::lower() { global_tmps_buffer_size_ = std::max(irpass::offload(ir).total_size, (size_t)(1)); if (print_ir) { - TC_TRACE("Offloaded:"); + TI_TRACE("Offloaded:"); irpass::re_id(ir); irpass::print(ir); } irpass::full_simplify(ir); if (print_ir) { - TC_TRACE("Simplified II:"); + TI_TRACE("Simplified II:"); irpass::re_id(ir); irpass::print(ir); } irpass::demote_atomics(ir); if (print_ir) { - TC_TRACE("Atomics demoted:"); + TI_TRACE("Atomics demoted:"); irpass::re_id(ir); irpass::print(ir); } diff --git a/taichi/backends/codegen_x86.cpp b/taichi/backends/codegen_x86.cpp index d0fd469ca4b08..22704dc7db275 100644 --- a/taichi/backends/codegen_x86.cpp +++ b/taichi/backends/codegen_x86.cpp @@ -46,7 +46,7 @@ class CPUIRCodeGen : public IRVisitor { } void visit(RandStmt *stmt) { - TC_ASSERT(stmt->ret_type.data_type == DataType::f32); + TI_ASSERT(stmt->ret_type.data_type == DataType::f32); emit("const auto {} = {}::rand();", stmt->raw_name(), stmt->ret_data_type_name()); } @@ -122,7 +122,7 @@ class CPUIRCodeGen : public IRVisitor { } void visit(StructForStmt *for_stmt) { - TC_ASSERT_INFO(current_struct_for == nullptr, + TI_ASSERT_INFO(current_struct_for == nullptr, "Struct for cannot be nested."); current_struct_for = for_stmt; emit("{{"); @@ -209,7 +209,7 @@ class CPUIRCodeGen : public IRVisitor { emit("{} = {}_;", loop_var->raw_name(), loop_var->raw_name()); } else { // reversed loop - TC_ASSERT(for_stmt->vectorize == 1); + TI_ASSERT(for_stmt->vectorize == 1); emit("for (int {}_ = {} - 1; {}_ >= {}; {}_ = {}_ - {}) {{", loop_var->raw_name(), for_stmt->end->raw_name(), loop_var->raw_name(), for_stmt->begin->raw_name(), @@ -279,7 +279,7 @@ class CPUIRCodeGen : public IRVisitor { } void visit(SNodeOpStmt *stmt) { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED /* stmt->ret_type.data_type = DataType::i32; if (stmt->op_type == SNodeOpType::length) { @@ -298,7 +298,7 @@ class CPUIRCodeGen : public IRVisitor { make_list(indices, "")); } if (stmt->op_type == SNodeOpType::append) { - TC_ASSERT(stmt->val->width() == 1); + TI_ASSERT(stmt->val->width() == 1); emit("{}_tmp->append({}({}[{}]));", snode->node_type_name, snode->ch[0]->node_type_name, stmt->val->raw_name(), l); } else if (stmt->op_type == SNodeOpType::clear) { @@ -319,7 +319,7 @@ class CPUIRCodeGen : public IRVisitor { emit("activate_{}(root, {});", snode->node_type_name, make_list(indices, "")); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } emit("}}"); } @@ -332,9 +332,9 @@ class CPUIRCodeGen : public IRVisitor { if (mask) { emit("if ({}[{}]) ", mask->raw_name(), l); } else { - TC_ASSERT(stmt->val->ret_type.data_type == DataType::f32 || + TI_ASSERT(stmt->val->ret_type.data_type == DataType::f32 || stmt->val->ret_type.data_type == DataType::i32); - TC_ASSERT(stmt->op_type == AtomicOpType::add); + TI_ASSERT(stmt->op_type == AtomicOpType::add); emit("atomic_add({}[{}], {}[{}]);", stmt->dest->raw_name(), l, stmt->val->raw_name(), l); } @@ -352,7 +352,7 @@ class CPUIRCodeGen : public IRVisitor { std::vector indices(max_num_indices, "0"); // = "(root, "; for (int i = 0; i < (int)stmt->indices.size(); i++) { if (snode->physical_index_position[i] != -1) { - // TC_ASSERT(snode->physical_index_position[i] != -1); + // TI_ASSERT(snode->physical_index_position[i] != -1); indices[snode->physical_index_position[i]] = stmt->indices[i]->raw_name() + fmt::format("[{}]", l); } @@ -378,7 +378,7 @@ class CPUIRCodeGen : public IRVisitor { all_offsets_zero = false; } if (identical_indices) { - // TC_WARN("Weakened addressing"); + // TI_WARN("Weakened addressing"); weakened = true; std::string cond; @@ -427,7 +427,7 @@ class CPUIRCodeGen : public IRVisitor { if (!current_program->config.force_vectorized_global_store) { for (int i = 0; i < stmt->data->ret_type.width; i++) { if (stmt->parent->mask()) { - TC_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->width() == 1); emit("if ({}[{}])", stmt->parent->mask()->raw_name(), i); } emit("*({} *){}[{}] = {}[{}];", @@ -443,7 +443,7 @@ class CPUIRCodeGen : public IRVisitor { const int width = stmt->width(); if (get_current_program().config.attempt_vectorized_load_cpu && width >= 4 && stmt->ptr->is()) { - TC_ASSERT(stmt->ret_type.data_type == DataType::i32 || + TI_ASSERT(stmt->ret_type.data_type == DataType::i32 || stmt->ret_type.data_type == DataType::f32); auto shuffle = stmt->ptr->as(); @@ -505,8 +505,8 @@ class CPUIRCodeGen : public IRVisitor { } void visit(ExternalPtrStmt *stmt) { - TC_ASSERT(stmt->width() == 1); - TC_ASSERT(stmt->indices.size() == 1); + TI_ASSERT(stmt->width() == 1); + TI_ASSERT(stmt->indices.size() == 1); auto dt = stmt->ret_type.data_type; emit("const {} *{}[1] = {{&{}[{}]}};", data_type_name(dt), stmt->raw_name(), stmt->base_ptrs[0]->raw_name(), stmt->indices[0]->raw_name()); @@ -529,7 +529,7 @@ class CPUIRCodeGen : public IRVisitor { void visit(AssertStmt *stmt) { emit("#if defined(TL_DEBUG)"); - emit(R"(TC_ASSERT_INFO({}, "{}");)", stmt->val->raw_name(), stmt->text); + emit(R"(TI_ASSERT_INFO({}, "{}");)", stmt->val->raw_name(), stmt->text); emit("#endif"); } @@ -605,7 +605,7 @@ class CPUIRCodeGen : public IRVisitor { }; void CPUCodeGen::lower_cpp() { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } void CPUCodeGen::lower_llvm() { @@ -617,27 +617,27 @@ void CPUCodeGen::lower_llvm() { print_ir = prog->config.print_ir; } if (print_ir) { - TC_TRACE("Initial IR:"); + TI_TRACE("Initial IR:"); irpass::re_id(ir); irpass::print(ir); } if (kernel->grad) { irpass::reverse_segments(ir); if (print_ir) { - TC_TRACE("Segment reversed (for autodiff):"); + TI_TRACE("Segment reversed (for autodiff):"); irpass::re_id(ir); irpass::print(ir); } } irpass::lower(ir); if (print_ir) { - TC_TRACE("Lowered:"); + TI_TRACE("Lowered:"); irpass::re_id(ir); irpass::print(ir); } irpass::typecheck(ir); if (print_ir) { - TC_TRACE("Typechecked:"); + TI_TRACE("Typechecked:"); irpass::re_id(ir); irpass::print(ir); } @@ -645,112 +645,112 @@ void CPUCodeGen::lower_llvm() { irpass::demote_dense_struct_fors(ir); irpass::typecheck(ir); if (print_ir) { - TC_TRACE("Dense Struct-for demoted:"); + TI_TRACE("Dense Struct-for demoted:"); irpass::print(ir); } } irpass::slp_vectorize(ir); if (print_ir) { - TC_TRACE("SLPed:"); + TI_TRACE("SLPed:"); irpass::re_id(ir); irpass::print(ir); } irpass::loop_vectorize(ir); if (print_ir) { - TC_TRACE("LoopVeced:"); + TI_TRACE("LoopVeced:"); irpass::re_id(ir); irpass::print(ir); } irpass::vector_split(ir, prog->config.max_vector_width, prog->config.serial_schedule); if (print_ir) { - TC_TRACE("LoopSplitted:"); + TI_TRACE("LoopSplitted:"); irpass::re_id(ir); irpass::print(ir); } irpass::simplify(ir); if (print_ir) { - TC_TRACE("Simplified I:"); + TI_TRACE("Simplified I:"); irpass::re_id(ir); irpass::print(ir); } if (kernel->grad) { // irpass::re_id(ir); - // TC_TRACE("Primal:"); + // TI_TRACE("Primal:"); // irpass::print(ir); irpass::demote_atomics(ir); irpass::simplify(ir); irpass::make_adjoint(ir); irpass::typecheck(ir); if (print_ir) { - TC_TRACE("Adjoint:"); + TI_TRACE("Adjoint:"); irpass::re_id(ir); irpass::print(ir); } } irpass::lower_access(ir, true); if (print_ir) { - TC_TRACE("Access Lowered:"); + TI_TRACE("Access Lowered:"); irpass::re_id(ir); irpass::print(ir); } irpass::die(ir); if (print_ir) { - TC_TRACE("DIEd:"); + TI_TRACE("DIEd:"); irpass::re_id(ir); irpass::print(ir); } irpass::simplify(ir); if (print_ir) { - TC_TRACE("Simplified II:"); + TI_TRACE("Simplified II:"); irpass::re_id(ir); irpass::print(ir); } irpass::die(ir); if (print_ir) { - TC_TRACE("DIEd:"); + TI_TRACE("DIEd:"); irpass::re_id(ir); irpass::print(ir); } irpass::flag_access(ir); if (print_ir) { - TC_TRACE("Access Flagged:"); + TI_TRACE("Access Flagged:"); irpass::re_id(ir); irpass::print(ir); } irpass::constant_fold(ir); if (print_ir) { - TC_TRACE("Constant folded:"); + TI_TRACE("Constant folded:"); irpass::re_id(ir); irpass::print(ir); } irpass::offload(ir); if (print_ir) { - TC_TRACE("Offloaded:"); + TI_TRACE("Offloaded:"); irpass::re_id(ir); irpass::print(ir); } irpass::full_simplify(ir); if (print_ir) { - TC_TRACE("Simplified III:"); + TI_TRACE("Simplified III:"); irpass::re_id(ir); irpass::print(ir); } irpass::demote_atomics(ir); if (print_ir) { - TC_TRACE("Atomics demoted:"); + TI_TRACE("Atomics demoted:"); irpass::re_id(ir); irpass::print(ir); } } void CPUCodeGen::lower() { - TC_PROFILER(__FUNCTION__) + TI_PROFILER(__FUNCTION__) if (prog->config.use_llvm) { lower_llvm(); } else { diff --git a/taichi/backends/kernel.cpp b/taichi/backends/kernel.cpp index 2b8572d32e071..4b27ec84535f6 100644 --- a/taichi/backends/kernel.cpp +++ b/taichi/backends/kernel.cpp @@ -12,12 +12,12 @@ FunctionType KernelCodeGen::compile(taichi::Tlang::Program &prog, this->kernel = &kernel; lower(); if (prog.config.use_llvm) { - TC_PROFILER("codegen llvm") + TI_PROFILER("codegen llvm") return codegen_llvm(); } else { codegen(); generate_binary(""); - // TC_P(Time::get_time() - t); + // TI_P(Time::get_time() - t); return load_function(); } } diff --git a/taichi/backends/kernel.h b/taichi/backends/kernel.h index 09d49e18b93df..8417b60553b3b 100644 --- a/taichi/backends/kernel.h +++ b/taichi/backends/kernel.h @@ -33,7 +33,7 @@ class KernelCodeGen : public CodeGenBase { virtual void codegen() = 0; virtual FunctionType codegen_llvm() { - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; return nullptr; } diff --git a/taichi/backends/llvm_codegen_utils.cpp b/taichi/backends/llvm_codegen_utils.cpp index 3b8a1d9365860..0dfbd707baf86 100644 --- a/taichi/backends/llvm_codegen_utils.cpp +++ b/taichi/backends/llvm_codegen_utils.cpp @@ -13,26 +13,26 @@ void check_func_call_signature(llvm::Value *func, std::vector arglist) { auto func_type = func->getType()->getPointerElementType(); int num_params = func_type->getFunctionNumParams(); - TC_ASSERT(num_params == arglist.size()); + TI_ASSERT(num_params == arglist.size()); for (int i = 0; i < (int)arglist.size(); i++) { auto required = func_type->getFunctionParamType(i); auto provided = arglist[i]->getType(); - // TC_INFO(" required from context {}", (void *)&required->getContext()); - // TC_INFO(" provided from context {}", (void *)&provided->getContext()); + // TI_INFO(" required from context {}", (void *)&required->getContext()); + // TI_INFO(" provided from context {}", (void *)&provided->getContext()); if (required != provided) { - // TC_INFO("Function : {}", std::string(func->getName())); - // TC_INFO(" Type : {}", type_name(func->getType())); + // TI_INFO("Function : {}", std::string(func->getName())); + // TI_INFO(" Type : {}", type_name(func->getType())); if (&required->getContext() != &provided->getContext()) { - TC_INFO(" parameter {} types are from different contexts", i); - TC_INFO(" required from context {}", + TI_INFO(" parameter {} types are from different contexts", i); + TI_INFO(" required from context {}", (void *)&required->getContext()); - TC_INFO(" provided from context {}", + TI_INFO(" provided from context {}", (void *)&provided->getContext()); } - TC_INFO(" parameter {} mismatch: required={}, provided={}", i, + TI_INFO(" parameter {} mismatch: required={}, provided={}", i, type_name(required), type_name(provided)); - TC_ERROR("Bad function signature."); + TI_ERROR("Bad function signature."); } } } diff --git a/taichi/backends/llvm_codegen_utils.h b/taichi/backends/llvm_codegen_utils.h index 853c3caeac5ee..e72cfb8b8ed6e 100644 --- a/taichi/backends/llvm_codegen_utils.h +++ b/taichi/backends/llvm_codegen_utils.h @@ -71,7 +71,7 @@ class ModuleBuilder { llvm::Type *get_runtime_type(const std::string &name) { auto ty = module->getTypeByName("struct." + name); if (!ty) { - TC_ERROR("Runtime type {} not found.", name); + TI_ERROR("Runtime type {} not found.", name); } return ty; } @@ -79,7 +79,7 @@ class ModuleBuilder { llvm::Function *get_runtime_function(const std::string &name) { auto f = module->getFunction(name); if (!f) { - TC_ERROR("Runtime function {} not found.", name); + TI_ERROR("Runtime function {} not found.", name); } f->removeAttribute(AttributeList::FunctionIndex, llvm::Attribute::OptimizeNone); diff --git a/taichi/backends/llvm_jit.h b/taichi/backends/llvm_jit.h index 2d3d138f08086..f166231ffca06 100644 --- a/taichi/backends/llvm_jit.h +++ b/taichi/backends/llvm_jit.h @@ -108,7 +108,7 @@ class TaichiLLVMJIT { return JTMB.takeError(); jtmb = std::make_unique(std::move(*JTMB)); } else { - TC_ASSERT(arch == Arch::cuda); + TI_ASSERT(arch == Arch::cuda); Triple triple("nvptx64", "nvidia", "cuda"); jtmb = std::make_unique(triple); } @@ -186,8 +186,8 @@ class TaichiLLVMJIT { inline void *jit_lookup_name(TaichiLLVMJIT *jit, const std::string &name) { auto ExprSymbol = jit->lookup(name); if (!ExprSymbol) - TC_ERROR("Function \"{}\" not found", name); + TI_ERROR("Function \"{}\" not found", name); return (void *)(llvm::cantFail(ExprSymbol.getAddress())); } -TLANG_NAMESPACE_END \ No newline at end of file +TLANG_NAMESPACE_END diff --git a/taichi/backends/llvm_jit_ptx.cpp b/taichi/backends/llvm_jit_ptx.cpp index 4986c7e9317b8..f0405b39ec1aa 100644 --- a/taichi/backends/llvm_jit_ptx.cpp +++ b/taichi/backends/llvm_jit_ptx.cpp @@ -34,7 +34,7 @@ std::string compile_module_to_ptx(std::unique_ptr &module) { std::string err_str; const llvm::Target *target = TargetRegistry::lookupTarget(triple.str(), err_str); - TC_ERROR_UNLESS(target, err_str); + TI_ERROR_UNLESS(target, err_str); bool fast_math = get_current_program().config.fast_math; @@ -63,7 +63,7 @@ std::string compile_module_to_ptx(std::unique_ptr &module) { triple.str(), cuda_context->get_mcpu(), cuda_mattrs(), options, llvm::Reloc::PIC_, llvm::CodeModel::Small, CodeGenOpt::Aggressive)); - TC_ERROR_UNLESS(target_machine.get(), "Could not allocate target machine!"); + TI_ERROR_UNLESS(target_machine.get(), "Could not allocate target machine!"); module->setDataLayout(target_machine->createDataLayout()); @@ -127,7 +127,7 @@ std::string compile_module_to_ptx(std::unique_ptr &module) { module_pass_manager, ostream, nullptr, TargetMachine::CGFT_AssemblyFile, true); - TC_ERROR_IF(fail, "Failed to set up passes to emit PTX source\n"); + TI_ERROR_IF(fail, "Failed to set up passes to emit PTX source\n"); // Run optimization passes function_pass_manager.doInitialization(); @@ -153,7 +153,7 @@ CUDAContext::CUDAContext() { char name[128]; check_cuda_error(cuDeviceGetName(name, 128, device)); - TC_TRACE("Using CUDA Device [id=0]: {}", name); + TI_TRACE("Using CUDA Device [id=0]: {}", name); int cc_major, cc_minor; check_cuda_error(cuDeviceGetAttribute( @@ -161,7 +161,7 @@ CUDAContext::CUDAContext() { check_cuda_error(cuDeviceGetAttribute( &cc_minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device)); - TC_TRACE("CUDA Device Compute Capability: {}.{}", cc_major, cc_minor); + TI_TRACE("CUDA Device Compute Capability: {}.{}", cc_major, cc_minor); check_cuda_error(cuCtxCreate(&context, 0, device)); check_cuda_error(cudaMalloc(&context_buffer, sizeof(Context))); @@ -173,13 +173,13 @@ CUmodule CUDAContext::compile(const std::string &ptx) { make_current(); // Create module for object CUmodule cudaModule; - TC_TRACE("PTX size: {:.2f}KB", ptx.size() / 1024.0); + TI_TRACE("PTX size: {:.2f}KB", ptx.size() / 1024.0); auto t = Time::get_time(); - TC_TRACE("Loading module..."); + TI_TRACE("Loading module..."); auto _ = std::lock_guard(cuda_context->lock); check_cuda_error( cuModuleLoadDataEx(&cudaModule, ptx.c_str(), 0, nullptr, nullptr)); - TC_TRACE("CUDA module load time : {}ms", (Time::get_time() - t) * 1000); + TI_TRACE("CUDA module load time : {}ms", (Time::get_time() - t) * 1000); cudaModules.push_back(cudaModule); return cudaModule; } @@ -192,7 +192,7 @@ CUfunction CUDAContext::get_function(CUmodule module, auto t = Time::get_time(); check_cuda_error(cuModuleGetFunction(&func, module, func_name.c_str())); t = Time::get_time() - t; - TC_TRACE("Kernel {} compilation time: {}ms", func_name, t * 1000); + TI_TRACE("Kernel {} compilation time: {}ms", func_name, t * 1000); return func; } @@ -230,7 +230,7 @@ void CUDAContext::launch(CUfunction func, check_cuda_error(cudaDeviceSynchronize()); auto err = cudaGetLastError(); if (err) { - TC_ERROR("CUDA Kernel Launch Error: {}", cudaGetErrorString(err)); + TI_ERROR("CUDA Kernel Launch Error: {}", cudaGetErrorString(err)); } } } @@ -246,13 +246,13 @@ CUDAContext::~CUDAContext() { #else std::string compile_module_to_ptx(std::unique_ptr &module) { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } int compile_ptx_and_launch(const std::string &ptx, const std::string &kernel_name, void *) { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } #endif diff --git a/taichi/backends/loopgen.cpp b/taichi/backends/loopgen.cpp index d99f1ae94e6ea..655f5e2803a64 100644 --- a/taichi/backends/loopgen.cpp +++ b/taichi/backends/loopgen.cpp @@ -8,14 +8,14 @@ #if defined(__host__) #undef __host__ #endif -#if defined(CUDA_FOUND) +#if defined(TI_WITH_CUDA) #include #endif TLANG_NAMESPACE_BEGIN LoopGenerator::LoopGenerator(taichi::Tlang::CodeGenBase *gen) : gen(gen) { -#if defined(CUDA_FOUND) +#if defined(TI_WITH_CUDA) int num_SMs; cudaDeviceGetAttribute(&num_SMs, cudaDevAttrMultiProcessorCount, 0); grid_dim = num_SMs * 32; // each SM can have 16-32 resident blocks @@ -41,8 +41,8 @@ void LoopGenerator::emit_listgen_func(SNode *snode, } else { child_block_size /= child_block_division; } - // TC_P(child_block_size); - // TC_P(child_block_division); + // TI_P(child_block_size); + // TI_P(child_block_division); auto parent = snode->parent; @@ -106,7 +106,7 @@ void LoopGenerator::emit_listgen_func(SNode *snode, emit("auto list_element = ({}::child_type *)leaves[leaf_loop].ptr;", parent->node_type_name); auto chid = parent->child_id(snode); - TC_ASSERT(chid != -1); + TI_ASSERT(chid != -1); emit("auto {}_cache = list_element->get{}();", snode->node_type_name, chid); for (int i = 0; i < max_num_indices; i++) { emit("auto {} = leaves[leaf_loop].indices[{}];", @@ -128,7 +128,7 @@ void LoopGenerator::emit_listgen_func(SNode *snode, emit("if (cid >= input_meta.end_loop) break;"); emit("if (!{}_cache->is_active(cid)) continue;", parent->node_type_name); if (deactivate) { - TC_ASSERT(child_block_division == 1); + TI_ASSERT(child_block_division == 1); emit("{}_cache->deactivate(cid);", parent->node_type_name); } } else { diff --git a/taichi/backends/loopgen.h b/taichi/backends/loopgen.h index cd93c1cd31be0..1f200b1103143 100644 --- a/taichi/backends/loopgen.h +++ b/taichi/backends/loopgen.h @@ -34,7 +34,7 @@ class LoopGenerator { } void generate_loop_header(SNode *snode, StructForStmt *stmt) { - TC_ASSERT(snode->type != SNodeType::place) + TI_ASSERT(snode->type != SNodeType::place) if (snode->parent != nullptr) { generate_loop_header(snode->parent, stmt); } @@ -72,7 +72,7 @@ class LoopGenerator { snode->extractors[i].num_bits, snode->extractors[i].start); } else { - TC_ASSERT(snode->num_active_indices <= 3); + TI_ASSERT(snode->num_active_indices <= 3); uint32 mask = mask_base[snode->num_active_indices - 1] << (snode->num_active_indices - morton_id - 1); morton_id++; @@ -111,7 +111,7 @@ class LoopGenerator { snode->node_type_name); } if (snode->_multi_threaded) { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } if (snode->type == SNodeType::hash) { diff --git a/taichi/backends/struct.cpp b/taichi/backends/struct.cpp index 59fdc7257f710..e89d7e70b521b 100644 --- a/taichi/backends/struct.cpp +++ b/taichi/backends/struct.cpp @@ -9,14 +9,14 @@ TLANG_NAMESPACE_BEGIN StructCompiler::StructCompiler(Program *prog) : CodeGenBase(), loopgen(this), prog(prog) { creator = [] { - TC_ERROR("Not Specified"); + TI_ERROR("Not Specified"); return nullptr; }; profiler_clear = [] { - TC_WARN("Profiler not yet implemented in this backend."); + TI_WARN("Profiler not yet implemented in this backend."); }; profiler_print = [] { - TC_WARN("Profiler not yet implemented in this backend."); + TI_WARN("Profiler not yet implemented in this backend."); }; if (get_current_program().config.arch == Arch::x86_64) @@ -42,7 +42,7 @@ void StructCompiler::collect_snodes(SNode &snode) { } void StructCompiler::infer_snode_properties(SNode &snode) { - // TC_P(snode.type_name()); + // TI_P(snode.type_name()); for (int ch_id = 0; ch_id < (int)snode.ch.size(); ch_id++) { auto &ch = snode.ch[ch_id]; ch->parent = &snode; @@ -71,7 +71,7 @@ void StructCompiler::infer_snode_properties(SNode &snode) { if (ch_id == 0) { snode.total_bit_start = total_bits_start_inferred; } else if (snode.parent != nullptr) { // root is ok - // TC_ASSERT(snode.total_bit_start == total_bits_start_inferred); + // TI_ASSERT(snode.total_bit_start == total_bits_start_inferred); } // infer extractors int acc_offsets = 0; @@ -82,9 +82,9 @@ void StructCompiler::infer_snode_properties(SNode &snode) { snode.extractors[i].acc_offset = acc_offsets; } else if (snode.parent != nullptr) { // root is OK /* - TC_ASSERT_INFO(snode.extractors[i].start == inferred, + TI_ASSERT_INFO(snode.extractors[i].start == inferred, "Inconsistent bit configuration"); - TC_ASSERT_INFO(snode.extractors[i].dest_offset == + TI_ASSERT_INFO(snode.extractors[i].dest_offset == snode.total_bit_start + acc_offsets, "Inconsistent bit configuration"); */ @@ -98,14 +98,14 @@ void StructCompiler::infer_snode_properties(SNode &snode) { active_extractor_counder += 1; SNode *p = snode.parent; while (p) { - TC_ASSERT_INFO( + TI_ASSERT_INFO( p->extractors[i].num_bits == 0, "Dynamic SNode must have a standalone dimensionality."); p = p->parent; } } } - TC_ASSERT_INFO(active_extractor_counder == 1, + TI_ASSERT_INFO(active_extractor_counder == 1, "Dynamic SNode can have only one index extractor."); } } @@ -126,7 +126,7 @@ void StructCompiler::infer_snode_properties(SNode &snode) { if (snode.ch.empty()) { if (snode.type != SNodeType::place && snode.type != SNodeType::root) { - TC_ERROR("{} node must have at least one child.", + TI_ERROR("{} node must have at least one child.", snode_type_name(snode.type)); } } @@ -136,7 +136,7 @@ void StructCompiler::generate_types(SNode &snode) { auto type = snode.type; if (snode.type != SNodeType::place && snode.ch.empty()) { - TC_ERROR("Non-place node should have at least one child."); + TI_ERROR("Non-place node should have at least one child."); } // create children type that supports forking... @@ -145,12 +145,12 @@ void StructCompiler::generate_types(SNode &snode) { emit("{} member{};", snode.ch[i]->node_type_name, i); } if (snode.ch.size() == 1 && snode.ch[0]->type == SNodeType::place) { - emit("TC_DEVICE {}_ch({} v) {{*get0()=v;}}", snode.node_type_name, + emit("TI_DEVICE {}_ch({} v) {{*get0()=v;}}", snode.node_type_name, snode.ch[0]->node_type_name); - emit("TC_DEVICE {}_ch() = default;", snode.node_type_name); + emit("TI_DEVICE {}_ch() = default;", snode.node_type_name); } for (int i = 0; i < (int)snode.ch.size(); i++) { - emit("TC_DEVICE {} *get{}() {{return &member{};}} ", + emit("TI_DEVICE {} *get{}() {{return &member{};}} ", snode.ch[i]->node_type_name, i, i); } emit("}};"); @@ -172,15 +172,15 @@ void StructCompiler::generate_types(SNode &snode) { emit("using {} = hash<{}_ch>;", snode.node_type_name, snode.node_type_name); } else if (type == SNodeType::place) { emit( - "struct {} {{ using val_type = {}; val_type val; TC_DEVICE operator " + "struct {} {{ using val_type = {}; val_type val; TI_DEVICE operator " "{}() {{return val;}} " - "TC_DEVICE {}(){{}} TC_DEVICE {}({} val) " + "TI_DEVICE {}(){{}} TI_DEVICE {}({} val) " ": val(val){{ }} }};", snode.node_type_name, snode.data_type_name(), snode.data_type_name(), snode.node_type_name, snode.node_type_name, snode.data_type_name()); } else { - TC_P(snode.type_name()); - TC_NOT_IMPLEMENTED; + TI_P(snode.type_name()); + TI_NOT_IMPLEMENTED; } if (snode.has_null()) { @@ -203,14 +203,14 @@ void StructCompiler::generate_leaf_accessors(SNode &snode) { if (!is_leaf) { // Chain accessors for non-leaf nodes - TC_ASSERT(snode.ch.size() > 0); + TI_ASSERT(snode.ch.size() > 0); for (int i = 0; i < (int)snode.ch.size(); i++) { auto ch = snode.ch[i]; emit( "TLANG_ACCESSOR {} *access_{}({} *parent, int i " ") {{", ch->node_type_name, ch->node_type_name, snode.node_type_name); - // emit("#if defined(TC_STRUCT)"); + // emit("#if defined(TI_STRUCT)"); // emit("parent->activate(i, index);"); // emit("#endif"); emit("auto lookup = parent->look_up(i); "); @@ -225,7 +225,7 @@ void StructCompiler::generate_leaf_accessors(SNode &snode) { } // SNode::place & indirect // emit end2end accessors for leaf (place) nodes, using chain accessors - TC_ASSERT(max_num_indices == 4); + TI_ASSERT(max_num_indices == 4); constexpr int mode_weak_access = 0; constexpr int mode_strong_access = 1; constexpr int mode_activate = 2; @@ -245,7 +245,7 @@ void StructCompiler::generate_leaf_accessors(SNode &snode) { auto ret_type = mode == mode_query ? "bool" : fmt::format("{} *", snode.node_type_name); emit( - "TLANG_ACCESSOR TC_EXPORT {} {}_{}(void *root, int i0=0, int i1=0, " + "TLANG_ACCESSOR TI_EXPORT {} {}_{}(void *root, int i0=0, int i1=0, " "int " "i2=0, " "int i3=0) {{", @@ -268,7 +268,7 @@ void StructCompiler::generate_leaf_accessors(SNode &snode) { emit("tmp = (tmp << {}) + ((i{} >> {}) & ((1 << {}) - 1));", e.num_bits, j, e.start, e.num_bits); } else { - TC_WARN("Emitting shortcut indexing"); + TI_WARN("Emitting shortcut indexing"); emit("tmp = i{};", j); } } @@ -278,7 +278,7 @@ void StructCompiler::generate_leaf_accessors(SNode &snode) { if (force_activate) emit("#if 1"); else - emit("#if defined(TC_STRUCT)"); + emit("#if defined(TI_STRUCT)"); } if (stack[i]->type != SNodeType::place) { if (mode == mode_query) { @@ -372,13 +372,13 @@ void StructCompiler::run(SNode &root, bool host) { for (int i = 0; i < (int)snodes.size(); i++) { if (snodes[i]->type != SNodeType::place) emit( - "TC_EXPORT AllocatorStat stat_{}() {{return " + "TI_EXPORT AllocatorStat stat_{}() {{return " "Managers::get_allocator<{}>()->get_stat();}} ", snodes[i]->node_type_name, snodes[i]->node_type_name); if (snodes[i]->type == SNodeType::pointer || snodes[i]->type == SNodeType::hash) { emit( - "TC_EXPORT void clear_{}(int flags) {{" + "TI_EXPORT void clear_{}(int flags) {{" "Managers::get_allocator<{}>()->clear(flags);}} ", snodes[i]->node_type_name, snodes[i]->node_type_name); } @@ -386,12 +386,12 @@ void StructCompiler::run(SNode &root, bool host) { root_type = root.node_type_name; generate_leaf_accessors(root); - emit("#if defined(TC_STRUCT)"); - emit("TC_EXPORT void *create_data_structure() {{"); + emit("#if defined(TI_STRUCT)"); + emit("TI_EXPORT void *create_data_structure() {{"); emit("Managers::initialize();"); - TC_ASSERT((int)snodes.size() <= max_num_snodes); + TI_ASSERT((int)snodes.size() <= max_num_snodes); for (int i = 0; i < (int)snodes.size(); i++) { // if (snodes[i]->type == SNodeType::pointer || // snodes[i]->type == SNodeType::hashed) { @@ -424,10 +424,10 @@ void StructCompiler::run(SNode &root, bool host) { // emit("CPUProfiler profiler;"); emit("#endif"); - emit("TC_EXPORT void release_data_structure(void *ds) {{delete ({} *)ds;}}", + emit("TI_EXPORT void release_data_structure(void *ds) {{delete ({} *)ds;}}", root_type); - emit("TC_EXPORT void profiler_print()"); + emit("TI_EXPORT void profiler_print()"); emit("{{"); emit("#if defined(TLANG_GPU)"); emit("CUDAProfiler::get_instance().print();"); @@ -436,7 +436,7 @@ void StructCompiler::run(SNode &root, bool host) { emit("#endif"); emit("}}"); - emit("TC_EXPORT void profiler_clear()"); + emit("TI_EXPORT void profiler_clear()"); emit("{{"); emit("#if defined(TLANG_GPU)"); emit("CUDAProfiler::get_instance().clear();"); @@ -452,7 +452,7 @@ void StructCompiler::run(SNode &root, bool host) { emit("}} }}"); write_source(); - generate_binary("-DTC_STRUCT"); + generate_binary("-DTI_STRUCT"); load_dll(); creator = load_function("create_data_structure"); profiler_print = load_function("profiler_print"); diff --git a/taichi/backends/struct_llvm.cpp b/taichi/backends/struct_llvm.cpp index 366afaaa61daf..eea70fba98438 100644 --- a/taichi/backends/struct_llvm.cpp +++ b/taichi/backends/struct_llvm.cpp @@ -11,7 +11,7 @@ TLANG_NAMESPACE_BEGIN void assert_failed_host(const char *msg) { - TC_ERROR("Assertion failure: {}", msg); + TI_ERROR("Assertion failure: {}", msg); } StructCompilerLLVM::StructCompilerLLVM(Program *prog, Arch arch) @@ -19,7 +19,7 @@ StructCompilerLLVM::StructCompilerLLVM(Program *prog, Arch arch) ModuleBuilder(prog->get_llvm_context(arch)->get_init_module()), arch(arch) { creator = [] { - TC_WARN("Data structure creation not implemented"); + TI_WARN("Data structure creation not implemented"); return nullptr; }; tlctx = prog->get_llvm_context(arch); @@ -54,7 +54,7 @@ void StructCompilerLLVM::generate_types(SNode &snode) { llvm::Type *body_type = nullptr, *aux_type = nullptr; if (type == SNodeType::dense) { - TC_ASSERT(snode._morton == false); + TI_ASSERT(snode._morton == false); body_type = llvm::ArrayType::get(ch_type, snode.max_num_elements()); if (snode._bitmasked) { aux_type = llvm::ArrayType::get(Type::getInt32Ty(*llvm_ctx), @@ -72,12 +72,14 @@ void StructCompilerLLVM::generate_types(SNode &snode) { } else if (snode.dt == DataType::f64){ body_type = llvm::Type::getDoubleTy(*ctx); } else { - TC_NOT_IMPLEMENTED + TI_NOT_IMPLEMENTED } } else if (type == SNodeType::pointer) { // mutex - aux_type = llvm::PointerType::getInt64Ty(*ctx); - body_type = llvm::PointerType::getInt8PtrTy(*ctx); + aux_type = llvm::ArrayType::get(llvm::PointerType::getInt64Ty(*ctx), + snode.max_num_elements()); + body_type = llvm::ArrayType::get(llvm::PointerType::getInt8PtrTy(*ctx), + snode.max_num_elements()); } else if (type == SNodeType::dynamic) { // mutex and n (number of elements) aux_type = @@ -85,8 +87,8 @@ void StructCompilerLLVM::generate_types(SNode &snode) { llvm::PointerType::getInt32Ty(*ctx)}); body_type = llvm::PointerType::getInt8PtrTy(*ctx); } else { - TC_P(snode.type_name()); - TC_NOT_IMPLEMENTED; + TI_P(snode.type_name()); + TI_NOT_IMPLEMENTED; } if (aux_type != nullptr) { llvm_type = llvm::StructType::create(*ctx, {aux_type, body_type}, ""); @@ -96,7 +98,7 @@ void StructCompilerLLVM::generate_types(SNode &snode) { snode.has_aux_structure = false; } - TC_ASSERT(llvm_type != nullptr); + TI_ASSERT(llvm_type != nullptr); snode_attr[snode].llvm_type = llvm_type; snode_attr[snode].llvm_aux_type = aux_type; snode_attr[snode].llvm_body_type = body_type; @@ -230,11 +232,11 @@ void StructCompilerLLVM::run(SNode &root, bool host) { generate_leaf_accessors(root); if (prog->config.print_struct_llvm_ir) { - TC_INFO("Struct Module IR"); + TI_INFO("Struct Module IR"); module->print(errs(), nullptr); } - TC_ASSERT((int)snodes.size() <= max_num_snodes); + TI_ASSERT((int)snodes.size() <= max_num_snodes); auto root_size = tlctx->jit->getDataLayout().getTypeAllocSize(snode_attr[root].llvm_type); @@ -284,7 +286,7 @@ void StructCompilerLLVM::run(SNode &root, bool host) { auto root_id = root.id; auto prog = this->prog; creator = [=]() { - TC_TRACE("Allocating data structure of size {} B", root_size); + TI_TRACE("Allocating data structure of size {} B", root_size); auto root = initialize_runtime( &prog->llvm_runtime, prog, (int)snodes.size(), root_size, (void *)&taichi_allocate_aligned, logger.get_level() <= 1); @@ -309,11 +311,11 @@ void StructCompilerLLVM::run(SNode &root, bool host) { tlctx->get_type_size(snode_attr[snodes[i]].llvm_element_type) * snodes[i]->chunk_size; } - TC_TRACE("Initializing allocator for snode {} (node size {})", + TI_TRACE("Initializing allocator for snode {} (node size {})", snodes[i]->id, node_size); auto rt = prog->llvm_runtime; initialize_allocator(rt, i, node_size); - TC_TRACE("Allocating ambient element for snode {} (node size {})", + TI_TRACE("Allocating ambient element for snode {} (node size {})", snodes[i]->id, node_size); allocate_ambient(rt, i); } diff --git a/taichi/backends/struct_metal.cpp b/taichi/backends/struct_metal.cpp index 92dd8e035d7b0..4fc12240bef9b 100644 --- a/taichi/backends/struct_metal.cpp +++ b/taichi/backends/struct_metal.cpp @@ -4,7 +4,7 @@ TLANG_NAMESPACE_BEGIN namespace metal { MetalStructCompiler::CompiledResult MetalStructCompiler::run(SNode &node) { - TC_ASSERT(node.type == SNodeType::root); + TI_ASSERT(node.type == SNodeType::root); collect_snodes(node); // The host side has run this! // infer_snode_properties(node); @@ -85,9 +85,9 @@ void MetalStructCompiler::generate_types(const SNode &snode) { emit(" device byte* addr_;"); emit("}};"); } else { - TC_ERROR("SNodeType={} not supported on Metal", + TI_ERROR("SNodeType={} not supported on Metal", snode_type_name(snode.type)); - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; } emit(""); } diff --git a/taichi/common.h b/taichi/common.h index 5c3816ca49d1c..f2bc3be219cdc 100644 --- a/taichi/common.h +++ b/taichi/common.h @@ -16,12 +16,12 @@ #include #include -#if !defined(TC_INCLUDED) +#if !defined(TI_INCLUDED) #ifdef _WIN64 -#define TC_FORCE_INLINE __forceinline +#define TI_FORCE_INLINE __forceinline #else -#define TC_FORCE_INLINE inline __attribute__((always_inline)) +#define TI_FORCE_INLINE inline __attribute__((always_inline)) #endif #include #include @@ -44,7 +44,7 @@ using int32 = std::int32_t; using int64 = std::int64_t; namespace taichi { -TC_FORCE_INLINE uint32 rand_int() noexcept { +TI_FORCE_INLINE uint32 rand_int() noexcept { static unsigned int x = 123456789, y = 362436069, z = 521288629, w = 88675123; unsigned int t = x ^ (x << 11); x = y; @@ -53,30 +53,30 @@ TC_FORCE_INLINE uint32 rand_int() noexcept { return (w = (w ^ (w >> 19)) ^ (t ^ (t >> 8))); } -TC_FORCE_INLINE uint64 rand_int64() noexcept { +TI_FORCE_INLINE uint64 rand_int64() noexcept { return ((uint64)rand_int() << 32) + rand_int(); } template -TC_FORCE_INLINE T rand() noexcept; +TI_FORCE_INLINE T rand() noexcept; template <> -TC_FORCE_INLINE float rand() noexcept { +TI_FORCE_INLINE float rand() noexcept { return rand_int() * (1.0f / 4294967296.0f); } template <> -TC_FORCE_INLINE double rand() noexcept { +TI_FORCE_INLINE double rand() noexcept { return rand_int() * (1.0 / 4294967296.0); } template <> -TC_FORCE_INLINE int rand() noexcept { +TI_FORCE_INLINE int rand() noexcept { return rand_int(); } template -TC_FORCE_INLINE T rand() noexcept; +TI_FORCE_INLINE T rand() noexcept; } // namespace taichi #endif diff --git a/taichi/common/asset_manager.cpp b/taichi/common/asset_manager.cpp index 04f9cf9f3d35f..c322d66fef59d 100644 --- a/taichi/common/asset_manager.cpp +++ b/taichi/common/asset_manager.cpp @@ -1,11 +1,11 @@ #include "util.h" #include "asset_manager.h" -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN AssetManager &AssetManager::get_instance() { static AssetManager manager; return manager; } -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/taichi/common/asset_manager.h b/taichi/common/asset_manager.h index a089289770299..b68234da72d0e 100644 --- a/taichi/common/asset_manager.h +++ b/taichi/common/asset_manager.h @@ -12,7 +12,7 @@ #include #include -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN class AssetManager { private: @@ -27,10 +27,10 @@ class AssetManager { // Note: this is not thread safe! template std::shared_ptr get_asset_(int id) { - TC_ASSERT_INFO(id_to_asset.find(id) != id_to_asset.end(), + TI_ASSERT_INFO(id_to_asset.find(id) != id_to_asset.end(), "Asset not found"); auto ptr = id_to_asset[id]; - TC_ASSERT_INFO(!ptr.expired(), "Asset has been expired"); + TI_ASSERT_INFO(!ptr.expired(), "Asset has been expired"); return std::static_pointer_cast(ptr.lock()); } @@ -38,7 +38,7 @@ class AssetManager { int insert_asset_(const std::shared_ptr &ptr) { if (asset_to_id.find(ptr.get()) != asset_to_id.end()) { int existing_id = asset_to_id.find(ptr.get())->second; - TC_ASSERT_INFO(id_to_asset[existing_id].expired(), + TI_ASSERT_INFO(id_to_asset[existing_id].expired(), "Asset already exists"); asset_to_id.erase(ptr.get()); id_to_asset.erase(existing_id); @@ -62,4 +62,4 @@ class AssetManager { static AssetManager &get_instance(); }; -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/taichi/common/bit.h b/taichi/common/bit.h index fe2d4bfb0466c..de84f45d097ca 100644 --- a/taichi/common/bit.h +++ b/taichi/common/bit.h @@ -5,23 +5,23 @@ #pragma once #include "util.h" -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN namespace bit { -TC_FORCE_INLINE constexpr bool is_power_of_two(int32 x) { +TI_FORCE_INLINE constexpr bool is_power_of_two(int32 x) { return x != 0 && (x & (x - 1)) == 0; } -TC_FORCE_INLINE constexpr bool is_power_of_two(uint32 x) { +TI_FORCE_INLINE constexpr bool is_power_of_two(uint32 x) { return x != 0 && (x & (x - 1)) == 0; } -TC_FORCE_INLINE constexpr bool is_power_of_two(int64 x) { +TI_FORCE_INLINE constexpr bool is_power_of_two(int64 x) { return x != 0 && (x & (x - 1)) == 0; } -TC_FORCE_INLINE constexpr bool is_power_of_two(uint64 x) { +TI_FORCE_INLINE constexpr bool is_power_of_two(uint64 x) { return x != 0 && (x & (x - 1)) == 0; } @@ -47,25 +47,25 @@ struct Bits { } template - TC_FORCE_INLINE T get() const { + TI_FORCE_INLINE T get() const { return (data >> start) & (((T)1 << bits) - 1); } template - TC_FORCE_INLINE void set(T val) { + TI_FORCE_INLINE void set(T val) { data = (data & ~mask()) | ((val << start) & mask()); } - TC_FORCE_INLINE T operator()(T) const { + TI_FORCE_INLINE T operator()(T) const { return data; } - TC_FORCE_INLINE T get() const { + TI_FORCE_INLINE T get() const { return data; } - TC_FORCE_INLINE void set(const T &data) { + TI_FORCE_INLINE void set(const T &data) { this->data = data; } }; @@ -78,7 +78,7 @@ constexpr int bit_length() { return std::is_same() ? 1 : sizeof(T) * 8; } -#define TC_BIT_FIELD(T, name, start) \ +#define TI_BIT_FIELD(T, name, start) \ T get_##name() const { \ return (T)Base::get()>(); \ } \ @@ -87,7 +87,7 @@ constexpr int bit_length() { } template -TC_FORCE_INLINE constexpr T product(const std::array arr) { +TI_FORCE_INLINE constexpr T product(const std::array arr) { T ret(1); for (int i = 0; i < N; i++) { ret *= arr[i]; @@ -103,11 +103,11 @@ constexpr std::size_t least_pot_bound(std::size_t v) { return ret; } -TC_FORCE_INLINE constexpr uint32 pot_mask(int x) { +TI_FORCE_INLINE constexpr uint32 pot_mask(int x) { return (1u << x) - 1; } -TC_FORCE_INLINE constexpr uint32 log2int(uint64 value) { +TI_FORCE_INLINE constexpr uint32 log2int(uint64 value) { int ret = 0; value >>= 1; while (value) { @@ -118,18 +118,18 @@ TC_FORCE_INLINE constexpr uint32 log2int(uint64 value) { } template -constexpr TC_FORCE_INLINE copy_refcv_t &&reinterpret_bits(T &&t) { - TC_STATIC_ASSERT(sizeof(G) == sizeof(T)); +constexpr TI_FORCE_INLINE copy_refcv_t &&reinterpret_bits(T &&t) { + TI_STATIC_ASSERT(sizeof(G) == sizeof(T)); return std::forward>(*reinterpret_cast(&t)); }; -TC_FORCE_INLINE constexpr float64 compress(float32 h, float32 l) { +TI_FORCE_INLINE constexpr float64 compress(float32 h, float32 l) { uint64 data = ((uint64)reinterpret_bits(h) << 32) + reinterpret_bits(l); return reinterpret_bits(data); } -TC_FORCE_INLINE constexpr std::tuple extract(float64 x) { +TI_FORCE_INLINE constexpr std::tuple extract(float64 x) { auto data = reinterpret_bits(x); return std::make_tuple(reinterpret_bits((uint32)(data >> 32)), reinterpret_bits((uint32)(data & (-1)))); @@ -137,4 +137,4 @@ TC_FORCE_INLINE constexpr std::tuple extract(float64 x) { } // namespace bit -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/taichi/common/dict.h b/taichi/common/dict.h index 046e2b23bfb32..a5cde797563f3 100644 --- a/taichi/common/dict.h +++ b/taichi/common/dict.h @@ -18,11 +18,11 @@ #include "util.h" #include "asset_manager.h" -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN // Declare and then load // Load to `this` -#define TC_LOAD_CONFIG(name, default_val) \ +#define TI_LOAD_CONFIG(name, default_val) \ this->name = config.get(#name, default_val) class Dict { @@ -30,7 +30,7 @@ class Dict { std::map data; public: - TC_IO_DEF(data); + TI_IO_DEF(data); Dict() = default; @@ -74,7 +74,7 @@ class Dict { void check_string_integral(const std::string &str) const { if (!is_string_integral(str)) { - TC_ERROR( + TI_ERROR( "Getting integral value out of non-integral string '{}' is not " "allowed.", str); @@ -281,7 +281,7 @@ class Dict { std::string get_string(std::string key) const { if (data.find(key) == data.end()) { - TC_ERROR("No key named '{}' found.", key); + TI_ERROR("No key named '{}' found.", key); } return data.find(key)->second; } @@ -360,4 +360,4 @@ inline bool Dict::get(std::string key) const { using Config = Dict; -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/taichi/common/interface.h b/taichi/common/interface.h index 8f9415574cbb8..a9cc21f17be61 100644 --- a/taichi/common/interface.h +++ b/taichi/common/interface.h @@ -14,43 +14,43 @@ #include #include -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN template -TC_EXPORT std::shared_ptr create_instance(const std::string &alias); +TI_EXPORT std::shared_ptr create_instance(const std::string &alias); template -TC_EXPORT std::shared_ptr create_instance(const std::string &alias, +TI_EXPORT std::shared_ptr create_instance(const std::string &alias, const Config &config); template -TC_EXPORT std::unique_ptr create_instance_unique(const std::string &alias); +TI_EXPORT std::unique_ptr create_instance_unique(const std::string &alias); template -TC_EXPORT std::unique_ptr create_instance_unique(const std::string &alias, +TI_EXPORT std::unique_ptr create_instance_unique(const std::string &alias, const Config &config); template -TC_EXPORT std::unique_ptr create_instance_unique_ctor( +TI_EXPORT std::unique_ptr create_instance_unique_ctor( const std::string &alias, const Config &config); template -TC_EXPORT T *create_instance_raw(const std::string &alias); +TI_EXPORT T *create_instance_raw(const std::string &alias); template -TC_EXPORT T *create_instance_raw(const std::string &alias, +TI_EXPORT T *create_instance_raw(const std::string &alias, const Config &config); template -TC_EXPORT T *create_instance_placement(const std::string &alias, void *place); +TI_EXPORT T *create_instance_placement(const std::string &alias, void *place); template -TC_EXPORT T *create_instance_placement(const std::string &alias, +TI_EXPORT T *create_instance_placement(const std::string &alias, void *place, const Config &config); template -TC_EXPORT std::vector get_implementation_names(); +TI_EXPORT std::vector get_implementation_names(); class Unit { public: @@ -65,12 +65,12 @@ class Unit { } virtual std::string get_name() const { - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; return ""; } virtual std::string general_action(const Config &config) { - TC_NOT_IMPLEMENTED; + TI_NOT_IMPLEMENTED; return ""; } @@ -78,8 +78,8 @@ class Unit { } }; -#define TC_IMPLEMENTATION_HOLDER_NAME(T) ImplementationHolder_##T -#define TC_IMPLEMENTATION_HOLDER_PTR(T) instance_ImplementationHolder_##T +#define TI_IMPLEMENTATION_HOLDER_NAME(T) ImplementationHolder_##T +#define TI_IMPLEMENTATION_HOLDER_PTR(T) instance_ImplementationHolder_##T class ImplementationHolderBase { public: @@ -114,12 +114,12 @@ class InterfaceHolder { } }; -#define TC_INTERFACE(T) \ +#define TI_INTERFACE(T) \ extern void *get_implementation_holder_instance_##T(); \ - class TC_IMPLEMENTATION_HOLDER_NAME(T) final \ + class TI_IMPLEMENTATION_HOLDER_NAME(T) final \ : public ImplementationHolderBase { \ public: \ - TC_IMPLEMENTATION_HOLDER_NAME(T)(const std::string &name) { \ + TI_IMPLEMENTATION_HOLDER_NAME(T)(const std::string &name) { \ this->name = name; \ } \ using FactoryMethod = std::function()>; \ @@ -225,66 +225,66 @@ class InterfaceHolder { "Implementation [" + name + "::" + alias + "] not found!"); \ return (factory->second)(place); \ } \ - static TC_IMPLEMENTATION_HOLDER_NAME(T) * get_instance() { \ - return static_cast( \ + static TI_IMPLEMENTATION_HOLDER_NAME(T) * get_instance() { \ + return static_cast( \ get_implementation_holder_instance_##T()); \ } \ }; \ - extern TC_IMPLEMENTATION_HOLDER_NAME(T) * TC_IMPLEMENTATION_HOLDER_PTR(T); + extern TI_IMPLEMENTATION_HOLDER_NAME(T) * TI_IMPLEMENTATION_HOLDER_PTR(T); -#define TC_INTERFACE_DEF(class_name, base_alias) \ +#define TI_INTERFACE_DEF(class_name, base_alias) \ template <> \ - TC_EXPORT std::shared_ptr create_instance( \ + TI_EXPORT std::shared_ptr create_instance( \ const std::string &alias) { \ - return TC_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance()->create( \ + return TI_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance()->create( \ alias); \ } \ template <> \ - TC_EXPORT std::shared_ptr create_instance( \ + TI_EXPORT std::shared_ptr create_instance( \ const std::string &alias, const Config &config) { \ auto instance = create_instance(alias); \ instance->initialize(config); \ return instance; \ } \ template <> \ - TC_EXPORT std::unique_ptr create_instance_unique( \ + TI_EXPORT std::unique_ptr create_instance_unique( \ const std::string &alias) { \ - return TC_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ + return TI_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ ->create_unique(alias); \ } \ template <> \ - TC_EXPORT std::unique_ptr create_instance_unique( \ + TI_EXPORT std::unique_ptr create_instance_unique( \ const std::string &alias, const Config &config) { \ auto instance = create_instance_unique(alias); \ instance->initialize(config); \ return instance; \ } \ template <> \ - TC_EXPORT std::unique_ptr create_instance_unique_ctor( \ + TI_EXPORT std::unique_ptr create_instance_unique_ctor( \ const std::string &alias, const Dict &config) { \ - return TC_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ + return TI_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ ->create_unique_ctor(alias, config); \ } \ template <> \ - TC_EXPORT class_name *create_instance_raw(const std::string &alias) { \ - return TC_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ + TI_EXPORT class_name *create_instance_raw(const std::string &alias) { \ + return TI_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ ->create_raw(alias); \ } \ template <> \ - TC_EXPORT class_name *create_instance_placement(const std::string &alias, \ + TI_EXPORT class_name *create_instance_placement(const std::string &alias, \ void *place) { \ - return TC_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ + return TI_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ ->create_placement(alias, place); \ } \ template <> \ - TC_EXPORT class_name *create_instance_placement( \ + TI_EXPORT class_name *create_instance_placement( \ const std::string &alias, void *place, const Config &config) { \ auto instance = create_instance_placement(alias, place); \ instance->initialize(config); \ return instance; \ } \ template <> \ - TC_EXPORT class_name *create_instance_raw(const std::string &alias, \ + TI_EXPORT class_name *create_instance_raw(const std::string &alias, \ const Config &config) { \ auto instance = create_instance_raw(alias); \ instance->initialize(config); \ @@ -292,17 +292,17 @@ class InterfaceHolder { } \ template <> \ std::vector get_implementation_names() { \ - return TC_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ + return TI_IMPLEMENTATION_HOLDER_NAME(class_name)::get_instance() \ ->get_implementation_names(); \ } \ - TC_IMPLEMENTATION_HOLDER_NAME(class_name) * \ - TC_IMPLEMENTATION_HOLDER_PTR(class_name) = nullptr; \ + TI_IMPLEMENTATION_HOLDER_NAME(class_name) * \ + TI_IMPLEMENTATION_HOLDER_PTR(class_name) = nullptr; \ void *get_implementation_holder_instance_##class_name() { \ - if (!TC_IMPLEMENTATION_HOLDER_PTR(class_name)) { \ - TC_IMPLEMENTATION_HOLDER_PTR(class_name) = \ - new TC_IMPLEMENTATION_HOLDER_NAME(class_name)(base_alias); \ + if (!TI_IMPLEMENTATION_HOLDER_PTR(class_name)) { \ + TI_IMPLEMENTATION_HOLDER_PTR(class_name) = \ + new TI_IMPLEMENTATION_HOLDER_NAME(class_name)(base_alias); \ } \ - return TC_IMPLEMENTATION_HOLDER_PTR(class_name); \ + return TI_IMPLEMENTATION_HOLDER_PTR(class_name); \ } \ class InterfaceInjector_##class_name { \ public: \ @@ -329,25 +329,25 @@ class InterfaceHolder { } \ } ImplementationInjector_##base_class_name##class_name##instance(base_alias); -#define TC_IMPLEMENTATION(base_class_name, class_name, alias) \ +#define TI_IMPLEMENTATION(base_class_name, class_name, alias) \ class ImplementationInjector_##base_class_name##class_name { \ public: \ ImplementationInjector_##base_class_name##class_name() { \ - TC_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance() \ + TI_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance() \ ->insert(alias); \ } \ } ImplementationInjector_##base_class_name##class_name##instance; -#define TC_IMPLEMENTATION_NEW(base_class_name, class_name) \ +#define TI_IMPLEMENTATION_NEW(base_class_name, class_name) \ class ImplementationInjector_##base_class_name##class_name { \ public: \ ImplementationInjector_##base_class_name##class_name() { \ - TC_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance() \ + TI_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance() \ ->insert_new(class_name::get_name_static()); \ } \ } ImplementationInjector_##base_class_name##class_name##instance; -#define TC_NAME(alias) \ +#define TI_NAME(alias) \ virtual std::string get_name() const override { \ return get_name_static(); \ } \ @@ -355,4 +355,4 @@ class InterfaceHolder { return alias; \ } -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/taichi/common/loader.h b/taichi/common/loader.h index 769c01c3fcb12..9f17a3721ea65 100644 --- a/taichi/common/loader.h +++ b/taichi/common/loader.h @@ -7,28 +7,28 @@ #include "interface.h" -TC_NAMESPACE_BEGIN +TI_NAMESPACE_BEGIN -#define TC_IMPLEMENTATION_LOADER(base_class_name, class_name, alias) \ +#define TI_IMPLEMENTATION_LOADER(base_class_name, class_name, alias) \ class ImplementationLoader_##base_class_name##class_name { \ public: \ ImplementationLoader_##base_class_name##class_name() { \ - TC_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance() \ + TI_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance() \ ->insert(alias); \ } \ ~ImplementationLoader_##base_class_name##class_name() { \ - TC_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance()->remove( \ + TI_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance()->remove( \ alias); \ } \ } ImplementationLoader_##base_class_name##class_name##instance; -#define TC_IMPLEMENTATION_UPDATER(base_class_name, class_name, alias) \ +#define TI_IMPLEMENTATION_UPDATER(base_class_name, class_name, alias) \ class ImplementationUpdater_##base_class_name##class_name { \ public: \ ImplementationUpdater_##base_class_name##class_name() { \ - TC_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance() \ + TI_IMPLEMENTATION_HOLDER_NAME(base_class_name)::get_instance() \ ->update(alias); \ } \ } ImplementationUpdater_##base_class_name##class_name##instance; -TC_NAMESPACE_END +TI_NAMESPACE_END diff --git a/taichi/common/meta.h b/taichi/common/meta.h index 13932bd61a4ec..07d76fef10a1f 100644 --- a/taichi/common/meta.h +++ b/taichi/common/meta.h @@ -17,7 +17,7 @@ namespace taichi { namespace meta { template