From d1f2b1314813c7e40b5f14bf2ee676d4f9c0b544 Mon Sep 17 00:00:00 2001 From: Taichi Gardener Date: Wed, 23 Jun 2021 10:34:28 +0000 Subject: [PATCH] Auto Format --- examples/autodiff_minimization.py | 3 ++- examples/autodiff_regression.py | 8 ++++--- examples/cg_possion.py | 4 +++- examples/comet.py | 4 +++- examples/euler.py | 4 +++- examples/export_ply.py | 8 ++++--- examples/fullscreen.py | 3 ++- examples/game_of_life.py | 3 ++- examples/gui_image_io.py | 3 ++- examples/mciso.py | 4 +++- examples/mciso_advanced.py | 3 ++- examples/mgpcg.py | 1 + examples/mgpcg_advanced.py | 6 +++-- examples/mnist.py | 6 +++-- examples/mpm3d.py | 3 ++- examples/nbody_oscillator.py | 3 ++- examples/odop_solar.py | 3 ++- examples/particle_renderer.py | 11 ++++++---- examples/pbf2d.py | 6 +++-- examples/physarum.py | 2 ++ examples/quadtree.py | 3 ++- examples/renderer_utils.py | 3 ++- examples/revert_image_color.py | 1 + examples/sdf2d.py | 5 ++++- examples/sdf_renderer.py | 6 +++-- examples/simple_uv.py | 3 ++- examples/taichi_autodiff.py | 3 ++- examples/taichi_bitmasked.py | 3 ++- examples/tree_gravity.py | 4 +++- examples/vortex_rings.py | 6 +++-- examples/waterwave.py | 3 ++- python/taichi/lang/kernel_impl.py | 3 ++- .../metal/shaders/snode_bit_pointer.metal.h | 6 +++-- taichi/backends/wasm/codegen_wasm.cpp | 14 ++++++------ taichi/ir/expression_ops.h | 2 +- taichi/program/async_utils.h | 4 ++-- taichi/program/callable.cpp | 2 +- taichi/runtime/llvm/node_pointer.h | 21 +++++++++--------- taichi/transforms/auto_diff.cpp | 22 +++++++++---------- taichi/transforms/inlining.cpp | 6 ++--- tests/python/bls_test_template.py | 6 +++-- tests/python/fuse_test_template.py | 3 ++- tests/python/test_assert.py | 3 ++- tests/python/test_async.py | 3 ++- tests/python/test_bit_array.py | 3 ++- tests/python/test_bit_array_vectorization.py | 3 ++- tests/python/test_bit_operations.py | 8 ++++--- tests/python/test_bit_struct.py | 3 ++- tests/python/test_bls_assume_in_range.py | 1 + tests/python/test_cast.py | 3 ++- tests/python/test_cli.py | 16 ++++++++------ tests/python/test_constant_func.py | 8 ++++--- tests/python/test_custom_float.py | 4 +++- tests/python/test_custom_float_exponents.py | 3 ++- tests/python/test_custom_float_shared_exp.py | 5 +++-- .../test_custom_float_time_integration.py | 4 +++- tests/python/test_custom_type_atomics.py | 3 ++- tests/python/test_debug.py | 3 ++- tests/python/test_dynamic.py | 3 ++- tests/python/test_element_wise.py | 5 +++-- tests/python/test_external_func.py | 3 ++- tests/python/test_field.py | 3 ++- tests/python/test_fuse_dense.py | 5 +++-- tests/python/test_fuse_dynamic.py | 4 +++- .../python/test_get_external_tensor_shape.py | 3 ++- tests/python/test_gui.py | 5 +++-- tests/python/test_image_io.py | 8 ++++--- tests/python/test_kernel_arg_errors.py | 3 ++- tests/python/test_lang.py | 3 ++- tests/python/test_linalg.py | 8 ++++--- tests/python/test_listgen.py | 3 ++- tests/python/test_matrix.py | 8 ++++--- tests/python/test_matrix_different_type.py | 3 ++- tests/python/test_mpm88.py | 4 +++- tests/python/test_mpm_particle_list.py | 3 ++- tests/python/test_native_functions.py | 3 ++- tests/python/test_ndrange.py | 3 ++- tests/python/test_numpy.py | 3 ++- tests/python/test_numpy_io.py | 3 ++- tests/python/test_print.py | 3 ++- tests/python/test_reduction.py | 3 ++- tests/python/test_runtime.py | 8 +++++-- tests/python/test_sfg.py | 3 ++- tests/python/test_ssa.py | 6 +++-- tests/python/test_static.py | 3 ++- tests/python/test_svd.py | 3 ++- tests/python/test_syntax_errors.py | 3 ++- tests/python/test_tensor_reflection.py | 3 ++- tests/python/test_test.py | 3 ++- tests/python/test_torch_ad.py | 3 ++- tests/python/test_torch_io.py | 3 ++- tests/python/test_trailing_bits.py | 3 ++- tests/python/test_types.py | 3 ++- tests/python/test_unary_ops.py | 3 ++- 94 files changed, 274 insertions(+), 155 deletions(-) diff --git a/examples/autodiff_minimization.py b/examples/autodiff_minimization.py index 0fd325529c6503..a3917306cc6514 100644 --- a/examples/autodiff_minimization.py +++ b/examples/autodiff_minimization.py @@ -1,6 +1,7 @@ -import taichi as ti import random +import taichi as ti + n = 8 x = ti.field(dtype=ti.f32, shape=n, needs_grad=True) y = ti.field(dtype=ti.f32, shape=n) diff --git a/examples/autodiff_regression.py b/examples/autodiff_regression.py index 1762f2b0a9b3bb..1a27a003e4f1de 100644 --- a/examples/autodiff_regression.py +++ b/examples/autodiff_regression.py @@ -1,9 +1,11 @@ -import taichi as ti -import taichi as tc -import matplotlib.pyplot as plt import random + +import matplotlib.pyplot as plt import numpy as np +import taichi as ti +import taichi as tc + tc.set_gdb_trigger(True) number_coeffs = 4 diff --git a/examples/cg_possion.py b/examples/cg_possion.py index 73d01f818c93be..5e65d6e7d1263e 100644 --- a/examples/cg_possion.py +++ b/examples/cg_possion.py @@ -1,5 +1,7 @@ -import taichi as ti import math + +import taichi as ti + ti.init(arch=ti.cuda) n = 512 diff --git a/examples/comet.py b/examples/comet.py index d63d4979089925..9fe8442dc2909b 100644 --- a/examples/comet.py +++ b/examples/comet.py @@ -1,5 +1,7 @@ -import taichi as ti import taichi_glsl as tl + +import taichi as ti + ti.init(arch=[ti.cuda, ti.metal]) dim = 3 diff --git a/examples/euler.py b/examples/euler.py index c981ecc93a210b..994ecf74522aea 100644 --- a/examples/euler.py +++ b/examples/euler.py @@ -1,5 +1,7 @@ -import taichi as ti import matplotlib.cm as cm + +import taichi as ti + # A compressible euler equation solver using two methods # 1: 2nd order muscl # 2: thinc BVD, ref: "Limiter-free discontinuity-capturing scheme diff --git a/examples/export_ply.py b/examples/export_ply.py index 272ab63d5fcfb4..a5c045d245985a 100644 --- a/examples/export_ply.py +++ b/examples/export_ply.py @@ -1,7 +1,9 @@ -import taichi as ti -import numpy as np -import random import os +import random + +import numpy as np + +import taichi as ti ti.init(arch=ti.cpu) diff --git a/examples/fullscreen.py b/examples/fullscreen.py index d13f1eab36daa6..412747be1868b6 100644 --- a/examples/fullscreen.py +++ b/examples/fullscreen.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + ti.init(ti.gpu) res = (1920, 1080) diff --git a/examples/game_of_life.py b/examples/game_of_life.py index 6436f53cf9dc3f..e5481b072e4299 100644 --- a/examples/game_of_life.py +++ b/examples/game_of_life.py @@ -1,9 +1,10 @@ # Game of Life written in 100 lines of Taichi # In memory of John Horton Conway (1937 - 2020) -import taichi as ti import numpy as np +import taichi as ti + ti.init() n = 64 diff --git a/examples/gui_image_io.py b/examples/gui_image_io.py index c949a26f1c2d3e..b4893a463e9ec8 100644 --- a/examples/gui_image_io.py +++ b/examples/gui_image_io.py @@ -1,6 +1,7 @@ -import taichi as ti import os +import taichi as ti + pixel = ti.field(ti.u8, shape=(512, 512, 3)) diff --git a/examples/mciso.py b/examples/mciso.py index 8fa4bb4e9c7d41..4bea2d4df26a01 100644 --- a/examples/mciso.py +++ b/examples/mciso.py @@ -1,5 +1,7 @@ -import taichi as ti import numpy as np + +import taichi as ti + ti.init() N = 128 diff --git a/examples/mciso_advanced.py b/examples/mciso_advanced.py index 2eb3580e03fa6d..97a5dd0464fbae 100644 --- a/examples/mciso_advanced.py +++ b/examples/mciso_advanced.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + @ti.data_oriented class MCISO: diff --git a/examples/mgpcg.py b/examples/mgpcg.py index d54a2ac7dc457d..f60ba02e01a3ec 100644 --- a/examples/mgpcg.py +++ b/examples/mgpcg.py @@ -1,5 +1,6 @@ # Solve Poisson's equation on an NxN grid using MGPCG import numpy as np + import taichi as ti real = ti.f32 diff --git a/examples/mgpcg_advanced.py b/examples/mgpcg_advanced.py index 613e967db02409..33f02a98211fc3 100644 --- a/examples/mgpcg_advanced.py +++ b/examples/mgpcg_advanced.py @@ -1,7 +1,9 @@ +import math +import time + import numpy as np + import taichi as ti -import time -import math @ti.data_oriented diff --git a/examples/mnist.py b/examples/mnist.py index 728ba638434e99..19369c1e5b5ab5 100644 --- a/examples/mnist.py +++ b/examples/mnist.py @@ -1,7 +1,9 @@ -import numpy as np +import pickle import random + +import numpy as np + import taichi as ti -import pickle # ti.runtime.print_preprocessed = True # ti.cfg.print_ir = True diff --git a/examples/mpm3d.py b/examples/mpm3d.py index 05deb659a1e353..0e63972e0120bd 100644 --- a/examples/mpm3d.py +++ b/examples/mpm3d.py @@ -1,8 +1,9 @@ export_file = '' # use '/tmp/mpm3d.ply' for exporting result to disk -import taichi as ti import numpy as np +import taichi as ti + ti.init(arch=ti.gpu) #dim, n_grid, steps, dt = 2, 128, 20, 2e-4 diff --git a/examples/nbody_oscillator.py b/examples/nbody_oscillator.py index 693c1dcc8d6f3b..323f5d438eeacf 100644 --- a/examples/nbody_oscillator.py +++ b/examples/nbody_oscillator.py @@ -1,6 +1,7 @@ -import taichi as ti import math +import taichi as ti + ti.init(arch=ti.gpu) N = 1600 diff --git a/examples/odop_solar.py b/examples/odop_solar.py index a9f12f62f0becc..72a72f4b53d364 100644 --- a/examples/odop_solar.py +++ b/examples/odop_solar.py @@ -1,6 +1,7 @@ -import taichi as ti import math +import taichi as ti + ti.init() diff --git a/examples/particle_renderer.py b/examples/particle_renderer.py index b76b48c3062dda..078c6bd3b85c84 100644 --- a/examples/particle_renderer.py +++ b/examples/particle_renderer.py @@ -1,9 +1,12 @@ -import taichi as ti -import numpy as np import math import time -from renderer_utils import out_dir, ray_aabb_intersection, inf, eps, \ - intersect_sphere, sphere_aabb_intersect_motion, inside_taichi + +import numpy as np +from renderer_utils import (eps, inf, inside_taichi, intersect_sphere, out_dir, + ray_aabb_intersection, + sphere_aabb_intersect_motion) + +import taichi as ti ti.init(arch=ti.cuda, device_memory_GB=4) diff --git a/examples/pbf2d.py b/examples/pbf2d.py index bd130476d77c80..3ea258b3a898b3 100644 --- a/examples/pbf2d.py +++ b/examples/pbf2d.py @@ -1,10 +1,12 @@ # Macklin, M. and Müller, M., 2013. Position based fluids. ACM Transactions on Graphics (TOG), 32(4), p.104. # Taichi implementation by Ye Kuang (k-ye) -import taichi as ti -import numpy as np import math +import numpy as np + +import taichi as ti + ti.init(arch=ti.gpu) screen_res = (800, 400) diff --git a/examples/physarum.py b/examples/physarum.py index 3d4b6c827c9942..9eadf30211d0a5 100644 --- a/examples/physarum.py +++ b/examples/physarum.py @@ -3,7 +3,9 @@ See https://sagejenson.com/physarum for the details.''' import numpy as np + import taichi as ti + ti.init(arch=ti.gpu) PARTICLE_N = 1024 diff --git a/examples/quadtree.py b/examples/quadtree.py index b7e98ef79cc1c4..28cfbec9b4ddd4 100644 --- a/examples/quadtree.py +++ b/examples/quadtree.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + ti.init(arch=ti.cpu) RES = 1024 diff --git a/examples/renderer_utils.py b/examples/renderer_utils.py index e2308a1ed67ca6..de3b782ce4212c 100644 --- a/examples/renderer_utils.py +++ b/examples/renderer_utils.py @@ -1,6 +1,7 @@ -import taichi as ti import math +import taichi as ti + eps = 1e-4 inf = 1e10 diff --git a/examples/revert_image_color.py b/examples/revert_image_color.py index 59f6fc5e4170b7..430f6edf267c38 100644 --- a/examples/revert_image_color.py +++ b/examples/revert_image_color.py @@ -1,4 +1,5 @@ import taichi as ti + ti.init() ## Load the image: diff --git a/examples/sdf2d.py b/examples/sdf2d.py index 584c5d39279c0a..6b301fb2fc6d45 100644 --- a/examples/sdf2d.py +++ b/examples/sdf2d.py @@ -1,6 +1,9 @@ -import taichi as ti from math import tau + from renderer_utils import reflect, refract + +import taichi as ti + ti.init(arch=ti.opengl) N = 512 diff --git a/examples/sdf_renderer.py b/examples/sdf_renderer.py index f5a261c02b9633..284e00d4b586c9 100644 --- a/examples/sdf_renderer.py +++ b/examples/sdf_renderer.py @@ -1,8 +1,10 @@ -import taichi as ti -import time import math +import time + import numpy as np +import taichi as ti + ti.init(arch=ti.gpu) res = 1280, 720 color_buffer = ti.Vector.field(3, dtype=ti.f32, shape=res) diff --git a/examples/simple_uv.py b/examples/simple_uv.py index e57997bad6a9fc..288c81c8552529 100644 --- a/examples/simple_uv.py +++ b/examples/simple_uv.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + ti.init() res = (512, 512) diff --git a/examples/taichi_autodiff.py b/examples/taichi_autodiff.py index 482735212c6d8a..579245f32ff4ba 100644 --- a/examples/taichi_autodiff.py +++ b/examples/taichi_autodiff.py @@ -1,6 +1,7 @@ -import taichi as ti import matplotlib.pyplot as plt +import taichi as ti + N = 2048 x, y = ti.field(ti.f32), ti.field(ti.f32) diff --git a/examples/taichi_bitmasked.py b/examples/taichi_bitmasked.py index 9baba8f2918ddf..a7bd086606e77b 100644 --- a/examples/taichi_bitmasked.py +++ b/examples/taichi_bitmasked.py @@ -1,6 +1,7 @@ -import taichi as ti import math +import taichi as ti + ti.init(arch=ti.gpu) n = 256 diff --git a/examples/tree_gravity.py b/examples/tree_gravity.py index 2435c121aedb4e..07e661b70cedcb 100644 --- a/examples/tree_gravity.py +++ b/examples/tree_gravity.py @@ -1,7 +1,9 @@ # N-body gravity simulation in 300 lines of Taichi, tree method, no multipole, O(N log N) # Author: archibate <1931127624@qq.com>, all left reserved -import taichi as ti import taichi_glsl as tl + +import taichi as ti + ti.init() if not hasattr(ti, 'jkl'): ti.jkl = ti.indices(1, 2, 3) diff --git a/examples/vortex_rings.py b/examples/vortex_rings.py index 49a3ca0f0a1555..ac8eaedb761c80 100644 --- a/examples/vortex_rings.py +++ b/examples/vortex_rings.py @@ -1,8 +1,10 @@ # C++ reference and tutorial (Chinese): https://zhuanlan.zhihu.com/p/26882619 -import taichi as ti -import numpy as np import math +import numpy as np + +import taichi as ti + ti.init(arch=ti.gpu) eps = 0.01 diff --git a/examples/waterwave.py b/examples/waterwave.py index 81830c84f94b65..76845ff15cf2ae 100644 --- a/examples/waterwave.py +++ b/examples/waterwave.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + ti.init(arch=ti.gpu) light_color = 1 diff --git a/python/taichi/lang/kernel_impl.py b/python/taichi/lang/kernel_impl.py index 73f7cb7f768003..eb2a07c389d085 100644 --- a/python/taichi/lang/kernel_impl.py +++ b/python/taichi/lang/kernel_impl.py @@ -5,7 +5,6 @@ import re import numpy as np -import taichi as ti from taichi.core import primitive_types from taichi.core.util import ti_core as _ti_core from taichi.lang import impl, util @@ -16,6 +15,8 @@ from taichi.lang.transformer import ASTTransformer from taichi.misc.util import obsolete +import taichi as ti + def _remove_indent(lines): lines = lines.split('\n') diff --git a/taichi/backends/metal/shaders/snode_bit_pointer.metal.h b/taichi/backends/metal/shaders/snode_bit_pointer.metal.h index 5e29603b39b9ac..63310fe1ff39f5 100644 --- a/taichi/backends/metal/shaders/snode_bit_pointer.metal.h +++ b/taichi/backends/metal/shaders/snode_bit_pointer.metal.h @@ -37,7 +37,8 @@ STR( }; // |f| should already be scaled. |C| is the compute type. - template C mtl_float_to_custom_int(float f) { + template + C mtl_float_to_custom_int(float f) { // Branch free implementation of `f + sign(f) * 0.5`. // See rounding_prepare_f* in taichi/runtime/llvm/runtime.cpp const int32_t delta_bits = @@ -133,7 +134,8 @@ STR( return static_cast(step1 >> (N - bits)); } - template C mtl_get_full_bits(SNodeBitPointer bp) { + template + C mtl_get_full_bits(SNodeBitPointer bp) { return static_cast(*(bp.base)); }) METAL_END_SRC_DEF diff --git a/taichi/backends/wasm/codegen_wasm.cpp b/taichi/backends/wasm/codegen_wasm.cpp index e43b0186a08f24..c72ec01adbff07 100644 --- a/taichi/backends/wasm/codegen_wasm.cpp +++ b/taichi/backends/wasm/codegen_wasm.cpp @@ -100,11 +100,11 @@ class CodeGenLLVMWASM : public CodeGenLLVM { /** * Extracts the original function name decorated by @ti.kernel - * + * * @param kernel_name The format is defined in * https://github.com/taichi-dev/taichi/blob/734da3f8f4439ce7f6a5337df7c54fb6dc34def8/python/taichi/lang/kernel_impl.py#L360-L362 */ - std::string extract_original_kernel_name(const std::string& kernel_name) { + std::string extract_original_kernel_name(const std::string &kernel_name) { if (kernel->is_evaluator) return kernel_name; int pos = kernel_name.length() - 1; @@ -200,9 +200,9 @@ class CodeGenLLVMWASM : public CodeGenLLVM { return task_kernel_name; } - // Context's address is pass by kernel_args[0] which is supposed to be 0 in default. - // Runtime's address will be set to kernel_args[0] after set_root() call. - // The objects of Context and Runtime are overlapped with each other. + // Context's address is pass by kernel_args[0] which is supposed to be 0 in + // default. Runtime's address will be set to kernel_args[0] after set_root() + // call. The objects of Context and Runtime are overlapped with each other. // // Context Runtime Root Buffer // +-----------+ +-------------+ +-------------+ @@ -240,7 +240,7 @@ class CodeGenLLVMWASM : public CodeGenLLVM { llvm::Value *runtime_address_val_ptr = builder->CreatePointerCast( runtime_address_ptr, llvm::Type::getInt32PtrTy(*llvm_context)); llvm::Value *runtime_address_val = builder->CreatePtrToInt( - kernel_args[0], llvm::Type::getInt32Ty(*llvm_context)); + kernel_args[0], llvm::Type::getInt32Ty(*llvm_context)); builder->CreateStore(runtime_address_val, runtime_address_val_ptr); llvm::Value *runtime_ptr = @@ -251,7 +251,7 @@ class CodeGenLLVMWASM : public CodeGenLLVM { llvm::Value *root_base_ptr = builder->CreatePointerCast( kernel_args[0], llvm::Type::getInt32PtrTy(*llvm_context)); - llvm::Value* root_base_val = builder->CreateLoad(root_base_ptr); + llvm::Value *root_base_val = builder->CreateLoad(root_base_ptr); llvm::Value *root_val = builder->CreateAdd(root_base_val, kernel_args[1]); llvm::Value *root_ptr = builder->CreateIntToPtr( root_val, llvm::Type::getInt8PtrTy(*llvm_context)); diff --git a/taichi/ir/expression_ops.h b/taichi/ir/expression_ops.h index 3c349e65dbcdbd..7ff0bd1813d4e1 100644 --- a/taichi/ir/expression_ops.h +++ b/taichi/ir/expression_ops.h @@ -82,7 +82,7 @@ DEFINE_EXPRESSION_OP_BINARY(&&, bit_and) DEFINE_EXPRESSION_OP_BINARY(||, bit_or) // DEFINE_EXPRESSION_OP_BINARY(&, bit_and) // DEFINE_EXPRESSION_OP_BINARY(|, bit_or) -DEFINE_EXPRESSION_OP_BINARY (^, bit_xor) +DEFINE_EXPRESSION_OP_BINARY(^, bit_xor) DEFINE_EXPRESSION_OP_BINARY(<<, bit_shl) DEFINE_EXPRESSION_OP_BINARY(>>, bit_sar) DEFINE_EXPRESSION_OP_BINARY(<, cmp_lt) diff --git a/taichi/program/async_utils.h b/taichi/program/async_utils.h index a23380c45fef62..8952f8b9404d95 100644 --- a/taichi/program/async_utils.h +++ b/taichi/program/async_utils.h @@ -167,8 +167,8 @@ TLANG_NAMESPACE_END namespace std { template <> struct hash { - std::size_t operator()(const taichi::lang::IRHandle &ir_handle) const - noexcept { + std::size_t operator()( + const taichi::lang::IRHandle &ir_handle) const noexcept { return ir_handle.hash(); } }; diff --git a/taichi/program/callable.cpp b/taichi/program/callable.cpp index 095c04fc8e94ab..c97be6f1c9bbd8 100644 --- a/taichi/program/callable.cpp +++ b/taichi/program/callable.cpp @@ -9,7 +9,7 @@ int Callable::insert_arg(const DataType &dt, bool is_external_array) { return (int)args.size() - 1; } -int Callable::insert_ret(const DataType& dt) { +int Callable::insert_ret(const DataType &dt) { rets.emplace_back(dt->get_compute_type()); return (int)rets.size() - 1; } diff --git a/taichi/runtime/llvm/node_pointer.h b/taichi/runtime/llvm/node_pointer.h index 0886dc0357f35e..2e2554ed128fc3 100644 --- a/taichi/runtime/llvm/node_pointer.h +++ b/taichi/runtime/llvm/node_pointer.h @@ -48,16 +48,17 @@ void Pointer_activate(Ptr meta_, Ptr node, int i) { // The cuda_ calls will return 0 or do noop on CPUs u32 mask = cuda_active_mask(); if (is_representative(mask, (u64)lock)) { - locked_task(lock, - [&] { - auto rt = meta->context->runtime; - auto alloc = rt->node_allocators[meta->snode_id]; - auto allocated = (u64)alloc->allocate(); - // TODO: Not sure if we really need atomic_exchange here, - // just to be safe. - atomic_exchange_u64((u64 *)data_ptr, allocated); - }, - [&]() { return *data_ptr == nullptr; }); + locked_task( + lock, + [&] { + auto rt = meta->context->runtime; + auto alloc = rt->node_allocators[meta->snode_id]; + auto allocated = (u64)alloc->allocate(); + // TODO: Not sure if we really need atomic_exchange here, + // just to be safe. + atomic_exchange_u64((u64 *)data_ptr, allocated); + }, + [&]() { return *data_ptr == nullptr; }); } warp_barrier(mask); } diff --git a/taichi/transforms/auto_diff.cpp b/taichi/transforms/auto_diff.cpp index a3d7d816693603..04c5a4c00adfb0 100644 --- a/taichi/transforms/auto_diff.cpp +++ b/taichi/transforms/auto_diff.cpp @@ -192,18 +192,16 @@ class ReplaceLocalVarWithStacks : public BasicStmtVisitor { void visit(AllocaStmt *alloc) override { TI_ASSERT(alloc->width() == 1); - bool load_only = irpass::analysis::gather_statements( - alloc->parent, - [&](Stmt *s) { - if (auto store = s->cast()) - return store->dest == alloc; - else if (auto atomic = s->cast()) { - return atomic->dest == alloc; - } else { - return false; - } - }) - .empty(); + bool load_only = + irpass::analysis::gather_statements(alloc->parent, [&](Stmt *s) { + if (auto store = s->cast()) + return store->dest == alloc; + else if (auto atomic = s->cast()) { + return atomic->dest == alloc; + } else { + return false; + } + }).empty(); if (!load_only) { auto dtype = alloc->ret_type; auto stack_alloca = Stmt::make(dtype, ad_stack_size); diff --git a/taichi/transforms/inlining.cpp b/taichi/transforms/inlining.cpp index 46ba5bcfd7e8ae..0ed35dd9a4f058 100644 --- a/taichi/transforms/inlining.cpp +++ b/taichi/transforms/inlining.cpp @@ -34,9 +34,9 @@ class Inliner : public BasicStmtVisitor { modifier.replace_with(stmt, std::move(inlined_ir->as()->statements)); } else { - if (irpass::analysis::gather_statements( - inlined_ir.get(), [&](Stmt *s) { return s->is(); }) - .size() > 1) { + if (irpass::analysis::gather_statements(inlined_ir.get(), [&](Stmt *s) { + return s->is(); + }).size() > 1) { TI_WARN( "Multiple returns in function \"{}\" may not be handled properly.", func->get_name()); diff --git a/tests/python/bls_test_template.py b/tests/python/bls_test_template.py index aa3eaab05b99d0..476a4e9c0325d4 100644 --- a/tests/python/bls_test_template.py +++ b/tests/python/bls_test_template.py @@ -1,7 +1,9 @@ -import taichi as ti -import numpy as np import random +import numpy as np + +import taichi as ti + def bls_test_template(dim, N, diff --git a/tests/python/fuse_test_template.py b/tests/python/fuse_test_template.py index 71a1d36d4b30a9..1a8b1a8738ed11 100644 --- a/tests/python/fuse_test_template.py +++ b/tests/python/fuse_test_template.py @@ -1,6 +1,7 @@ -import taichi as ti import time +import taichi as ti + def template_fuse_dense_x2y2z(size=1024**3, repeat=10, diff --git a/tests/python/test_assert.py b/tests/python/test_assert.py index db36d8798483ec..a960d3663a1bb1 100644 --- a/tests/python/test_assert.py +++ b/tests/python/test_assert.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + @ti.require(ti.extension.assertion) @ti.all_archs_with(debug=True, gdb_trigger=False) diff --git a/tests/python/test_async.py b/tests/python/test_async.py index 7f0ca451d0ec86..bc2e8867f8dd6f 100644 --- a/tests/python/test_async.py +++ b/tests/python/test_async.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + @ti.test(require=ti.extension.async_mode, async_mode=True) def test_simple(): diff --git a/tests/python/test_bit_array.py b/tests/python/test_bit_array.py index ab42d7eb062f7e..af5cefd74ef2a9 100644 --- a/tests/python/test_bit_array.py +++ b/tests/python/test_bit_array.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + @ti.test(require=ti.extension.quant, debug=True) def test_1D_bit_array(): diff --git a/tests/python/test_bit_array_vectorization.py b/tests/python/test_bit_array_vectorization.py index 06ca0e26ad0eaf..9fad21763e919c 100644 --- a/tests/python/test_bit_array_vectorization.py +++ b/tests/python/test_bit_array_vectorization.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + @ti.test(require=ti.extension.quant, debug=True, cfg_optimization=False) def test_vectorized_struct_for(): diff --git a/tests/python/test_bit_operations.py b/tests/python/test_bit_operations.py index f069547a63c399..28956f1a2f19ae 100644 --- a/tests/python/test_bit_operations.py +++ b/tests/python/test_bit_operations.py @@ -1,9 +1,11 @@ -import taichi as ti -import numpy as np import operator as ops -from taichi import allclose + +import numpy as np import pytest +import taichi as ti +from taichi import allclose + @ti.test() def test_bit_shl(): diff --git a/tests/python/test_bit_struct.py b/tests/python/test_bit_struct.py index ab048f6cad17a9..cb4e1c53ee1c54 100644 --- a/tests/python/test_bit_struct.py +++ b/tests/python/test_bit_struct.py @@ -1,7 +1,8 @@ -import taichi as ti import numpy as np from pytest import approx +import taichi as ti + @ti.test(require=ti.extension.quant_basic, debug=True) def test_simple_array(): diff --git a/tests/python/test_bls_assume_in_range.py b/tests/python/test_bls_assume_in_range.py index 6750f10b4e2b03..7fdb836a7b85b2 100644 --- a/tests/python/test_bls_assume_in_range.py +++ b/tests/python/test_bls_assume_in_range.py @@ -1,4 +1,5 @@ import taichi as ti + from .bls_test_template import bls_particle_grid diff --git a/tests/python/test_cast.py b/tests/python/test_cast.py index fdcbbe91811470..b97df4d8dd5557 100644 --- a/tests/python/test_cast.py +++ b/tests/python/test_cast.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + @ti.all_archs def test_cast_f32(): diff --git a/tests/python/test_cli.py b/tests/python/test_cli.py index 569660b9c0411b..989a437166cecf 100644 --- a/tests/python/test_cli.py +++ b/tests/python/test_cli.py @@ -1,12 +1,14 @@ -import taichi as ti -from taichi.main import TaichiMain -from contextlib import contextmanager -import sys -import copy -import pytest -from unittest.mock import patch import argparse +import copy +import sys +from contextlib import contextmanager from pathlib import Path +from unittest.mock import patch + +import pytest +from taichi.main import TaichiMain + +import taichi as ti @contextmanager diff --git a/tests/python/test_constant_func.py b/tests/python/test_constant_func.py index 8fb3dd09de4699..1464a18abb425a 100644 --- a/tests/python/test_constant_func.py +++ b/tests/python/test_constant_func.py @@ -1,9 +1,11 @@ -import taichi as ti -import numpy as np import operator as ops -from taichi import allclose + +import numpy as np import pytest +import taichi as ti +from taichi import allclose + binary_func_table = [ (ops.add, ) * 2, (ops.sub, ) * 2, diff --git a/tests/python/test_custom_float.py b/tests/python/test_custom_float.py index dfb713fde377c4..5e3280007b711a 100644 --- a/tests/python/test_custom_float.py +++ b/tests/python/test_custom_float.py @@ -1,7 +1,9 @@ -import taichi as ti import math + from pytest import approx +import taichi as ti + @ti.test(require=ti.extension.quant_basic) def test_custom_float(): diff --git a/tests/python/test_custom_float_exponents.py b/tests/python/test_custom_float_exponents.py index 8f1ae915969f66..6ecd324e94e75b 100644 --- a/tests/python/test_custom_float_exponents.py +++ b/tests/python/test_custom_float_exponents.py @@ -1,8 +1,9 @@ -import taichi as ti import numpy as np import pytest from pytest import approx +import taichi as ti + @ti.test(require=ti.extension.quant) def test_custom_float_unsigned(): diff --git a/tests/python/test_custom_float_shared_exp.py b/tests/python/test_custom_float_shared_exp.py index 2ba91f930e71f4..d27fad67622d9a 100644 --- a/tests/python/test_custom_float_shared_exp.py +++ b/tests/python/test_custom_float_shared_exp.py @@ -1,6 +1,7 @@ -import taichi as ti -from pytest import approx import pytest +from pytest import approx + +import taichi as ti @pytest.mark.parametrize('exponent_bits', [5, 6, 7, 8]) diff --git a/tests/python/test_custom_float_time_integration.py b/tests/python/test_custom_float_time_integration.py index b02d59384709e9..4b8271a478c4ac 100644 --- a/tests/python/test_custom_float_time_integration.py +++ b/tests/python/test_custom_float_time_integration.py @@ -1,8 +1,10 @@ -import taichi as ti import math + import pytest from pytest import approx +import taichi as ti + @pytest.mark.parametrize('use_cft,use_exponent,use_shared_exp', [(False, False, False), (True, False, False), diff --git a/tests/python/test_custom_type_atomics.py b/tests/python/test_custom_type_atomics.py index 89fe66c89ca55b..01449c082fb19e 100644 --- a/tests/python/test_custom_type_atomics.py +++ b/tests/python/test_custom_type_atomics.py @@ -1,6 +1,7 @@ -import taichi as ti from pytest import approx +import taichi as ti + @ti.test(require=ti.extension.quant_basic, debug=True) def test_custom_int_atomics(): diff --git a/tests/python/test_debug.py b/tests/python/test_debug.py index 19a61d74602790..8b71544dd3b470 100644 --- a/tests/python/test_debug.py +++ b/tests/python/test_debug.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + def test_cpu_debug_snode_reader(): ti.init(arch=ti.x64, debug=True) diff --git a/tests/python/test_dynamic.py b/tests/python/test_dynamic.py index ec646f5b3d97cc..8ae306fe5774c0 100644 --- a/tests/python/test_dynamic.py +++ b/tests/python/test_dynamic.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + def ti_support_dynamic(test): return ti.archs_excluding(ti.cc)(test) diff --git a/tests/python/test_element_wise.py b/tests/python/test_element_wise.py index fe7588b358e10a..7895c69f974398 100644 --- a/tests/python/test_element_wise.py +++ b/tests/python/test_element_wise.py @@ -1,8 +1,9 @@ -import taichi as ti import numpy as np -from taichi import allclose import pytest +import taichi as ti +from taichi import allclose + def _c_mod(a, b): return a - b * int(float(a) / b) diff --git a/tests/python/test_external_func.py b/tests/python/test_external_func.py index 7c67dacfe984f3..6a0347128a1329 100644 --- a/tests/python/test_external_func.py +++ b/tests/python/test_external_func.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + @ti.require(ti.extension.extfunc) @ti.archs_excluding(ti.cpu) diff --git a/tests/python/test_field.py b/tests/python/test_field.py index a672ea83c271ee..363e71f5d8cab8 100644 --- a/tests/python/test_field.py +++ b/tests/python/test_field.py @@ -2,9 +2,10 @@ To test our new `ti.field` API is functional (#1500) ''' -import taichi as ti import pytest +import taichi as ti + data_types = [ti.i32, ti.f32, ti.i64, ti.f64] field_shapes = [(), 8, (6, 12)] vector_dims = [3] diff --git a/tests/python/test_fuse_dense.py b/tests/python/test_fuse_dense.py index 59fdad84c67713..3afbf48912d678 100644 --- a/tests/python/test_fuse_dense.py +++ b/tests/python/test_fuse_dense.py @@ -1,6 +1,7 @@ import taichi as ti -from .fuse_test_template import template_fuse_dense_x2y2z, \ - template_fuse_reduction + +from .fuse_test_template import (template_fuse_dense_x2y2z, + template_fuse_reduction) @ti.test(require=ti.extension.async_mode, async_mode=True) diff --git a/tests/python/test_fuse_dynamic.py b/tests/python/test_fuse_dynamic.py index 216cec4f6f8f6f..dedae5000d353c 100644 --- a/tests/python/test_fuse_dynamic.py +++ b/tests/python/test_fuse_dynamic.py @@ -1,7 +1,9 @@ -import taichi as ti import time + import pytest +import taichi as ti + def benchmark_fuse_dynamic_x2y2z(size=1024**2, repeat=10, first_n=100): x = ti.field(ti.i32) diff --git a/tests/python/test_get_external_tensor_shape.py b/tests/python/test_get_external_tensor_shape.py index 70328d059a39b4..7e9f04e2b20ca8 100644 --- a/tests/python/test_get_external_tensor_shape.py +++ b/tests/python/test_get_external_tensor_shape.py @@ -1,7 +1,8 @@ -import taichi as ti import numpy as np import pytest +import taichi as ti + if ti.has_pytorch(): import torch diff --git a/tests/python/test_gui.py b/tests/python/test_gui.py index 77aef8e51e5a40..836111f1248c7d 100644 --- a/tests/python/test_gui.py +++ b/tests/python/test_gui.py @@ -1,8 +1,9 @@ -from taichi import make_temp_file -import taichi as ti import numpy as np import pytest +import taichi as ti +from taichi import make_temp_file + @ti.host_arch_only @pytest.mark.parametrize('dtype', [ti.u8, ti.f32]) diff --git a/tests/python/test_image_io.py b/tests/python/test_image_io.py index 1a2768dc7ac3f8..eda380c5788f96 100644 --- a/tests/python/test_image_io.py +++ b/tests/python/test_image_io.py @@ -1,8 +1,10 @@ -import taichi as ti +import os + import numpy as np -from taichi import make_temp_file import pytest -import os + +import taichi as ti +from taichi import make_temp_file # jpg is also supported but hard to test here since it's lossy: diff --git a/tests/python/test_kernel_arg_errors.py b/tests/python/test_kernel_arg_errors.py index 4b6404cfe77686..9df816cd92f199 100644 --- a/tests/python/test_kernel_arg_errors.py +++ b/tests/python/test_kernel_arg_errors.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + @ti.test(arch=ti.cpu) def test_pass_float_as_i32(): diff --git a/tests/python/test_lang.py b/tests/python/test_lang.py index 6aa6c13cfeaf04..04027a9af78547 100644 --- a/tests/python/test_lang.py +++ b/tests/python/test_lang.py @@ -1,7 +1,8 @@ -import taichi as ti import numpy as np import pytest +import taichi as ti + @ti.all_archs def test_nested_subscript(): diff --git a/tests/python/test_linalg.py b/tests/python/test_linalg.py index 433c8b3aadd304..cf5431f5115419 100644 --- a/tests/python/test_linalg.py +++ b/tests/python/test_linalg.py @@ -1,8 +1,10 @@ -import taichi as ti +import math + import numpy as np -from taichi import approx import pytest -import math + +import taichi as ti +from taichi import approx @ti.all_archs diff --git a/tests/python/test_listgen.py b/tests/python/test_listgen.py index 59c73165607f0d..0711d3eaa2ffcd 100644 --- a/tests/python/test_listgen.py +++ b/tests/python/test_listgen.py @@ -1,6 +1,7 @@ -import taichi as ti from random import randrange +import taichi as ti + @ti.all_archs def test_listgen(): diff --git a/tests/python/test_matrix.py b/tests/python/test_matrix.py index 626e4326778061..5b26574e192492 100644 --- a/tests/python/test_matrix.py +++ b/tests/python/test_matrix.py @@ -1,9 +1,11 @@ -import taichi as ti -import numpy as np -from taichi import approx import operator + +import numpy as np import pytest +import taichi as ti +from taichi import approx + operation_types = [operator.add, operator.sub, operator.matmul] test_matrix_arrays = [ np.array([[1, 2], [3, 4]]), diff --git a/tests/python/test_matrix_different_type.py b/tests/python/test_matrix_different_type.py index 318067d9014520..ee2de6f7f20c1a 100644 --- a/tests/python/test_matrix_different_type.py +++ b/tests/python/test_matrix_different_type.py @@ -1,6 +1,7 @@ -import taichi as ti from pytest import approx +import taichi as ti + # TODO: test more matrix operations @ti.test() diff --git a/tests/python/test_mpm88.py b/tests/python/test_mpm88.py index eb9176cf4021de..bfbadbcbf36dce 100644 --- a/tests/python/test_mpm88.py +++ b/tests/python/test_mpm88.py @@ -1,7 +1,9 @@ +import os + import pytest + import taichi as ti from taichi import approx -import os def run_mpm88_test(): diff --git a/tests/python/test_mpm_particle_list.py b/tests/python/test_mpm_particle_list.py index 8c8976735f1018..61e3febb05b71a 100644 --- a/tests/python/test_mpm_particle_list.py +++ b/tests/python/test_mpm_particle_list.py @@ -1,6 +1,7 @@ -import taichi as ti import random +import taichi as ti + @ti.data_oriented class MPMSolver: diff --git a/tests/python/test_native_functions.py b/tests/python/test_native_functions.py index 25c037bb79ff63..cc356de3500671 100644 --- a/tests/python/test_native_functions.py +++ b/tests/python/test_native_functions.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + @ti.all_archs def test_abs(): diff --git a/tests/python/test_ndrange.py b/tests/python/test_ndrange.py index cc32561bc6765e..d73529dfc9c4e2 100644 --- a/tests/python/test_ndrange.py +++ b/tests/python/test_ndrange.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + @ti.all_archs def test_1d(): diff --git a/tests/python/test_numpy.py b/tests/python/test_numpy.py index cd82ceaaff270e..2f9f7de926853c 100644 --- a/tests/python/test_numpy.py +++ b/tests/python/test_numpy.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + def with_data_type(dt): val = ti.field(ti.i32) diff --git a/tests/python/test_numpy_io.py b/tests/python/test_numpy_io.py index e95aac89d4a465..68fbd48ab7d815 100644 --- a/tests/python/test_numpy_io.py +++ b/tests/python/test_numpy_io.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + @ti.all_archs def test_to_numpy_2d(): diff --git a/tests/python/test_print.py b/tests/python/test_print.py index 3a7d6061e606fc..724b56a6b6f503 100644 --- a/tests/python/test_print.py +++ b/tests/python/test_print.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + # Not really testable.. # Just making sure it does not crash diff --git a/tests/python/test_reduction.py b/tests/python/test_reduction.py index e7db7c730efda6..331f9868f1dbf6 100644 --- a/tests/python/test_reduction.py +++ b/tests/python/test_reduction.py @@ -1,6 +1,7 @@ -import taichi as ti from pytest import approx +import taichi as ti + def _test_reduction_single(dtype, criterion): N = 1024 * 1024 diff --git a/tests/python/test_runtime.py b/tests/python/test_runtime.py index 1a8d8bf6214d81..c1c6b8a45f979e 100644 --- a/tests/python/test_runtime.py +++ b/tests/python/test_runtime.py @@ -1,8 +1,12 @@ -import taichi as ti -import sys, os, copy +import copy +import os +import sys from contextlib import contextmanager + import pytest +import taichi as ti + @contextmanager def patch_os_environ_helper(custom_environ: dict, excludes: dict): diff --git a/tests/python/test_sfg.py b/tests/python/test_sfg.py index f08b822f7f3950..d4221913dc540a 100644 --- a/tests/python/test_sfg.py +++ b/tests/python/test_sfg.py @@ -1,7 +1,8 @@ -import taichi as ti import numpy as np import pytest +import taichi as ti + @ti.test(require=[ti.extension.async_mode, ti.extension.sparse], async_mode=True) diff --git a/tests/python/test_ssa.py b/tests/python/test_ssa.py index 5bfa7d1e44d7f5..8e3b0038a03b36 100644 --- a/tests/python/test_ssa.py +++ b/tests/python/test_ssa.py @@ -3,10 +3,12 @@ 1. Ensure working well when computation result is assigned to self. 2. Prevent duplicate-evaluation on expression with side-effect like random. ''' -import taichi as ti +import math + import numpy as np + +import taichi as ti from taichi import approx -import math @ti.all_archs diff --git a/tests/python/test_static.py b/tests/python/test_static.py index bb3e769efda327..08f36eab9fea73 100644 --- a/tests/python/test_static.py +++ b/tests/python/test_static.py @@ -1,7 +1,8 @@ -import taichi as ti import numpy as np import pytest +import taichi as ti + @pytest.mark.parametrize('val', [0, 1]) @ti.test(ti.cpu) diff --git a/tests/python/test_svd.py b/tests/python/test_svd.py index 18a6e77a2062d9..883a7bd4772632 100644 --- a/tests/python/test_svd.py +++ b/tests/python/test_svd.py @@ -1,5 +1,6 @@ -import taichi as ti import numpy as np + +import taichi as ti from taichi import approx diff --git a/tests/python/test_syntax_errors.py b/tests/python/test_syntax_errors.py index 27e8f146cc725d..aabbe2d0c59ce2 100644 --- a/tests/python/test_syntax_errors.py +++ b/tests/python/test_syntax_errors.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + @ti.must_throw(ti.TaichiSyntaxError) def test_try(): diff --git a/tests/python/test_tensor_reflection.py b/tests/python/test_tensor_reflection.py index 313c4cce8a2f44..5bac5cdbbe31a2 100644 --- a/tests/python/test_tensor_reflection.py +++ b/tests/python/test_tensor_reflection.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + @ti.all_archs def test_POT(): diff --git a/tests/python/test_test.py b/tests/python/test_test.py index 7652a72fd0399f..855879dde405ff 100644 --- a/tests/python/test_test.py +++ b/tests/python/test_test.py @@ -4,9 +4,10 @@ TODO: Skips these tests after all tests are using @ti.test ''' -import taichi as ti import pytest +import taichi as ti + ### `ti.test` diff --git a/tests/python/test_torch_ad.py b/tests/python/test_torch_ad.py index adeec4a583722a..73198fcf2bd381 100644 --- a/tests/python/test_torch_ad.py +++ b/tests/python/test_torch_ad.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + if ti.has_pytorch(): import torch diff --git a/tests/python/test_torch_io.py b/tests/python/test_torch_io.py index f11991d3ab547e..1a06c2177deb82 100644 --- a/tests/python/test_torch_io.py +++ b/tests/python/test_torch_io.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + if ti.has_pytorch(): import torch diff --git a/tests/python/test_trailing_bits.py b/tests/python/test_trailing_bits.py index 0578753d1df659..ece5a1d738d2a0 100644 --- a/tests/python/test_trailing_bits.py +++ b/tests/python/test_trailing_bits.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + def _test_trailing_bits(): ti.init(arch=ti.cpu, debug=True, print_ir=True) diff --git a/tests/python/test_types.py b/tests/python/test_types.py index b9bfc594c760f5..50bd28ef5162cd 100644 --- a/tests/python/test_types.py +++ b/tests/python/test_types.py @@ -1,6 +1,7 @@ -import taichi as ti import pytest +import taichi as ti + _TI_TYPES = [ti.i8, ti.i16, ti.i32, ti.u8, ti.u16, ti.u32, ti.f32] _TI_64_TYPES = [ti.i64, ti.u64, ti.f64] diff --git a/tests/python/test_unary_ops.py b/tests/python/test_unary_ops.py index 7f00ea8bbdbc8f..774efdd03424b9 100644 --- a/tests/python/test_unary_ops.py +++ b/tests/python/test_unary_ops.py @@ -1,6 +1,7 @@ -import taichi as ti import numpy as np +import taichi as ti + def _test_op(dt, taichi_op, np_op): print('arch={} default_fp={}'.format(ti.cfg.arch, ti.cfg.default_fp))