From b075fef57dce266cbe26008645293a32082ffa49 Mon Sep 17 00:00:00 2001 From: Lin Jiang Date: Thu, 16 Nov 2023 15:24:31 +0800 Subject: [PATCH 1/6] [Doc] Remove the limit on kernel argument size --- docs/lang/articles/kernels/kernel_function.md | 34 +++++++++++++------ 1 file changed, 24 insertions(+), 10 deletions(-) diff --git a/docs/lang/articles/kernels/kernel_function.md b/docs/lang/articles/kernels/kernel_function.md index 3bf2ff39f7023..3cd3875ab4e3c 100644 --- a/docs/lang/articles/kernels/kernel_function.md +++ b/docs/lang/articles/kernels/kernel_function.md @@ -154,6 +154,16 @@ You can also use argument packs if you want to pass many arguments to a kernel. When defining the arguments of a kernel in Taichi, please make sure that each of the arguments has type hint. +:::caution WARNING + +We have removed the limit on the size of the argument in Taichi v1.7.0. +However, please keep in mind that the size of arguments in a kernel should be small. +When you pass a large argument to a kernel, the compile time will increase significantly. +If you find yourself passing a large argument to a kernel, you may want to consider using a `ti.field()` or a `ti.types.ndarray()` instead. + +We have not tested arguments with a very large size (>4KB), and we do not guarantee that it will work properly. +::: + ### Return value In Taichi, a kernel can have multiple return values, and the return values can either be a scalar, `ti.types.matrix()`, or `ti.types.vector()`. @@ -188,6 +198,16 @@ When defining the return value of a kernel in Taichi, it is important to follow - Use type hint to specify the return value of a kernel. - Make sure that you have at most one return statement in a kernel. +:::caution WARNING + +We have removed the limit on the size of the return values in Taichi v1.7.0. +However, please keep in mind that the size of return values in a kernel should be small. +When you pass a large argument to a kernel, the compile time will increase significantly. +If you find your return value is very large, you may want to consider using a `ti.field()` or a `ti.types.ndarray()` instead. + +We have not tested return values with a very large size (>4KB), and we do not guarantee that it will work properly. +::: + #### Automatic type cast In the following code snippet, the return value is automatically cast into the hinted type: @@ -264,11 +284,8 @@ All Taichi inline functions are force-inlined. This means that if you call a Tai ### Arguments -A Taichi inline function can accept multiple arguments, which may include scalar, `ti.types.matrix()`, `ti.types.vector()`, `ti.types.struct()`, `ti.types.ndarray()`, `ti.field()`, and `ti.template()` types. Note that some of the restrictions on kernel arguments do not apply to Taichi functions: - -- It is not strictly required to type hint the function arguments (but it is still recommended). -- You can pass an unlimited number of elements in the function arguments. - +A Taichi inline function can accept multiple arguments, which may include scalar, `ti.types.matrix()`, `ti.types.vector()`, `ti.types.struct()`, `ti.types.ndarray()`, `ti.field()`, and `ti.template()` types. +Note that unlike Taichi kernels, it is not strictly required to type hint the function arguments (but it is still recommended). ### Return values @@ -281,10 +298,8 @@ Return values of a Taichi inline function can be scalars, `ti.types.matrix()`, ` ### Arguments -A Taichi real function can accept multiple arguments, which may include scalar, `ti.types.matrix()`, `ti.types.vector()`, `ti.types.struct()`, `ti.types.ndarray()`, `ti.field()`, and `ti.template()` types. Note the following: - -- You must type hint the function arguments. -- You can pass an unlimited number of elements in the function arguments. +A Taichi real function can accept multiple arguments, which may include scalar, `ti.types.matrix()`, `ti.types.vector()`, `ti.types.struct()`, `ti.types.ndarray()`, `ti.field()`, and `ti.template()` types. +Note that you must type hint the function arguments. ### Return values @@ -302,7 +317,6 @@ Return values of a Taichi inline function can be scalars, `ti.types.matrix()`, ` | Type hint arguments | Mandatory | Recommended | Mandatory | | Type hint return values | Mandatory | Recommended | Mandatory | | Return type | | | | -| Maximum number of elements in arguments | | Unlimited | Unlimited | | Maximum number of return statements | 1 | 1 | Unlimited | From bc023dac613f30f53810761fbadc2b8547129243 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Thu, 16 Nov 2023 07:25:55 +0000 Subject: [PATCH 2/6] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- docs/lang/articles/kernels/kernel_function.md | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/docs/lang/articles/kernels/kernel_function.md b/docs/lang/articles/kernels/kernel_function.md index 3cd3875ab4e3c..a07db9b2832d4 100644 --- a/docs/lang/articles/kernels/kernel_function.md +++ b/docs/lang/articles/kernels/kernel_function.md @@ -156,8 +156,8 @@ When defining the arguments of a kernel in Taichi, please make sure that each of :::caution WARNING -We have removed the limit on the size of the argument in Taichi v1.7.0. -However, please keep in mind that the size of arguments in a kernel should be small. +We have removed the limit on the size of the argument in Taichi v1.7.0. +However, please keep in mind that the size of arguments in a kernel should be small. When you pass a large argument to a kernel, the compile time will increase significantly. If you find yourself passing a large argument to a kernel, you may want to consider using a `ti.field()` or a `ti.types.ndarray()` instead. @@ -284,7 +284,7 @@ All Taichi inline functions are force-inlined. This means that if you call a Tai ### Arguments -A Taichi inline function can accept multiple arguments, which may include scalar, `ti.types.matrix()`, `ti.types.vector()`, `ti.types.struct()`, `ti.types.ndarray()`, `ti.field()`, and `ti.template()` types. +A Taichi inline function can accept multiple arguments, which may include scalar, `ti.types.matrix()`, `ti.types.vector()`, `ti.types.struct()`, `ti.types.ndarray()`, `ti.field()`, and `ti.template()` types. Note that unlike Taichi kernels, it is not strictly required to type hint the function arguments (but it is still recommended). ### Return values @@ -298,7 +298,7 @@ Return values of a Taichi inline function can be scalars, `ti.types.matrix()`, ` ### Arguments -A Taichi real function can accept multiple arguments, which may include scalar, `ti.types.matrix()`, `ti.types.vector()`, `ti.types.struct()`, `ti.types.ndarray()`, `ti.field()`, and `ti.template()` types. +A Taichi real function can accept multiple arguments, which may include scalar, `ti.types.matrix()`, `ti.types.vector()`, `ti.types.struct()`, `ti.types.ndarray()`, `ti.field()`, and `ti.template()` types. Note that you must type hint the function arguments. ### Return values From e9a500ac2982ba9fa5910ea5ec8e164ef40d3b19 Mon Sep 17 00:00:00 2001 From: Lin Jiang Date: Thu, 16 Nov 2023 15:29:25 +0800 Subject: [PATCH 3/6] fix --- docs/lang/articles/kernels/kernel_function.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/lang/articles/kernels/kernel_function.md b/docs/lang/articles/kernels/kernel_function.md index 3cd3875ab4e3c..825a888adaec3 100644 --- a/docs/lang/articles/kernels/kernel_function.md +++ b/docs/lang/articles/kernels/kernel_function.md @@ -202,7 +202,7 @@ When defining the return value of a kernel in Taichi, it is important to follow We have removed the limit on the size of the return values in Taichi v1.7.0. However, please keep in mind that the size of return values in a kernel should be small. -When you pass a large argument to a kernel, the compile time will increase significantly. +When the return value of the kernel is very large, the compile time will increase significantly. If you find your return value is very large, you may want to consider using a `ti.field()` or a `ti.types.ndarray()` instead. We have not tested return values with a very large size (>4KB), and we do not guarantee that it will work properly. From 53c9044e38f921b2052ad28d7fbe41c19375720a Mon Sep 17 00:00:00 2001 From: Lin Jiang Date: Thu, 16 Nov 2023 15:40:22 +0800 Subject: [PATCH 4/6] add test --- tests/python/test_argument.py | 16 ++++++++++++++++ tests/python/test_return.py | 16 ++++++++++++++++ 2 files changed, 32 insertions(+) diff --git a/tests/python/test_argument.py b/tests/python/test_argument.py index a957765375fc1..4089a4d6b75a2 100644 --- a/tests/python/test_argument.py +++ b/tests/python/test_argument.py @@ -274,3 +274,19 @@ def foo(a: ti.f32) -> ti.f32: return bar(a) assert foo(1.5) == 1.0 + + +@test_utils.test() +def test_arg_4k(): + vec1024 = ti.types.vector(1024, ti.i32) + + @ti.kernel + def bar(a: vec1024)->ti.i32: + ret = 0 + for i in range(1024): + ret += a[i] + + return ret + + a = vec1024([i for i in range(1024)]) + assert bar(a) == 523776 diff --git a/tests/python/test_return.py b/tests/python/test_return.py index 9ea1df4b02821..0affd8230a258 100644 --- a/tests/python/test_return.py +++ b/tests/python/test_return.py @@ -343,3 +343,19 @@ def foo() -> tp: return bar() assert foo().a == 0 + + +@test_utils.test() +def test_ret_4k(): + vec1024 = ti.types.vector(1024, ti.i32) + + @ti.kernel + def foo()->vec1024: + ret = vec1024(0) + for i in range(1024): + ret[i] = i + return ret + + ret = foo() + for i in range(1024): + assert ret[i] == i From 77067001fb1e1a4209d164cd4b6718193f9a1286 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Thu, 16 Nov 2023 07:41:29 +0000 Subject: [PATCH 5/6] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- tests/python/test_argument.py | 2 +- tests/python/test_return.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/python/test_argument.py b/tests/python/test_argument.py index 4089a4d6b75a2..884940ee251a7 100644 --- a/tests/python/test_argument.py +++ b/tests/python/test_argument.py @@ -281,7 +281,7 @@ def test_arg_4k(): vec1024 = ti.types.vector(1024, ti.i32) @ti.kernel - def bar(a: vec1024)->ti.i32: + def bar(a: vec1024) -> ti.i32: ret = 0 for i in range(1024): ret += a[i] diff --git a/tests/python/test_return.py b/tests/python/test_return.py index 0affd8230a258..b40e64277be6a 100644 --- a/tests/python/test_return.py +++ b/tests/python/test_return.py @@ -350,7 +350,7 @@ def test_ret_4k(): vec1024 = ti.types.vector(1024, ti.i32) @ti.kernel - def foo()->vec1024: + def foo() -> vec1024: ret = vec1024(0) for i in range(1024): ret[i] = i From b2cb76666a7e71b2b4d9b988203fb0a00d57896c Mon Sep 17 00:00:00 2001 From: Lin Jiang Date: Thu, 16 Nov 2023 15:57:29 +0800 Subject: [PATCH 6/6] exclude amd --- tests/python/test_argument.py | 2 +- tests/python/test_return.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/python/test_argument.py b/tests/python/test_argument.py index 884940ee251a7..6c7bf5a77eef7 100644 --- a/tests/python/test_argument.py +++ b/tests/python/test_argument.py @@ -276,7 +276,7 @@ def foo(a: ti.f32) -> ti.f32: assert foo(1.5) == 1.0 -@test_utils.test() +@test_utils.test(exclude=[ti.amdgpu]) def test_arg_4k(): vec1024 = ti.types.vector(1024, ti.i32) diff --git a/tests/python/test_return.py b/tests/python/test_return.py index b40e64277be6a..69e58ce9ed92b 100644 --- a/tests/python/test_return.py +++ b/tests/python/test_return.py @@ -345,7 +345,7 @@ def foo() -> tp: assert foo().a == 0 -@test_utils.test() +@test_utils.test(exclude=[ti.amdgpu]) def test_ret_4k(): vec1024 = ti.types.vector(1024, ti.i32)