ti.solve with singular A silently break control flow on GPU with debug=False

[YouJiacheng]

import taichi as ti

ti.init(arch=ti.gpu)

bar = ti.Matrix.field(3, 3, ti.float32, shape=16)
zoo = ti.Vector.field(3, ti.float32, shape=16)

bar.fill(0)
zoo.fill(1)

bar[0] = ti.Matrix([[1, 0, 0], [0, 1, 0], [0, 0, 1]])

@ti.kernel
def foo():
    for i in bar:
        res = ti.solve(bar[i], zoo[i], ti.float32)
        print(i)
foo()

The output is

[Taichi] version 1.2.1, llvm 10.0.0, commit 12ab828a, win, python 3.10.6
[Taichi] Starting on arch=cuda
[W 11/15/22 20:39:33.769 20600] [C:/Users/Administrator/actions-runner/_work/taichi/taichi/taichi/runtime/llvm/llvm_offline_cache.cpp:load_meta_data@138] Lock C:/taichi_cache/ticache/llvm/metadata.lock failed
0
[W 11/15/22 20:39:36.234 20600] [taichi/util/offline_cache.h:run@185] Lock C:/taichi_cache/ticache/llvm/metadata.lock failed
[W 11/15/22 20:39:36.482 20600] [C:/Users/Administrator/actions-runner/_work/taichi/taichi/taichi/runtime/llvm/llvm_offline_cache.cpp:dump@316] Lock C:/taichi_cache/ticache/llvm/metadata.lock failed

print(i) is executed only if A is not singular.
This is far out of my expectation, and completely different from CPU.
According to the document:

taichi should ignore assertion and produce non-sense result(e.g. NaN), instead of breaking control flow.

With debug=True, I can get nice error message.
But all 16 print(i) are executed on CPU, while only 1 print(i)(i=0) is executed on GPU.