[IR] Continue statement behavior is confusing in struct for

[BillXu2000]

Describe the bug
(update) Same issue as #3271.

Detected by @g1n0st .
Continue statement in struct for scope will skip the whole lowest activated dense block.
I am not sure whether it is a feature or a bug.

To Reproduce

import taichi as ti

ti.init(arch=ti.cuda, kernel_profiler = False)

n = 128
x = ti.field(ti.f32)
block = ti.root.pointer(ti.i, n // 16)
block.dense(ti.i, 16).place(x)

@ti.kernel 
def test():
    x[0] = 1
    x[1] = 2
    cnt = 0
    for i in x:
        if x[i] == 1: continue
        if x[i] == 2: cnt += 1
    print(cnt) # output 0

test()

Log/Screenshots

{11:11}~ ➭ p test.py  
[Taichi] version 0.8.3, llvm 10.0.0, commit 04cf87de, linux, python 3.8.5
[Taichi] Starting on arch=x64
0

Additional comments

taichi/taichi/ir/statements.cpp

Line 7 in 3a9c1f3

bool ContinueStmt::as_return() const {

https://github.com/taichi-dev/taichi/blob/master/taichi/codegen/codegen_llvm.cpp#L802
A continue statement will emit a return statement in LLVM.

https://github.com/taichi-dev/taichi/blob/master/taichi/codegen/codegen_llvm.cpp#L1527
But every thread in struct for will loop over the whole lowest dense block.
I guess that a return statement in the kernel will force the thread to jump over the owest dense block loop, and resulting in this confusing skipping behavior.

[k-ye]

Closing #3271 because it's a duplicate

[k-ye]

skip the whole lowest activated dense block

Nice catch! This is a bug