Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Kernel compilation failure when doing bitwise comparisons with sparse data structures #704

Closed
KLozes opened this issue Apr 3, 2020 · 4 comments
Closed

Kernel compilation failure when doing bitwise comparisons with sparse data structures #704

KLozes opened this issue Apr 3, 2020 · 4 comments
Assignees
@yuanming-hu
Labels
bug We've confirmed that this is an BUG

Comments

@KLozes
Copy link
Collaborator

KLozes commented Apr 3, 2020

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()
@KLozes KLozes added the potential bug Something that looks like a bug but not yet confirmed label Apr 3, 2020
@archibate
Copy link
Collaborator

archibate commented Apr 4, 2020 •
edited
Loading

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

@archibate
Copy link
Collaborator

archibate commented Apr 4, 2020 •
edited
Loading

@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
Copy link
Member

yuanming-hu commented Apr 5, 2020 •
edited
Loading

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 added a commit to yuanming-hu/taichi that referenced this issue Apr 5, 2020
@yuanming-hu
[LLVM] Fix LLVM struct-for codegen crashing due to extra return taich…
03c9cdf 
…i-dev#704
@yuanming-hu yuanming-hu mentioned this issue Apr 5, 2020
Merged
@yuanming-hu
Copy link
Member

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

@yuanming-hu yuanming-hu added bug We've confirmed that this is an BUG and removed potential bug Something that looks like a bug but not yet confirmed labels Apr 5, 2020
yuanming-hu added a commit that referenced this issue Apr 5, 2020
@yuanming-hu
[LLVM] Fix LLVM struct-for codegen crashing due to extra return #704 (#…
fe153b4 
…707)

* [LLVM] Fix LLVM struct-for codegen crashing due to extra return #704

* update
@KLozes KLozes closed this as completed Apr 6, 2020
@yuanming-hu yuanming-hu mentioned this issue Nov 25, 2020
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@yuanming-hu yuanming-hu

Labels
bug We've confirmed that this is an BUG
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
3 participants
@yuanming-hu @archibate @KLozes