Taichi: Kernel compilation failure when doing bitwise comparisons with sparse data structures

[KLozes]

Describe the bug
Kernels fail to compile when using sparse data structures and doing more than 1 bitwise comparison in an if statement. Things seem to work fine with dense data structures and other comparisons. I believe this bug is very recent.

Log/Screenshots

[Taichi] mode=development
[Taichi] preparing sandbox at /tmp/taichi-42b68ux7
[Taichi] sandbox prepared
[Taichi] version 0.5.10, cpu only, commit 62e24dd4, python 3.6.8
[T 04/03/20 09:17:19.009] [memory_pool.cpp:MemoryPool@14] Memory pool created. Default buffer size per allocator = 1024 MB
[T 04/03/20 09:17:19.009] [llvm_context.cpp:TaichiLLVMContext@38] Creating Taichi llvm context for arch: x64
[T 04/03/20 09:17:19.009] [llvm_context.cpp:TaichiLLVMContext@62] Taichi llvm context created.
[T 04/03/20 09:17:19.009] [program.cpp:Program@111] Program arch=x64
[T 04/03/20 09:17:19.010] [/home/klozes/Documents/software/taichi/python/taichi/lang/kernel.py:materialize@220] Materializing layout...
[T 04/03/20 09:17:19.013] [llvm_context.cpp:compile_runtime_bitcode@128] Compiling runtime module bitcode...
[T 04/03/20 09:17:19.243] [llvm_context.cpp:compile_runtime_bitcode@143] runtime module bitcode compiled.
[T 04/03/20 09:17:19.377] [unified_allocator.cpp:UnifiedAllocator@52] Allocating virtual address space of size 1024 MB
[T 04/03/20 09:17:19.377] [unified_allocator.cpp:UnifiedAllocator@61] Memory allocated. Allocation time = 1.72e-05 s
[T 04/03/20 09:17:19.377] [program.cpp:initialize_runtime_system@174] Allocating data structure of size 16384 B
[T 04/03/20 09:17:19.387] [program.cpp:initialize_runtime_system@183] LLVMRuntime initialized
[T 04/03/20 09:17:19.387] [program.cpp:initialize_runtime_system@185] LLVMRuntime pointer fetched
[D 04/03/20 09:17:19.397] [memory_pool.cpp:daemon@105] Processing memory alloc request 0
[D 04/03/20 09:17:19.397] [memory_pool.cpp:daemon@112]   Allocating memory 8232 B (alignment 4096B) 
[D 04/03/20 09:17:19.397] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f0065048000
[D 04/03/20 09:17:19.398] [memory_pool.cpp:daemon@105] Processing memory alloc request 1
[D 04/03/20 09:17:19.398] [memory_pool.cpp:daemon@112]   Allocating memory 8232 B (alignment 4096B) 
[D 04/03/20 09:17:19.398] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f006504b000
[D 04/03/20 09:17:19.399] [memory_pool.cpp:daemon@105] Processing memory alloc request 2
[D 04/03/20 09:17:19.399] [memory_pool.cpp:daemon@112]   Allocating memory 8232 B (alignment 4096B) 
[D 04/03/20 09:17:19.399] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f006504e000
[D 04/03/20 09:17:19.400] [memory_pool.cpp:daemon@105] Processing memory alloc request 3
[D 04/03/20 09:17:19.400] [memory_pool.cpp:daemon@112]   Allocating memory 8232 B (alignment 4096B) 
[D 04/03/20 09:17:19.400] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f0065051000
[D 04/03/20 09:17:19.401] [memory_pool.cpp:daemon@105] Processing memory alloc request 4
[D 04/03/20 09:17:19.401] [memory_pool.cpp:daemon@112]   Allocating memory 8232 B (alignment 4096B) 
[D 04/03/20 09:17:19.401] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f0065054000
[D 04/03/20 09:17:19.403] [memory_pool.cpp:daemon@105] Processing memory alloc request 5
[D 04/03/20 09:17:19.403] [memory_pool.cpp:daemon@112]   Allocating memory 3145728 B (alignment 4096B) 
[D 04/03/20 09:17:19.403] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f0065057000
[T 04/03/20 09:17:19.403] [program.cpp:initialize_runtime_system@211] Initializing allocator for snode 1 (node size 128)
[D 04/03/20 09:17:19.415] [memory_pool.cpp:daemon@105] Processing memory alloc request 6
[D 04/03/20 09:17:19.415] [memory_pool.cpp:daemon@112]   Allocating memory 56 B (alignment 4096B) 
[D 04/03/20 09:17:19.415] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f0065357000
[D 04/03/20 09:17:19.416] [memory_pool.cpp:daemon@105] Processing memory alloc request 7
[D 04/03/20 09:17:19.416] [memory_pool.cpp:daemon@112]   Allocating memory 8232 B (alignment 4096B) 
[D 04/03/20 09:17:19.416] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f0065358000
[D 04/03/20 09:17:19.417] [memory_pool.cpp:daemon@105] Processing memory alloc request 8
[D 04/03/20 09:17:19.417] [memory_pool.cpp:daemon@112]   Allocating memory 8232 B (alignment 4096B) 
[D 04/03/20 09:17:19.417] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f006535b000
[D 04/03/20 09:17:19.418] [memory_pool.cpp:daemon@105] Processing memory alloc request 9
[D 04/03/20 09:17:19.418] [memory_pool.cpp:daemon@112]   Allocating memory 8232 B (alignment 4096B) 
[D 04/03/20 09:17:19.418] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f006535e000
[T 04/03/20 09:17:19.418] [program.cpp:initialize_runtime_system@216] Allocating ambient element for snode 1 (node size 128)
[D 04/03/20 09:17:19.424] [memory_pool.cpp:daemon@105] Processing memory alloc request 10
[D 04/03/20 09:17:19.424] [memory_pool.cpp:daemon@112]   Allocating memory 2097152 B (alignment 4096B) 
[D 04/03/20 09:17:19.424] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f0065361000
[T 04/03/20 09:17:19.430] [program.cpp:materialize_layout@249] materialize_layout called
[T 04/03/20 09:17:19.431] [/home/klozes/Documents/software/taichi/python/taichi/lang/kernel.py:__call__@414] Compiling kernel func1_c4_0...
[D 04/03/20 09:17:19.552] [memory_pool.cpp:daemon@105] Processing memory alloc request 11
[D 04/03/20 09:17:19.552] [memory_pool.cpp:daemon@112]   Allocating memory 3145728 B (alignment 4096B) 
[D 04/03/20 09:17:19.552] [memory_pool.cpp:daemon@114]   Allocated. Ptr = 0x7f0065561000
[T 04/03/20 09:17:19.566] [/home/klozes/Documents/software/taichi/python/taichi/lang/kernel.py:__call__@414] Compiling kernel func2_c6_0...
[T 04/03/20 09:17:19.695] [/home/klozes/Documents/software/taichi/python/taichi/lang/kernel.py:__call__@414] Compiling kernel func3_c8_0...
python3: /home/klozes/Documents/software/llvm-8.0.1.src/lib/Transforms/Utils/Local.cpp:611: bool llvm::SimplifyInstructionsInBlock(llvm::BasicBlock*, const llvm::TargetLibraryInfo*): Assertion `!BI->isTerminator()' failed.
[E 04/03/20 09:17:19.791] Received signal 6 (Aborted)


***********************************
* Taichi Compiler Stack Traceback *
***********************************
/tmp/taichi-42b68ux7/taichi_core.so: taichi::Logger::error(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, bool)
/tmp/taichi-42b68ux7/taichi_core.so: taichi::signal_handler(int)
/lib/x86_64-linux-gnu/libc.so.6(+0x3ef20) [0x7f00cdd8ff20]
/lib/x86_64-linux-gnu/libc.so.6: gsignal
/lib/x86_64-linux-gnu/libc.so.6: abort
/lib/x86_64-linux-gnu/libc.so.6(+0x3039a) [0x7f00cdd8139a]
/lib/x86_64-linux-gnu/libc.so.6(+0x30412) [0x7f00cdd81412]
/tmp/taichi-42b68ux7/taichi_core.so: llvm::SimplifyInstructionsInBlock(llvm::BasicBlock*, llvm::TargetLibraryInfo const*)
/tmp/taichi-42b68ux7/taichi_core.so(+0x1c427cc) [0x7f00a90e37cc]
/tmp/taichi-42b68ux7/taichi_core.so(+0x1c44d7f) [0x7f00a90e5d7f]
/tmp/taichi-42b68ux7/taichi_core.so: llvm::FPPassManager::runOnFunction(llvm::Function&)
/tmp/taichi-42b68ux7/taichi_core.so(+0xf1fde9) [0x7f00a83c0de9]
/tmp/taichi-42b68ux7/taichi_core.so: llvm::legacy::PassManagerImpl::run(llvm::Module&)
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::JITSessionCPU::global_optimize_module_cpu(std::unique_ptr<llvm::Module, std::default_delete<llvm::Module> >&)
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::JITSessionCPU::add_module(std::unique_ptr<llvm::Module, std::default_delete<llvm::Module> >)
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::TaichiLLVMContext::add_module(std::unique_ptr<llvm::Module, std::default_delete<llvm::Module> >)
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::CodeGenLLVM::compile_module_to_executable()
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::CodeGenLLVM::gen()
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::CodeGenCPU::codegen()
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::KernelCodeGen::compile()
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::Program::compile(taichi::lang::Kernel&)
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::Kernel::compile()
/tmp/taichi-42b68ux7/taichi_core.so: taichi::lang::Kernel::operator()()
/tmp/taichi-42b68ux7/taichi_core.so(+0x72ba64) [0x7f00a7bcca64]
/tmp/taichi-42b68ux7/taichi_core.so(+0x65a3ad) [0x7f00a7afb3ad]
python3(_PyCFunction_FastCallDict+0x35c) [0x566fec]
python3() [0x595231]
python3() [0x54ac01]
python3(_PyObject_FastCallKeywords+0x19c) [0x5aa69c]
python3() [0x50ab53]
python3(_PyEval_EvalFrameDefault+0x449) [0x50c549]
python3() [0x5081d5]
python3() [0x58952b]
python3(PyObject_Call+0x3e) [0x5a04ce]
python3(_PyEval_EvalFrameDefault+0x17f5) [0x50d8f5]
python3() [0x5081d5]
python3(_PyFunction_FastCallDict+0x2e2) [0x5095d2]
python3() [0x5951c1]
python3() [0x54ac01]
python3(PyObject_Call+0x3e) [0x5a04ce]
python3(_PyEval_EvalFrameDefault+0x17f5) [0x50d8f5]
python3() [0x5081d5]
python3() [0x50a020]
python3() [0x50aa1d]
python3(_PyEval_EvalFrameDefault+0x449) [0x50c549]
python3() [0x5081d5]
python3(PyEval_EvalCode+0x23) [0x50b3a3]
python3() [0x635082]
python3(PyRun_FileExFlags+0x97) [0x635137]
python3(PyRun_SimpleFileExFlags+0x17f) [0x6388ef]
python3(Py_Main+0x591) [0x639491]
python3(main+0xe0) [0x4b0f60]
/lib/x86_64-linux-gnu/libc.so.6: __libc_start_main
python3(_start+0x2a) [0x5b2eaa]
[sudo] password for klozes: 

To Reproduce

import taichi as ti

ti.init(arch=ti.x64, debug = True)

x = ti.var(dt=ti.i32)
y = ti.var(dt=ti.i32)
ti.root.pointer(ti.ij, 128//4).dense(ti.ij, 4).place(x, y)

@ti.kernel
def func1():
    for i,j in x:
        if x[i,j]&2==2: # compiles fine
            y[i,j] = 1

@ti.kernel
def func2():
    for i,j in x:
        if x[i,j]==2 or x[i,j]==4: # compiles fine
            y[i,j] = 1

@ti.kernel
def func3():
    for i,j in x:
        if x[i,j]&2==2 or x[i,j]&4==4: # does not compile
            y[i,j] = 1


func1()
func2()
func3()

[archibate]

You may also want to try ti.init(print_ir=True). It will print useful information about the IR compiling process.

[archibate]

@yuanming-hu I think we can have better error message (potentially indicating line number in IR) for assertion failures like this? So that CHI (#689) could become more mature.

[yuanming-hu]

Thanks for reporting!

Final IR of func3:

==========
kernel {
  $0 = offloaded clear_list S1pointer
  $1 = offloaded listgen S1pointer
  $2 = offloaded clear_list S2dense
  $3 = offloaded listgen S2dense
  $4 = offloaded struct_for(S2dense) block_dim=0 {
    <i32 x1> $5 = loop index 0
    <i32 x1> $6 = loop index 1
    <gen*x1> $7 = get root
    <i32 x1> $8 = const [0]
    <gen*x1> $9 = [S0root][root]::lookup($7, $8) activate = false
    <gen*x1> $10 = get child [S0root->S1pointer] $9
    <i32 x1> $11 = bit_extract($5 + 0, 2~7)
    <i32 x1> $12 = bit_extract($6 + 0, 2~7)
    <i32 x1> $13 = const [1]
    <i32 x1> $14 = const [32]
    <i32 x1> $15 = mul $11 $14
    <i32 x1> $16 = add $12 $15
    <gen*x1> $17 = [S1pointer][pointer]::lookup($10, $16) activate = false
    <gen*x1> $18 = get child [S1pointer->S2dense] $17
    <i32 x1> $19 = bit_extract($5 + 0, 0~2)
    <i32 x1> $20 = bit_extract($6 + 0, 0~2)
    <i32 x1> $21 = const [4]
    <i32 x1> $22 = mul $19 $21
    <i32 x1> $23 = add $20 $22
    <gen*x1> $24 = [S2dense][dense]::lookup($18, $23) activate = false
    <i32*x1> $25 = get child [S2dense->S3place_i32] $24
    <i32 x1> $26 = global load $25
    <i32 x1> $27 = const [2]
    <i32 x1> $28 = bit_and $26 $27
    <i32 x1> $29 = cmp_eq $28 $27
    <i32 x1> $30 = bit_and $13 $29
    <i32 x1> $31 = alloca
    <i32 x1> $32 = alloca
    if $30 {
      <i32 x1> $34 = const [1]
      <i32 x1> $35 = loop index 0
      <i32 x1> $36 = loop index 1
      <gen*x1> $37 = get root
      <i32 x1> $38 = const [0]
      <gen*x1> $39 = [S0root][root]::lookup($37, $38) activate = false
      <gen*x1> $40 = get child [S0root->S1pointer] $39
      <i32 x1> $41 = bit_extract($35 + 0, 2~7)
      <i32 x1> $42 = bit_extract($36 + 0, 2~7)
      <i32 x1> $43 = const [32]
      <i32 x1> $44 = mul $41 $43
      <i32 x1> $45 = add $42 $44
      <gen*x1> $46 = [S1pointer][pointer]::lookup($40, $45) activate = false
      <gen*x1> $47 = get child [S1pointer->S2dense] $46
      <i32 x1> $48 = bit_extract($35 + 0, 0~2)
      <i32 x1> $49 = bit_extract($36 + 0, 0~2)
      <i32 x1> $50 = const [4]
      <i32 x1> $51 = mul $48 $50
      <i32 x1> $52 = add $49 $51
      <gen*x1> $53 = [S2dense][dense]::lookup($47, $52) activate = false
      <i32*x1> $54 = get child [S2dense->S4place_i32] $53
      <i32*x1> $55 : global store [$54 <- $34]
    }
  }
}

LLVM complains

Terminator found in the middle of a basic block!
label %after_loop

LLVM IR:

define internal void @loop_body(%struct.Context*, %struct.Element*, i32, i32) {
allocs:
  %4 = alloca i32
  %5 = alloca %struct.PhysicalCoordinates
  %6 = alloca %struct.PointerMeta
  %7 = alloca %struct.DenseMeta
  %8 = alloca i32
  %9 = alloca i32
  %10 = alloca %struct.PointerMeta
  %11 = alloca %struct.DenseMeta
  br label %entry

entry:                                            ; preds = %allocs
  br label %loop_body

loop_body:                                        ; preds = %entry
  store i32 %2, i32* %4
  br label %test

test:                                             ; preds = %loop_body_tail, %loop_body
  %12 = load i32, i32* %4
  %13 = icmp slt i32 %12, %3
  br i1 %13, label %loop_body1, label %after_loop

loop_body1:                                       ; preds = %test
  %14 = call %struct.PhysicalCoordinates* @Element_get_ptr_pcoord(%struct.Element* %1)
  %15 = load i32, i32* %4
  call void @S2_refine_coordinates(%struct.PhysicalCoordinates* %14, %struct.PhysicalCoordinates* %5, i32 %15)
  br i1 true, label %bound_guarded_loop_body, label %loop_body_tail

after_loop:                                       ; preds = %test
  ret void
  ret void

loop_body_tail:                                   ; preds = %after_if, %loop_body1
  %16 = load i32, i32* %4
  %17 = add i32 %16, 1
  store i32 %17, i32* %4
  br label %test

bound_guarded_loop_body:                          ; preds = %loop_body1
  %18 = getelementptr %struct.PhysicalCoordinates, %struct.PhysicalCoordinates* %5, i32 0, i32 0, i32 0
  %19 = load i32, i32* %18
  %20 = getelementptr %struct.PhysicalCoordinates, %struct.PhysicalCoordinates* %5, i32 0, i32 0, i32 1
  %21 = load i32, i32* %20
  %22 = call %struct.LLVMRuntime* @Context_get_runtime(%struct.Context* %0)
  %23 = call i8* @LLVMRuntime_get_root(%struct.LLVMRuntime* %22)
  %24 = bitcast i8* %23 to %S0_ch*
  %25 = getelementptr %S0_ch, %S0_ch* %24, i32 0
  %26 = bitcast %S0_ch* %25 to i8*
  %27 = call i8* @get_ch_S0_to_S1(i8* %26)
  %28 = bitcast i8* %27 to %0*
  %29 = add i32 %19, 0
  %30 = lshr i32 %29, 2
  %31 = and i32 %30, 31
  %32 = add i32 %21, 0
  %33 = lshr i32 %32, 2
  %34 = and i32 %33, 31
  %35 = mul i32 %31, 32
  %36 = add i32 %34, %35
  %37 = bitcast %struct.PointerMeta* %6 to %struct.StructMeta*
  call void @StructMeta_set_snode_id(%struct.StructMeta* %37, i32 1)
  call void @StructMeta_set_element_size(%struct.StructMeta* %37, i64 128)
  call void @StructMeta_set_max_num_elements(%struct.StructMeta* %37, i32 1024)
  call void @StructMeta_set_context(%struct.StructMeta* %37, %struct.Context* %0)
  call void @StructMeta_set_lookup_element(%struct.StructMeta* %37, i8* (i8*, i8*, i32)* @Pointer_lookup_element)
  call void @StructMeta_set_is_active(%struct.StructMeta* %37, i32 (i8*, i8*, i32)* @Pointer_is_active)
  call void @StructMeta_set_get_num_elements(%struct.StructMeta* %37, i32 (i8*, i8*)* @Pointer_get_num_elements)
  call void @StructMeta_set_from_parent_element(%struct.StructMeta* %37, i8* (i8*)* @get_ch_S0_to_S1)
  call void @StructMeta_set_refine_coordinates(%struct.StructMeta* %37, void (%struct.PhysicalCoordinates*, %struct.PhysicalCoordinates*, i32)* @S1_refine_coordinates)
  %38 = bitcast %struct.PointerMeta* %6 to i8*
  %39 = bitcast %0* %28 to i8*
  %40 = call i8* @Pointer_lookup_element(i8* %38, i8* %39, i32 %36)
  %41 = call i8* @get_ch_S1_to_S2(i8* %40)
  %42 = bitcast i8* %41 to [16 x %S2_ch]*
  %43 = add i32 %19, 0
  %44 = lshr i32 %43, 0
  %45 = and i32 %44, 3
  %46 = add i32 %21, 0
  %47 = lshr i32 %46, 0
  %48 = and i32 %47, 3
  %49 = mul i32 %45, 4
  %50 = add i32 %48, %49
  %51 = bitcast %struct.DenseMeta* %7 to %struct.StructMeta*
  call void @StructMeta_set_snode_id(%struct.StructMeta* %51, i32 2)
  call void @StructMeta_set_element_size(%struct.StructMeta* %51, i64 8)
  call void @StructMeta_set_max_num_elements(%struct.StructMeta* %51, i32 16)
  call void @StructMeta_set_context(%struct.StructMeta* %51, %struct.Context* %0)
  call void @StructMeta_set_lookup_element(%struct.StructMeta* %51, i8* (i8*, i8*, i32)* @Dense_lookup_element)
  call void @StructMeta_set_is_active(%struct.StructMeta* %51, i32 (i8*, i8*, i32)* @Dense_is_active)
  call void @StructMeta_set_get_num_elements(%struct.StructMeta* %51, i32 (i8*, i8*)* @Dense_get_num_elements)
  call void @StructMeta_set_from_parent_element(%struct.StructMeta* %51, i8* (i8*)* @get_ch_S1_to_S2)
  call void @StructMeta_set_refine_coordinates(%struct.StructMeta* %51, void (%struct.PhysicalCoordinates*, %struct.PhysicalCoordinates*, i32)* @S2_refine_coordinates)
  call void @DenseMeta_set_morton_dim(%struct.DenseMeta* %7, i32 0)
  %52 = bitcast %struct.DenseMeta* %7 to i8*
  %53 = bitcast [16 x %S2_ch]* %42 to i8*
  %54 = call i8* @Dense_lookup_element(i8* %52, i8* %53, i32 %50)
  %55 = call i8* @get_ch_S2_to_S3(i8* %54)
  %56 = bitcast i8* %55 to i32*
  %57 = load i32, i32* %56
  %58 = and i32 %57, 2
  %59 = icmp eq i32 %58, 2
  %60 = sext i1 %59 to i32
  %61 = and i32 1, %60
  store i32 0, i32* %8
  store i32 0, i32* %9
  %62 = icmp ne i32 %61, 0
  br i1 %62, label %true_block, label %false_block

true_block:                                       ; preds = %bound_guarded_loop_body
  %63 = getelementptr %struct.PhysicalCoordinates, %struct.PhysicalCoordinates* %5, i32 0, i32 0, i32 0
  %64 = load i32, i32* %63
  %65 = getelementptr %struct.PhysicalCoordinates, %struct.PhysicalCoordinates* %5, i32 0, i32 0, i32 1
  %66 = load i32, i32* %65
  %67 = call %struct.LLVMRuntime* @Context_get_runtime(%struct.Context* %0)
  %68 = call i8* @LLVMRuntime_get_root(%struct.LLVMRuntime* %67)
  %69 = bitcast i8* %68 to %S0_ch*
  %70 = getelementptr %S0_ch, %S0_ch* %69, i32 0
  %71 = bitcast %S0_ch* %70 to i8*
  %72 = call i8* @get_ch_S0_to_S1(i8* %71)
  %73 = bitcast i8* %72 to %0*
  %74 = add i32 %64, 0
  %75 = lshr i32 %74, 2
  %76 = and i32 %75, 31
  %77 = add i32 %66, 0
  %78 = lshr i32 %77, 2
  %79 = and i32 %78, 31
  %80 = mul i32 %76, 32
  %81 = add i32 %79, %80
  %82 = bitcast %struct.PointerMeta* %10 to %struct.StructMeta*
  call void @StructMeta_set_snode_id(%struct.StructMeta* %82, i32 1)
  call void @StructMeta_set_element_size(%struct.StructMeta* %82, i64 128)
  call void @StructMeta_set_max_num_elements(%struct.StructMeta* %82, i32 1024)
  call void @StructMeta_set_context(%struct.StructMeta* %82, %struct.Context* %0)
  call void @StructMeta_set_lookup_element(%struct.StructMeta* %82, i8* (i8*, i8*, i32)* @Pointer_lookup_element)
  call void @StructMeta_set_is_active(%struct.StructMeta* %82, i32 (i8*, i8*, i32)* @Pointer_is_active)
  call void @StructMeta_set_get_num_elements(%struct.StructMeta* %82, i32 (i8*, i8*)* @Pointer_get_num_elements)
  call void @StructMeta_set_from_parent_element(%struct.StructMeta* %82, i8* (i8*)* @get_ch_S0_to_S1)
  call void @StructMeta_set_refine_coordinates(%struct.StructMeta* %82, void (%struct.PhysicalCoordinates*, %struct.PhysicalCoordinates*, i32)* @S1_refine_coordinates)
  %83 = bitcast %struct.PointerMeta* %10 to i8*
  %84 = bitcast %0* %73 to i8*
  %85 = call i8* @Pointer_lookup_element(i8* %83, i8* %84, i32 %81)
  %86 = call i8* @get_ch_S1_to_S2(i8* %85)
  %87 = bitcast i8* %86 to [16 x %S2_ch]*
  %88 = add i32 %64, 0
  %89 = lshr i32 %88, 0
  %90 = and i32 %89, 3
  %91 = add i32 %66, 0
  %92 = lshr i32 %91, 0
  %93 = and i32 %92, 3
  %94 = mul i32 %90, 4
  %95 = add i32 %93, %94
  %96 = bitcast %struct.DenseMeta* %11 to %struct.StructMeta*
  call void @StructMeta_set_snode_id(%struct.StructMeta* %96, i32 2)
  call void @StructMeta_set_element_size(%struct.StructMeta* %96, i64 8)
  call void @StructMeta_set_max_num_elements(%struct.StructMeta* %96, i32 16)
  call void @StructMeta_set_context(%struct.StructMeta* %96, %struct.Context* %0)
  call void @StructMeta_set_lookup_element(%struct.StructMeta* %96, i8* (i8*, i8*, i32)* @Dense_lookup_element)
  call void @StructMeta_set_is_active(%struct.StructMeta* %96, i32 (i8*, i8*, i32)* @Dense_is_active)
  call void @StructMeta_set_get_num_elements(%struct.StructMeta* %96, i32 (i8*, i8*)* @Dense_get_num_elements)
  call void @StructMeta_set_from_parent_element(%struct.StructMeta* %96, i8* (i8*)* @get_ch_S1_to_S2)
  call void @StructMeta_set_refine_coordinates(%struct.StructMeta* %96, void (%struct.PhysicalCoordinates*, %struct.PhysicalCoordinates*, i32)* @S2_refine_coordinates)
  call void @DenseMeta_set_morton_dim(%struct.DenseMeta* %11, i32 0)
  %97 = bitcast %struct.DenseMeta* %11 to i8*
  %98 = bitcast [16 x %S2_ch]* %87 to i8*
  %99 = call i8* @Dense_lookup_element(i8* %97, i8* %98, i32 %95)
  %100 = call i8* @get_ch_S2_to_S4(i8* %99)
  %101 = bitcast i8* %100 to i32*
  store i32 1, i32* %101
  br label %after_if

false_block:                                      ; preds = %bound_guarded_loop_body
  br label %after_if

after_if:                                         ; preds = %false_block, %true_block
  br label %loop_body_tail
}

Note

after_loop:                                       ; preds = %test
  ret void
  ret void

[yuanming-hu]

Fixed in #707:
https://github.com/taichi-dev/taichi/pull/707/files#diff-8e075ba664009f2366f44ec5c170d122L1365