taichi-dev / taichi

Productive, portable, and performant GPU programming in Python.
https://taichi-lang.org
Apache License 2.0
25.37k stars 2.27k forks source link

RuntimeError: [offload.cpp:visit@731] struct_for cannot be nested inside a kernel, stmt=$14 #8489

Open Cleaf-y opened 6 months ago

Cleaf-y commented 6 months ago

Environment

macos 14.4 aarch64, [Taichi] version 1.6.0, llvm 15.0.7, commit f1c6fbbd, osx, python 3.11.8

Issue

When dealing with SNode, I encounterd a strange problem with the following simplified code:

import taichi as ti

ti.init(arch=ti.cpu)

data = ti.field(dtype=ti.u8)
block = ti.root.pointer(ti.ij, (50, 50))
block.place(data)
data[0, 1] = 1

@ti.kernel
def test():
    for _ in range(100000):
        for i, j in data:
            print(i, j)

test()

Then it gots the error:

[E 03/17/24 20:16:46.217 526947] [offload.cpp:visit@731] struct_for cannot be nested inside a kernel, stmt=$14 Traceback (most recent call last): File "/Users/shuyechen/LabWork/iFIM/cross_validation.py", line 18, in test() File "/Users/shu/miniconda3/envs/fmm4/lib/python3.11/site-packages/taichi/lang/kernel_impl.py", line 974, in wrapped return primal(*args, *kwargs) ^^^^^^^^^^^^^^^^^^^^^^^ File "/Users/shu/miniconda3/envs/fmm4/lib/python3.11/site-packages/taichi/lang/kernel_impl.py", line 906, in call return self.runtime.compiled_functions[key](args) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/Users/shuminiconda3/envs/fmm4/lib/python3.11/site-packages/taichi/lang/kernel_impl.py", line 817, in func__ raise e from None File "/Users/shu/miniconda3/envs/fmm4/lib/python3.11/site-packages/taichi/lang/kernel_impl.py", line 814, in func__ t_kernel(launch_ctx) RuntimeError: [offload.cpp:visit@731] struct_for cannot be nested inside a kernel, stmt=$14

Could you please help me check this? I initially planning to use the SNode pointer to respresent a list of active points which will be calculated in one iteration. I tried to implement a count_available points function(which goes through all points, checks wheter it is acitve and update the counter and finally return it, looks reaaly dumb) to be used as the while loop condition, but got the same error, and I tried this for loop, still in vain. Maybe I use this in the wrong way? Thanks for your time!

Ives0721 commented 6 months ago

In the the-struct-for-statement doc of TaiChi v1.6, it had mention that :

The struct for statement must be at the outermost scope of the kernel, and it cannot be terminated by a break statement even when it is run serially.

So the loop style written in your example code (i.e. A struct-for nested in a range-for) is not allowed in TaiChi v1.6.