From 92450961f3258cab35b7907fa046d2c4fdaf5b60 Mon Sep 17 00:00:00 2001 From: Yi Xu Date: Thu, 27 Oct 2022 17:31:04 +0800 Subject: [PATCH 1/3] [bug] Fix false overflow alarm in struct fors under packed mode --- taichi/ir/snode.cpp | 4 ++-- taichi/transforms/demote_dense_struct_fors.cpp | 2 +- tests/python/test_struct_for_non_pot.py | 15 +++++++++++++++ 3 files changed, 18 insertions(+), 3 deletions(-) diff --git a/taichi/ir/snode.cpp b/taichi/ir/snode.cpp index 6a4fe1f3890f0..0cb333fd22039 100644 --- a/taichi/ir/snode.cpp +++ b/taichi/ir/snode.cpp @@ -91,8 +91,8 @@ SNode &SNode::create_node(std::vector axes, } if (acc_shape > std::numeric_limits::max()) { TI_WARN( - "Snode index might be out of int32 boundary but int64 indexing is not " - "supported yet."); + "SNode index might be out of int32 boundary but int64 indexing is not " + "supported yet. Struct fors might not work either."); } new_node.num_cells_per_container = acc_shape; // infer extractors (only for POT) diff --git a/taichi/transforms/demote_dense_struct_fors.cpp b/taichi/transforms/demote_dense_struct_fors.cpp index 55ba8c780c7f7..192aa7cf52742 100644 --- a/taichi/transforms/demote_dense_struct_fors.cpp +++ b/taichi/transforms/demote_dense_struct_fors.cpp @@ -26,7 +26,6 @@ void convert_to_range_for(OffloadedStmt *offloaded, bool packed) { snode = snode->parent; } std::reverse(snodes.begin(), snodes.end()); - TI_ASSERT(total_bits <= 30); // general shape calculation - no dependence on POT int64 total_n = 1; @@ -38,6 +37,7 @@ void convert_to_range_for(OffloadedStmt *offloaded, bool packed) { } total_n *= s->num_cells_per_container; } + TI_ASSERT(total_n <= std::numeric_limits::max()); offloaded->const_begin = true; offloaded->const_end = true; diff --git a/tests/python/test_struct_for_non_pot.py b/tests/python/test_struct_for_non_pot.py index c8eefbe3defb5..bb1a712b9bf87 100644 --- a/tests/python/test_struct_for_non_pot.py +++ b/tests/python/test_struct_for_non_pot.py @@ -66,3 +66,18 @@ def test_2d(): @test_utils.test(require=ti.extension.packed, packed=True) def test_2d_packed(): _test_2d() + + +@test_utils.test(packed=True) +def test_2d_overflow_if_not_packed(): + n, m = 40000000, 50 + arr = ti.field(ti.u8, (n, m)) + + @ti.kernel + def count() -> ti.i32: + res = 0 + for _ in ti.grouped(arr): + res += 1 + return res + + assert count() == n * m From e891f9ea9a70bc8f121fbe7821acc8bc4c8a6c15 Mon Sep 17 00:00:00 2001 From: Yi Xu Date: Fri, 28 Oct 2022 13:10:27 +0800 Subject: [PATCH 2/3] Fix tests --- tests/python/test_struct_for_non_pot.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/python/test_struct_for_non_pot.py b/tests/python/test_struct_for_non_pot.py index bb1a712b9bf87..7e865777b8ce1 100644 --- a/tests/python/test_struct_for_non_pot.py +++ b/tests/python/test_struct_for_non_pot.py @@ -68,9 +68,9 @@ def test_2d_packed(): _test_2d() -@test_utils.test(packed=True) +@test_utils.test(require=ti.extension.packed, packed=True) def test_2d_overflow_if_not_packed(): - n, m = 40000000, 50 + n, m = 2**25 + 1, 2**5 + 1 arr = ti.field(ti.u8, (n, m)) @ti.kernel From 7605bd1944a943c2a86cdc1cfc4afe17c133e16d Mon Sep 17 00:00:00 2001 From: Yi Xu Date: Fri, 28 Oct 2022 13:45:14 +0800 Subject: [PATCH 3/3] Fix tests --- tests/python/test_struct_for_non_pot.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/python/test_struct_for_non_pot.py b/tests/python/test_struct_for_non_pot.py index 7e865777b8ce1..dd35434dfe517 100644 --- a/tests/python/test_struct_for_non_pot.py +++ b/tests/python/test_struct_for_non_pot.py @@ -70,8 +70,8 @@ def test_2d_packed(): @test_utils.test(require=ti.extension.packed, packed=True) def test_2d_overflow_if_not_packed(): - n, m = 2**25 + 1, 2**5 + 1 - arr = ti.field(ti.u8, (n, m)) + n, m, p = 2**9 + 1, 2**9 + 1, 2**10 + 1 + arr = ti.field(ti.u8, (n, m, p)) @ti.kernel def count() -> ti.i32: @@ -80,4 +80,4 @@ def count() -> ti.i32: res += 1 return res - assert count() == n * m + assert count() == n * m * p