SIGN IN SIGN UP
taichi-dev / taichi UNCLAIMED

Productive, portable, and performant GPU programming in Python.

0 0 0 C++
import atexit
import functools
import math
import shutil
import threading
from os import listdir, rmdir, stat
from os.path import join
from tempfile import mkdtemp
import pytest
import taichi as ti
from tests import test_utils
OFFLINE_CACHE_TEMP_DIR = mkdtemp()
atexit.register(lambda: rmdir(OFFLINE_CACHE_TEMP_DIR))
supported_llvm_archs = {ti.cpu, ti.cuda, ti.amdgpu}
supported_gfx_archs = {ti.opengl, ti.vulkan, ti.metal, ti.dx11}
supported_archs_offline_cache = supported_llvm_archs | supported_gfx_archs
supported_archs_offline_cache = {v for v in supported_archs_offline_cache if v in test_utils.expected_archs()}
def is_offline_cache_file(filename):
suffixes = (".tic",)
return filename.endswith(suffixes)
def cache_files_size(path):
files = listdir(path)
result = 0
for file in files:
if is_offline_cache_file(file):
result += stat(join(path, file)).st_size
return result
def expected_num_cache_files(num_kernels: int = 0) -> int:
if num_kernels == 0:
return 0
# code files(*.tic) + metadata files(ticache.tcb)
return num_kernels + 1
def tmp_offline_cache_file_path():
return join(OFFLINE_CACHE_TEMP_DIR, str(threading.current_thread().ident))
def current_thread_ext_options():
return {
"offline_cache": True,
"offline_cache_file_path": tmp_offline_cache_file_path(),
"cuda_stack_limit": 1024,
"device_memory_GB": 0.2,
}
def cache_files_cnt():
try:
return len(listdir(tmp_offline_cache_file_path()))
except FileNotFoundError:
return 0
@ti.kernel
def kernel0() -> ti.i32:
return 1
def python_kernel0():
return 1
@ti.kernel
def kernel1(a: ti.i32, b: ti.i32, c: ti.f32) -> ti.f32:
return a / b + c * b - c + a**2 + ti.log(c)
def python_kernel1(a, b, c):
return a / b + c * b - c + a**2 + math.log(c)
@ti.kernel
def kernel2(n: ti.i32) -> ti.i32:
x = 0
for i in range(n):
ti.atomic_add(x, 1)
return x
def python_kernel2(n):
return n
def kernel3(a, mat):
mat_type = ti.types.matrix(mat.n, mat.m, ti.i32)
@ti.kernel
def kernel(u: ti.i32, v: mat_type) -> mat_type:
return u * v
return kernel(a, mat)
def python_kernel3(a, mat):
return a * mat
@ti.func
def func_sum(lo: ti.i32, hi: ti.i32) -> ti.i32:
res = 0
for i in range(lo, hi):
res += i
return res
@ti.func
def func_mul(lo: ti.i32, hi: ti.i32) -> ti.i32:
res = 1
for i in range(lo, hi):
res *= i
return res
@ti.kernel
def kernel4(lo: ti.i32, hi: ti.i32, n: ti.i32) -> ti.i32:
res = 0
for i in range(n):
res += func_sum(lo, hi)
return res
def python_kernel4(lo: ti.i32, hi: ti.i32, n: ti.i32):
res = 0
for i in range(n):
for j in range(lo, hi):
res += j
return res
@ti.kernel
def kernel5(lo: ti.i32, hi: ti.i32, n: ti.i32) -> ti.i32:
res = 1
for i in range(n):
res *= func_mul(lo, hi)
return res
def python_kernel5(lo: ti.i32, hi: ti.i32, n: ti.i32):
res = 1
for i in range(n):
for j in range(lo, hi):
res *= j
return res
simple_kernels_to_test = [
(kernel0, (), python_kernel0),
(kernel1, (100, 200, 10.2), python_kernel1),
(kernel2, (1024,), python_kernel2),
# FIXME: add this kernel back once we have a better way to compare matrices
# with test_utils.approx()
# (kernel3, (10, ti.Matrix([[1, 2], [256, 1024]],
# ti.i32)), python_kernel3),
# FIXME: add this kernel back once #6221 is fixed
# (kernel4, (1, 10, 2), python_kernel4),
(kernel5, (1, 2, 2), python_kernel5),
]
def _test_offline_cache_dec(func):
@functools.wraps(func)
def wrapped(*args, **kwargs):
test_utils.mkdir_p(tmp_offline_cache_file_path())
ret = None
try:
ret = func(*args, **kwargs)
except Exception as e:
raise e
finally:
ti.reset()
shutil.rmtree(tmp_offline_cache_file_path())
return ret
return wrapped
@_test_offline_cache_dec
def _test_offline_cache_for_a_kernel(curr_arch, kernel, args, result):
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
res1 = kernel(*args)
assert added_files() == expected_num_cache_files()
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
assert added_files() == expected_num_cache_files(1)
res2 = kernel(*args)
assert res1 == test_utils.approx(result) and res1 == test_utils.approx(res2)
ti.reset()
assert added_files() == expected_num_cache_files(1)
@_test_offline_cache_dec
def _test_closing_offline_cache_for_a_kernel(curr_arch, kernel, args, result):
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
def my_init():
ti.init(
arch=curr_arch,
enable_fallback=False,
offline_cache=False,
offline_cache_file_path=tmp_offline_cache_file_path(),
cuda_stack_limit=1024,
device_memory_GB=0.1,
)
my_init()
res1 = kernel(*args)
assert added_files() == expected_num_cache_files()
my_init()
assert added_files() == expected_num_cache_files()
res2 = kernel(*args)
assert res1 == test_utils.approx(result) and res1 == test_utils.approx(res2)
ti.reset()
assert added_files() == expected_num_cache_files()
@pytest.mark.parametrize("curr_arch", supported_archs_offline_cache)
def test_closing_offline_cache(curr_arch):
for kernel, args, get_res in simple_kernels_to_test:
_test_closing_offline_cache_for_a_kernel(curr_arch=curr_arch, kernel=kernel, args=args, result=get_res(*args))
@pytest.mark.parametrize("curr_arch", supported_archs_offline_cache)
def test_offline_cache_per_kernel(curr_arch):
for kernel, args, get_res in simple_kernels_to_test:
_test_offline_cache_for_a_kernel(curr_arch=curr_arch, kernel=kernel, args=args, result=get_res(*args))
@pytest.mark.parametrize("curr_arch", supported_archs_offline_cache)
@_test_offline_cache_dec
def test_multiple_ib_with_offline_cache(curr_arch):
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
def helper():
x = ti.field(float, (), needs_grad=True)
y = ti.field(float, (), needs_grad=True)
@ti.kernel
def compute_y():
for j in range(2):
for i in range(3):
y[None] += x[None]
for i in range(3):
y[None] += x[None]
x[None] = 1.0
with ti.ad.Tape(y):
compute_y()
assert y[None] == 12.0
assert x.grad[None] == 12.0
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
helper()
assert added_files() == expected_num_cache_files()
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
assert added_files() == expected_num_cache_files(9)
helper()
ti.reset()
assert added_files() == expected_num_cache_files(9)
@pytest.mark.parametrize("curr_arch", supported_archs_offline_cache)
@_test_offline_cache_dec
def test_calling_a_kernel_with_different_param_list(curr_arch):
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
mat_type = ti.types.matrix(2, 3, ti.i32)
@ti.kernel
def kernel(a: mat_type, b: mat_type) -> mat_type:
return a + 10 * b
def np_kernel(a, b):
return a + 10 * b
mat1 = ti.Matrix([[1, 2, 3], [3, 2, 1]], ti.i32)
mat2 = ti.Matrix([[1, 2, 3], [3, 2, 1]], ti.i32)
mat3 = ti.Matrix([[1, 2, 3], [3, 2, 1]], ti.i32)
np_mat1 = mat1.to_numpy()
np_mat2 = mat2.to_numpy()
np_mat3 = mat3.to_numpy()
assert added_files() == expected_num_cache_files()
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
assert (kernel(mat1, mat1).to_numpy() == np_kernel(np_mat1, np_mat1)).all()
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
assert added_files() == expected_num_cache_files(1)
assert (kernel(mat1, mat1).to_numpy() == np_kernel(np_mat1, np_mat1)).all()
assert (kernel(mat1, mat2).to_numpy() == np_kernel(np_mat1, np_mat2)).all()
assert (kernel(mat2, mat2).to_numpy() == np_kernel(np_mat2, np_mat2)).all()
assert (kernel(mat2, mat3).to_numpy() == np_kernel(np_mat2, np_mat3)).all()
ti.reset()
assert added_files() == expected_num_cache_files(1)
@pytest.mark.parametrize("curr_arch", supported_archs_offline_cache)
@_test_offline_cache_dec
def test_snode_reader_and_writer_with_offline_cache(curr_arch):
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
def helper():
x = ti.field(dtype=ti.f32, shape=())
y = ti.field(dtype=ti.f32, shape=())
x[None] = 3.14
y[None] = 4.14
assert x[None] == test_utils.approx(3.14)
assert y[None] == test_utils.approx(4.14)
x[None] = 6.28
y[None] = 7.28
assert x[None] == test_utils.approx(6.28)
assert y[None] == test_utils.approx(7.28)
assert added_files() == expected_num_cache_files()
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
helper()
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
assert added_files() == expected_num_cache_files(4)
helper()
ti.reset()
assert added_files() == expected_num_cache_files(4)
@pytest.mark.parametrize("curr_arch", supported_archs_offline_cache)
@_test_offline_cache_dec
def test_calling_many_kernels(curr_arch):
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
def helper():
for kernel, args, get_res in simple_kernels_to_test:
assert kernel(*args) == test_utils.approx(get_res(*args))
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
helper()
assert added_files() == expected_num_cache_files()
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
assert added_files() == expected_num_cache_files(len(simple_kernels_to_test))
helper()
ti.reset()
assert added_files() == expected_num_cache_files(len(simple_kernels_to_test))
@pytest.mark.parametrize("curr_arch", supported_archs_offline_cache)
@_test_offline_cache_dec
def test_offline_cache_with_different_snode_trees(curr_arch):
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
def helper():
x = ti.field(float, shape=5)
@ti.kernel
def trigger_compile():
x[0] += 1
# This case is used for testing SNodeTree storing order matters (i.e., use a ordered container such as vector instead of unordered_map or unordered_set) when generating kernel offline cache key
# The multiple `trigger_compile` equalivant to allocate each field to a different SNodeTree
# i.e.,
# x = ti.field(float)
# fb.dense(ti.i, 5).place(x)
# fb.finalize()
trigger_compile()
a = ti.field(float, shape=5)
trigger_compile()
b = ti.field(float, shape=10)
trigger_compile()
c = ti.field(float, shape=5)
trigger_compile()
d = ti.field(float, shape=10)
trigger_compile()
e = ti.field(float, shape=5)
trigger_compile()
f = ti.field(float, shape=10)
trigger_compile()
g = ti.field(float, shape=5)
trigger_compile()
h = ti.field(float, shape=10)
@ti.kernel
def kernel_forward():
for i in range(5):
a[i] += i
b[i] += i
c[i] += i
d[i] += i
e[i] += i
f[i] += i
g[i] += i
h[i] += i
kernel_forward()
assert added_files() == expected_num_cache_files(0)
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
helper()
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
assert added_files() == expected_num_cache_files(2)
helper()
# The number of cache file should not change
for _ in range(5):
ti.init(arch=curr_arch, enable_fallback=False, **current_thread_ext_options())
assert added_files() == expected_num_cache_files(2)
helper()
@pytest.mark.parametrize("curr_arch", supported_archs_offline_cache)
@_test_offline_cache_dec
def test_offline_cache_with_changing_compile_config(curr_arch):
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
@ti.kernel
def helper():
b = 200
c = 0
for i in range(b):
c += i
assert added_files() == expected_num_cache_files()
ti.init(arch=curr_arch, enable_fallback=False, opt_level=0, **current_thread_ext_options())
helper()
ti.init(arch=curr_arch, enable_fallback=False, opt_level=1, **current_thread_ext_options())
assert added_files() == expected_num_cache_files(1)
helper()
ti.reset()
assert added_files() == expected_num_cache_files(2)
ti.init(arch=curr_arch, enable_fallback=False, default_fp=ti.f32, **current_thread_ext_options())
helper()
ti.reset()
assert added_files() == expected_num_cache_files(2)
@pytest.mark.parametrize("curr_arch", supported_archs_offline_cache)
@pytest.mark.parametrize("factor", [0.0, 0.25, 0.85, 1.0])
@pytest.mark.parametrize("policy", ["never", "version", "lru", "fifo"])
@_test_offline_cache_dec
def test_offline_cache_cleaning(curr_arch, factor, policy):
def only_init(max_size):
ti.init(
arch=curr_arch,
enable_fallback=False,
offline_cache_cleaning_policy=policy,
offline_cache_max_size_of_files=max_size, # bytes
offline_cache_cleaning_factor=factor,
**current_thread_ext_options()
)
def run_simple_kernels(max_size):
only_init(max_size)
for kernel, args, get_res in simple_kernels_to_test:
assert kernel(*args) == test_utils.approx(get_res(*args))
kernel_count = len(simple_kernels_to_test)
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
assert added_files() == expected_num_cache_files()
run_simple_kernels(1024**3) # 1GB (>> size_of_cache_files)
ti.reset() # Dumping cache data
size_of_cache_files = cache_files_size(tmp_offline_cache_file_path())
assert added_files() == expected_num_cache_files(kernel_count)
only_init(size_of_cache_files * 2)
ti.reset()
assert added_files() == expected_num_cache_files(kernel_count)
only_init(1) # 1B (<< size_of_cache_files)
ti.reset()
rem = []
if policy in ["never", "version"]:
rem = kernel_count
else:
lo = -min(kernel_count - int(factor * kernel_count), kernel_count)
lo = kernel_count if lo == 0 else lo
rem = len(simple_kernels_to_test[lo:])
assert added_files() == expected_num_cache_files(rem)
# FIXME: Change to `supported_archs_offline_cache` after fixing bugs of real-function on gpu
@pytest.mark.run_in_serial
@pytest.mark.parametrize("curr_arch", {ti.cpu, ti.cuda} & supported_archs_offline_cache)
@_test_offline_cache_dec
[cuda] Only set CU_LIMIT_STACK_SIZE when necessary (#7906) If we set it to 8192 unconditionally it'll increase CUDA memory usage by ~800MB on my 3090. It's not necessary when deep recursively real function call isn't involved so let's default to not set it. Issue: # ### Brief Summary <!-- copilot:summary --> ### <samp>🤖 Generated by Copilot at 378709f</samp> Removed redundant and potentially harmful stack size limit settings for CUDA kernels using the LLVM backend. Added a single stack size limit setting based on user preference or default value in the `ti.init` function. These changes aim to improve the performance and compatibility of Taichi on CUDA devices. ### Walkthrough <!-- copilot:walkthrough --> ### <samp>🤖 Generated by Copilot at 378709f</samp> * Remove redundant and potentially harmful stack size limit settings for every kernel launch on CUDA devices using the LLVM backend ([link](https://github.com/taichi-dev/taichi/pull/7906/files?diff=unified&w=0#diff-d8cf951dea16cc738fc36faee3a158f37a605abccd312c4a6f18d73fd2504332L137-L138)) * Set the stack size limit only once per device, and only if the user specifies a non-zero value in the `ti.init` function, in the function that defines the Taichi runtime functions on CUDA devices using the LLVM backend ([link](https://github.com/taichi-dev/taichi/pull/7906/files?diff=unified&w=0#diff-b9155792159f392bd8bacd44cb1819be5239b022d707499fc364c0f93dd8c5e5R124-R127)) * Change the default value of `cuda_stack_limit` in the `CompileConfig` struct from 8192 to 0, to avoid unnecessary stack size limit settings unless the user explicitly requests them ([link](https://github.com/taichi-dev/taichi/pull/7906/files?diff=unified&w=0#diff-604e33db5b7c4f4b819098dee1f2634a36c0974822c1538aec0d57d28ca00dd9L102-R102))
2023-04-27 16:00:46 +08:00
@test_utils.test(cuda_stack_limit=8192)
def test_offline_cache_for_kernels_calling_real_func(curr_arch):
count_of_cache_file = cache_files_cnt()
def added_files():
return cache_files_cnt() - count_of_cache_file
def helper1():
[Lang] Move real function out of the experimental module (#8399) Issue: # ### Brief Summary <!-- copilot:summary --> ### <samp>🤖 Generated by Copilot at 8c46960</samp> This pull request updates the `real_func` examples and tests to use the stable `ti.real_func` decorator and the new vector, matrix, and field APIs. It also refactors some code for readability, performance, and consistency, and removes the deprecated `experimental` module. ### Walkthrough <!-- copilot:walkthrough --> ### <samp>🤖 Generated by Copilot at 8c46960</samp> * Remove `experimental` module and move `real_func` decorator to top-level `taichi` module ([link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-e0e39c57bc92c0e1da4c831f8820a4082a3cc26a9b88962bee72964dfa26a74aL13-R13), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-a157043b38542c8145447ff342fda65fe4d54fb777fe514daa70007e83e20dc1L1235-R1235)) * Update all examples and tests that use `real_func` decorator to import it from `taichi` module instead of `experimental` module ([link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-9cbae292e39bc9b37e4cfab4bfef6f363788d06ecb7757753ee5c0ba7ad0afcdL51-R51), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-9cbae292e39bc9b37e4cfab4bfef6f363788d06ecb7757753ee5c0ba7ad0afcdL58-R58), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-9cbae292e39bc9b37e4cfab4bfef6f363788d06ecb7757753ee5c0ba7ad0afcdL72-R82), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-9cbae292e39bc9b37e4cfab4bfef6f363788d06ecb7757753ee5c0ba7ad0afcdL136-R136), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f6098eab3d43899f2146cafc8454b2d9bae86064709eb2eeb43e0edbf459f22fL33-R33), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f6098eab3d43899f2146cafc8454b2d9bae86064709eb2eeb43e0edbf459f22fL55-R55), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f6098eab3d43899f2146cafc8454b2d9bae86064709eb2eeb43e0edbf459f22fL93-R98), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-86667a5be06e5995f4b15199aa3d4b655e64cceed6d108185d5cc71e03f0bb1fL43-R53), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-86667a5be06e5995f4b15199aa3d4b655e64cceed6d108185d5cc71e03f0bb1fL89-R94), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-86667a5be06e5995f4b15199aa3d4b655e64cceed6d108185d5cc71e03f0bb1fL108-R108), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL51-R51), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL61-R61), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL68-R68), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL78-R78), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL89-R89), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL116-R116), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL124-R124), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL66-R71), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL80-R80), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL86-R86), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL95-R95), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL105-R105), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL119-R119), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL125-R125), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL132-R132), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL344-R344), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f2d9c4a8c504d2073549cde99b31a733a2f39bba4ce9f062cbaa09e1be17aa87L243-R243), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L11-R11), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L29-R29), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L49-R49), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L153-R153), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L185-R185), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L194-R194), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L207-R207), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L221-R221), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L242-R242), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L259-R259), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L281-R281), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L295-R295), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L310-R310), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L329-R329), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L352-R352), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L451-R451), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L469-R469), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L484-R484), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L501-R501), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L515-R515), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L533-R533), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-ca3c8d1edb25b6a7f4affbb79b2e3e74f73b3757e5d465258ce42ea9eb09fbc0L1035-R1035), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-ca3c8d1edb25b6a7f4affbb79b2e3e74f73b3757e5d465258ce42ea9eb09fbc0L1052-R1052), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-ca3c8d1edb25b6a7f4affbb79b2e3e74f73b3757e5d465258ce42ea9eb09fbc0L1134-R1134), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-ca3c8d1edb25b6a7f4affbb79b2e3e74f73b3757e5d465258ce42ea9eb09fbc0L1149-R1149), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-85fb062b09589d588166180e3c7954753018dc39474d3846346f2a8bb3faac32L480-R480), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-85fb062b09589d588166180e3c7954753018dc39474d3846346f2a8bb3faac32L494-R494), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-8641463ffef94529453efd945e02d056e4339417cc65aad074d4e6b3ef093244L49-R53), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-e206c19691b4678b401a5e2787000e93ab0b0060cbf4ac2625b51f87599d98d8L218-R218), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-e206c19691b4678b401a5e2787000e93ab0b0060cbf4ac2625b51f87599d98d8L238-R238), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-e206c19691b4678b401a5e2787000e93ab0b0060cbf4ac2625b51f87599d98d8L258-R258)) * Delete `python/taichi/experimental.py` file as it is no longer needed ([link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-60267f25c1e1dfada9487235a4a793b63a49aa3d33c93b8449e24ac821a33174))
2023-10-31 14:11:53 +08:00
@ti.real_func
def sum(l: ti.i32, r: ti.i32) -> ti.i32:
if l == r:
return l
else:
return sum(l, (l + r) // 2) + sum((l + r) // 2 + 1, r)
@ti.kernel
def get_sum() -> ti.i32:
return sum(0, 99)
assert get_sum() == 99 * 50
def helper2():
[Lang] Move real function out of the experimental module (#8399) Issue: # ### Brief Summary <!-- copilot:summary --> ### <samp>🤖 Generated by Copilot at 8c46960</samp> This pull request updates the `real_func` examples and tests to use the stable `ti.real_func` decorator and the new vector, matrix, and field APIs. It also refactors some code for readability, performance, and consistency, and removes the deprecated `experimental` module. ### Walkthrough <!-- copilot:walkthrough --> ### <samp>🤖 Generated by Copilot at 8c46960</samp> * Remove `experimental` module and move `real_func` decorator to top-level `taichi` module ([link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-e0e39c57bc92c0e1da4c831f8820a4082a3cc26a9b88962bee72964dfa26a74aL13-R13), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-a157043b38542c8145447ff342fda65fe4d54fb777fe514daa70007e83e20dc1L1235-R1235)) * Update all examples and tests that use `real_func` decorator to import it from `taichi` module instead of `experimental` module ([link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-9cbae292e39bc9b37e4cfab4bfef6f363788d06ecb7757753ee5c0ba7ad0afcdL51-R51), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-9cbae292e39bc9b37e4cfab4bfef6f363788d06ecb7757753ee5c0ba7ad0afcdL58-R58), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-9cbae292e39bc9b37e4cfab4bfef6f363788d06ecb7757753ee5c0ba7ad0afcdL72-R82), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-9cbae292e39bc9b37e4cfab4bfef6f363788d06ecb7757753ee5c0ba7ad0afcdL136-R136), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f6098eab3d43899f2146cafc8454b2d9bae86064709eb2eeb43e0edbf459f22fL33-R33), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f6098eab3d43899f2146cafc8454b2d9bae86064709eb2eeb43e0edbf459f22fL55-R55), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f6098eab3d43899f2146cafc8454b2d9bae86064709eb2eeb43e0edbf459f22fL93-R98), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-86667a5be06e5995f4b15199aa3d4b655e64cceed6d108185d5cc71e03f0bb1fL43-R53), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-86667a5be06e5995f4b15199aa3d4b655e64cceed6d108185d5cc71e03f0bb1fL89-R94), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-86667a5be06e5995f4b15199aa3d4b655e64cceed6d108185d5cc71e03f0bb1fL108-R108), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL51-R51), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL61-R61), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL68-R68), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL78-R78), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL89-R89), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL116-R116), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f9f9a26653afb587b73d90d08c099a8236820fc72b4d247e001954d620073ebcL124-R124), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL66-R71), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL80-R80), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL86-R86), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL95-R95), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL105-R105), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL119-R119), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL125-R125), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL132-R132), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-52ca6a00561cb88b076a00d2ae2a30dca3d46713e52f0f5ca6b52ecfd95545daL344-R344), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-f2d9c4a8c504d2073549cde99b31a733a2f39bba4ce9f062cbaa09e1be17aa87L243-R243), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L11-R11), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L29-R29), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L49-R49), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L153-R153), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L185-R185), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L194-R194), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L207-R207), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L221-R221), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L242-R242), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L259-R259), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L281-R281), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L295-R295), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L310-R310), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L329-R329), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L352-R352), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L451-R451), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L469-R469), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L484-R484), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L501-R501), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L515-R515), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-d6f5c74e17462c8ff96d5bba06ebf81d16c015ca667fa945513a80be17bef017L533-R533), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-ca3c8d1edb25b6a7f4affbb79b2e3e74f73b3757e5d465258ce42ea9eb09fbc0L1035-R1035), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-ca3c8d1edb25b6a7f4affbb79b2e3e74f73b3757e5d465258ce42ea9eb09fbc0L1052-R1052), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-ca3c8d1edb25b6a7f4affbb79b2e3e74f73b3757e5d465258ce42ea9eb09fbc0L1134-R1134), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-ca3c8d1edb25b6a7f4affbb79b2e3e74f73b3757e5d465258ce42ea9eb09fbc0L1149-R1149), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-85fb062b09589d588166180e3c7954753018dc39474d3846346f2a8bb3faac32L480-R480), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-85fb062b09589d588166180e3c7954753018dc39474d3846346f2a8bb3faac32L494-R494), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-8641463ffef94529453efd945e02d056e4339417cc65aad074d4e6b3ef093244L49-R53), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-e206c19691b4678b401a5e2787000e93ab0b0060cbf4ac2625b51f87599d98d8L218-R218), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-e206c19691b4678b401a5e2787000e93ab0b0060cbf4ac2625b51f87599d98d8L238-R238), [link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-e206c19691b4678b401a5e2787000e93ab0b0060cbf4ac2625b51f87599d98d8L258-R258)) * Delete `python/taichi/experimental.py` file as it is no longer needed ([link](https://github.com/taichi-dev/taichi/pull/8399/files?diff=unified&w=0#diff-60267f25c1e1dfada9487235a4a793b63a49aa3d33c93b8449e24ac821a33174))
2023-10-31 14:11:53 +08:00
@ti.real_func
def sum(l: ti.i32, r: ti.i32) -> ti.i32:
if l == r:
return l
else:
return sum((l + r) // 2 + 1, r) + sum(l, (l + r) // 2)
@ti.kernel
def get_sum() -> ti.i32:
return sum(0, 99)
assert get_sum() == 99 * 50
assert added_files() == expected_num_cache_files()
def my_init():
ti.init(arch=curr_arch, enable_fallback=False, **{**current_thread_ext_options(), "cuda_stack_limit": 4096})
my_init()
helper1()
my_init()
assert added_files() == expected_num_cache_files(1)
helper1()
my_init()
assert added_files() == expected_num_cache_files(1)
helper2()
my_init()
assert added_files() == expected_num_cache_files(2)
helper2()
ti.reset()
assert added_files() == expected_num_cache_files(2)