From 8b125a6f18f41c3565d4c3d2238310854b9befe1 Mon Sep 17 00:00:00 2001 From: XiangyunYang Date: Fri, 4 Mar 2022 09:15:45 +0800 Subject: [PATCH 1/5] add error message --- python/taichi/lang/kernel_impl.py | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/python/taichi/lang/kernel_impl.py b/python/taichi/lang/kernel_impl.py index 2378295b61ddc..734cfbd406939 100644 --- a/python/taichi/lang/kernel_impl.py +++ b/python/taichi/lang/kernel_impl.py @@ -608,6 +608,20 @@ def func__(*args): if not self.is_grad and self.runtime.target_tape and not self.runtime.grad_replaced: self.runtime.target_tape.insert(self, args) + if actual_argument_slot > 8 and ( + _ti_core.arch_name(impl.current_cfg().arch) == 'opengl' + or _ti_core.arch_name(impl.current_cfg().arch) == 'cc'): + raise TaichiSyntaxError( + f"The number of elements in kernel arguments is too big! Do not exceed 8 on {_ti_core.arch_name(impl.current_cfg().arch)} backend." + ) + + if actual_argument_slot > 64 and ( + (_ti_core.arch_name(impl.current_cfg().arch) != 'opengl' + and _ti_core.arch_name(impl.current_cfg().arch) != 'cc')): + raise TaichiSyntaxError( + f"The number of elements in kernel arguments is too big! Do not exceed 64 on {_ti_core.arch_name(impl.current_cfg().arch)} backend." + ) + t_kernel(launch_ctx) ret = None From 18ae39e2486688a13799882b75f05ad254d9615c Mon Sep 17 00:00:00 2001 From: XiangyunYang Date: Fri, 4 Mar 2022 11:34:27 +0800 Subject: [PATCH 2/5] fix errortype --- python/taichi/lang/kernel_impl.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/taichi/lang/kernel_impl.py b/python/taichi/lang/kernel_impl.py index 734cfbd406939..5800e16cb202d 100644 --- a/python/taichi/lang/kernel_impl.py +++ b/python/taichi/lang/kernel_impl.py @@ -611,14 +611,14 @@ def func__(*args): if actual_argument_slot > 8 and ( _ti_core.arch_name(impl.current_cfg().arch) == 'opengl' or _ti_core.arch_name(impl.current_cfg().arch) == 'cc'): - raise TaichiSyntaxError( + raise RuntimeError( f"The number of elements in kernel arguments is too big! Do not exceed 8 on {_ti_core.arch_name(impl.current_cfg().arch)} backend." ) if actual_argument_slot > 64 and ( (_ti_core.arch_name(impl.current_cfg().arch) != 'opengl' and _ti_core.arch_name(impl.current_cfg().arch) != 'cc')): - raise TaichiSyntaxError( + raise RuntimeError( f"The number of elements in kernel arguments is too big! Do not exceed 64 on {_ti_core.arch_name(impl.current_cfg().arch)} backend." ) From f1817e989966865450385bd51e5a57705df47ff9 Mon Sep 17 00:00:00 2001 From: XiangyunYang Date: Fri, 4 Mar 2022 16:05:11 +0800 Subject: [PATCH 3/5] add test --- python/taichi/lang/kernel_impl.py | 6 ++--- tests/python/test_argument.py | 45 +++++++++++++++++++++++++++++++ 2 files changed, 48 insertions(+), 3 deletions(-) create mode 100644 tests/python/test_argument.py diff --git a/python/taichi/lang/kernel_impl.py b/python/taichi/lang/kernel_impl.py index 5800e16cb202d..99c5d1c9bc0c6 100644 --- a/python/taichi/lang/kernel_impl.py +++ b/python/taichi/lang/kernel_impl.py @@ -12,7 +12,7 @@ from taichi.lang.ast import (ASTTransformerContext, KernelSimplicityASTChecker, transform_tree) from taichi.lang.enums import Layout -from taichi.lang.exception import (TaichiCompilationError, +from taichi.lang.exception import (TaichiCompilationError, TaichiRuntimeError, TaichiRuntimeTypeError, TaichiSyntaxError) from taichi.lang.expr import Expr from taichi.lang.matrix import Matrix, MatrixType @@ -611,14 +611,14 @@ def func__(*args): if actual_argument_slot > 8 and ( _ti_core.arch_name(impl.current_cfg().arch) == 'opengl' or _ti_core.arch_name(impl.current_cfg().arch) == 'cc'): - raise RuntimeError( + raise TaichiRuntimeError( f"The number of elements in kernel arguments is too big! Do not exceed 8 on {_ti_core.arch_name(impl.current_cfg().arch)} backend." ) if actual_argument_slot > 64 and ( (_ti_core.arch_name(impl.current_cfg().arch) != 'opengl' and _ti_core.arch_name(impl.current_cfg().arch) != 'cc')): - raise RuntimeError( + raise TaichiRuntimeError( f"The number of elements in kernel arguments is too big! Do not exceed 64 on {_ti_core.arch_name(impl.current_cfg().arch)} backend." ) diff --git a/tests/python/test_argument.py b/tests/python/test_argument.py new file mode 100644 index 0000000000000..5faf406f53ee8 --- /dev/null +++ b/tests/python/test_argument.py @@ -0,0 +1,45 @@ +import pytest + +import taichi as ti +from tests import test_utils + + +@test_utils.test(arch=[ti.opengl, ti.cc]) +def test_exceed_max_eight(): + @ti.kernel + def foo1(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, + g: ti.i32, h: ti.i32) -> ti.i32: + return a + b + c + d + e + f + g + h + + assert foo1(1, 2, 3, 4, 5, 6, 7, 8) == 36 + + @ti.kernel + def foo2(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, + g: ti.i32, h: ti.i32, i: ti.i32) -> ti.i32: + return a + b + c + d + e + f + g + h + i + + with pytest.raises(ti.TaichiRuntimeError): + foo2(1, 2, 3, 4, 5, 6, 7, 8, 9) + + +@test_utils.test(exclude=[ti.opengl, ti.cc]) +def test_exceed_max_64(): + N = 64 + + @ti.kernel + def foo1(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + assert foo1(A) == 64 + + N = 65 + + @ti.kernel + def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + + with pytest.raises(ti.TaichiRuntimeError): + foo2(A) From 4b82fd2b0984aaebf3da3be86c63a432454f6b4a Mon Sep 17 00:00:00 2001 From: XiangyunYang Date: Fri, 4 Mar 2022 16:34:04 +0800 Subject: [PATCH 4/5] fix test --- tests/python/test_argument.py | 146 +++++++++++++++++++++++++++++++++- 1 file changed, 142 insertions(+), 4 deletions(-) diff --git a/tests/python/test_argument.py b/tests/python/test_argument.py index 5faf406f53ee8..40ed86e9db722 100644 --- a/tests/python/test_argument.py +++ b/tests/python/test_argument.py @@ -4,7 +4,7 @@ from tests import test_utils -@test_utils.test(arch=[ti.opengl, ti.cc]) +@test_utils.test(arch=ti.opengl) def test_exceed_max_eight(): @ti.kernel def foo1(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, @@ -18,11 +18,145 @@ def foo2(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, g: ti.i32, h: ti.i32, i: ti.i32) -> ti.i32: return a + b + c + d + e + f + g + h + i - with pytest.raises(ti.TaichiRuntimeError): + with pytest.raises( + ti.TaichiRuntimeError, + match= + "The number of elements in kernel arguments is too big! Do not exceed 8 on opengl backend." + ): foo2(1, 2, 3, 4, 5, 6, 7, 8, 9) -@test_utils.test(exclude=[ti.opengl, ti.cc]) +@test_utils.test(arch=ti.cc) +def test_exceed_max_eight(): + @ti.kernel + def foo1(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, + g: ti.i32, h: ti.i32) -> ti.i32: + return a + b + c + d + e + f + g + h + + assert foo1(1, 2, 3, 4, 5, 6, 7, 8) == 36 + + @ti.kernel + def foo2(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, + g: ti.i32, h: ti.i32, i: ti.i32) -> ti.i32: + return a + b + c + d + e + f + g + h + i + + with pytest.raises( + ti.TaichiRuntimeError, + match= + "The number of elements in kernel arguments is too big! Do not exceed 8 on cc backend." + ): + foo2(1, 2, 3, 4, 5, 6, 7, 8, 9) + + +@test_utils.test(arch=ti.cuda) +def test_exceed_max_64(): + N = 64 + + @ti.kernel + def foo1(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + assert foo1(A) == 64 + + N = 65 + + @ti.kernel + def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + + with pytest.raises( + ti.TaichiRuntimeError, + match= + "The number of elements in kernel arguments is too big! Do not exceed 64 on cuda backend." + ): + foo2(A) + + +@test_utils.test(arch=ti.metal) +def test_exceed_max_64(): + N = 64 + + @ti.kernel + def foo1(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + assert foo1(A) == 64 + + N = 65 + + @ti.kernel + def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + + with pytest.raises( + ti.TaichiRuntimeError, + match= + "The number of elements in kernel arguments is too big! Do not exceed 64 on metal backend." + ): + foo2(A) + + +@test_utils.test(arch=ti.vulkan) +def test_exceed_max_64(): + N = 64 + + @ti.kernel + def foo1(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + assert foo1(A) == 64 + + N = 65 + + @ti.kernel + def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + + with pytest.raises( + ti.TaichiRuntimeError, + match= + "The number of elements in kernel arguments is too big! Do not exceed 64 on vulkan backend." + ): + foo2(A) + + +@test_utils.test(arch=ti.x64) +def test_exceed_max_64(): + N = 64 + + @ti.kernel + def foo1(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + assert foo1(A) == 64 + + N = 65 + + @ti.kernel + def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: + return a.sum() + + A = ti.Vector([1] * N) + + with pytest.raises( + ti.TaichiRuntimeError, + match= + "The number of elements in kernel arguments is too big! Do not exceed 64 on x64 backend." + ): + foo2(A) + + +@test_utils.test(arch=ti.arm64) def test_exceed_max_64(): N = 64 @@ -41,5 +175,9 @@ def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: A = ti.Vector([1] * N) - with pytest.raises(ti.TaichiRuntimeError): + with pytest.raises( + ti.TaichiRuntimeError, + match= + "The number of elements in kernel arguments is too big! Do not exceed 64 on arm64 backend." + ): foo2(A) From d0bd5a4c5a26b0ce01e6f6c37959a764f6070034 Mon Sep 17 00:00:00 2001 From: XiangyunYang Date: Tue, 8 Mar 2022 15:27:17 +0800 Subject: [PATCH 5/5] fix test --- python/taichi/lang/kernel_impl.py | 8 +- tests/python/test_argument.py | 138 +----------------------------- 2 files changed, 8 insertions(+), 138 deletions(-) diff --git a/python/taichi/lang/kernel_impl.py b/python/taichi/lang/kernel_impl.py index 99c5d1c9bc0c6..100dde53c058e 100644 --- a/python/taichi/lang/kernel_impl.py +++ b/python/taichi/lang/kernel_impl.py @@ -609,15 +609,15 @@ def func__(*args): self.runtime.target_tape.insert(self, args) if actual_argument_slot > 8 and ( - _ti_core.arch_name(impl.current_cfg().arch) == 'opengl' - or _ti_core.arch_name(impl.current_cfg().arch) == 'cc'): + impl.current_cfg().arch == _ti_core.opengl + or impl.current_cfg().arch == _ti_core.cc): raise TaichiRuntimeError( f"The number of elements in kernel arguments is too big! Do not exceed 8 on {_ti_core.arch_name(impl.current_cfg().arch)} backend." ) if actual_argument_slot > 64 and ( - (_ti_core.arch_name(impl.current_cfg().arch) != 'opengl' - and _ti_core.arch_name(impl.current_cfg().arch) != 'cc')): + (impl.current_cfg().arch != _ti_core.opengl + and impl.current_cfg().arch != _ti_core.cc)): raise TaichiRuntimeError( f"The number of elements in kernel arguments is too big! Do not exceed 64 on {_ti_core.arch_name(impl.current_cfg().arch)} backend." ) diff --git a/tests/python/test_argument.py b/tests/python/test_argument.py index 40ed86e9db722..7fc38a7077b7d 100644 --- a/tests/python/test_argument.py +++ b/tests/python/test_argument.py @@ -4,7 +4,7 @@ from tests import test_utils -@test_utils.test(arch=ti.opengl) +@test_utils.test(arch=[ti.opengl, ti.cc]) def test_exceed_max_eight(): @ti.kernel def foo1(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, @@ -21,142 +21,12 @@ def foo2(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, with pytest.raises( ti.TaichiRuntimeError, match= - "The number of elements in kernel arguments is too big! Do not exceed 8 on opengl backend." + f"The number of elements in kernel arguments is too big! Do not exceed 8 on {ti.lang._ti_core.arch_name(ti.lang.impl.current_cfg().arch)} backend." ): foo2(1, 2, 3, 4, 5, 6, 7, 8, 9) -@test_utils.test(arch=ti.cc) -def test_exceed_max_eight(): - @ti.kernel - def foo1(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, - g: ti.i32, h: ti.i32) -> ti.i32: - return a + b + c + d + e + f + g + h - - assert foo1(1, 2, 3, 4, 5, 6, 7, 8) == 36 - - @ti.kernel - def foo2(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32, - g: ti.i32, h: ti.i32, i: ti.i32) -> ti.i32: - return a + b + c + d + e + f + g + h + i - - with pytest.raises( - ti.TaichiRuntimeError, - match= - "The number of elements in kernel arguments is too big! Do not exceed 8 on cc backend." - ): - foo2(1, 2, 3, 4, 5, 6, 7, 8, 9) - - -@test_utils.test(arch=ti.cuda) -def test_exceed_max_64(): - N = 64 - - @ti.kernel - def foo1(a: ti.types.vector(N, ti.i32)) -> ti.i32: - return a.sum() - - A = ti.Vector([1] * N) - assert foo1(A) == 64 - - N = 65 - - @ti.kernel - def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: - return a.sum() - - A = ti.Vector([1] * N) - - with pytest.raises( - ti.TaichiRuntimeError, - match= - "The number of elements in kernel arguments is too big! Do not exceed 64 on cuda backend." - ): - foo2(A) - - -@test_utils.test(arch=ti.metal) -def test_exceed_max_64(): - N = 64 - - @ti.kernel - def foo1(a: ti.types.vector(N, ti.i32)) -> ti.i32: - return a.sum() - - A = ti.Vector([1] * N) - assert foo1(A) == 64 - - N = 65 - - @ti.kernel - def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: - return a.sum() - - A = ti.Vector([1] * N) - - with pytest.raises( - ti.TaichiRuntimeError, - match= - "The number of elements in kernel arguments is too big! Do not exceed 64 on metal backend." - ): - foo2(A) - - -@test_utils.test(arch=ti.vulkan) -def test_exceed_max_64(): - N = 64 - - @ti.kernel - def foo1(a: ti.types.vector(N, ti.i32)) -> ti.i32: - return a.sum() - - A = ti.Vector([1] * N) - assert foo1(A) == 64 - - N = 65 - - @ti.kernel - def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: - return a.sum() - - A = ti.Vector([1] * N) - - with pytest.raises( - ti.TaichiRuntimeError, - match= - "The number of elements in kernel arguments is too big! Do not exceed 64 on vulkan backend." - ): - foo2(A) - - -@test_utils.test(arch=ti.x64) -def test_exceed_max_64(): - N = 64 - - @ti.kernel - def foo1(a: ti.types.vector(N, ti.i32)) -> ti.i32: - return a.sum() - - A = ti.Vector([1] * N) - assert foo1(A) == 64 - - N = 65 - - @ti.kernel - def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: - return a.sum() - - A = ti.Vector([1] * N) - - with pytest.raises( - ti.TaichiRuntimeError, - match= - "The number of elements in kernel arguments is too big! Do not exceed 64 on x64 backend." - ): - foo2(A) - - -@test_utils.test(arch=ti.arm64) +@test_utils.test(exclude=[ti.opengl, ti.cc]) def test_exceed_max_64(): N = 64 @@ -178,6 +48,6 @@ def foo2(a: ti.types.vector(N, ti.i32)) -> ti.i32: with pytest.raises( ti.TaichiRuntimeError, match= - "The number of elements in kernel arguments is too big! Do not exceed 64 on arm64 backend." + f"The number of elements in kernel arguments is too big! Do not exceed 64 on {ti.lang._ti_core.arch_name(ti.lang.impl.current_cfg().arch)} backend." ): foo2(A)