search

taichi-dev / taichi

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

Failed to export to C code. Any method to export sparse data structure to C? #2355

Open xuhao1 opened 3 years ago

xuhao1 commented 3 years ago

Describe the bug I am trying to export the taichi kernels to C code because my own taichi project TaichiSLAM requires massive real-time IO. I followed the instructions and tried to record my kernel, it raise

[I 05/20/21 14:57:28.509] [action_recorder.cpp:start_recording@27] ActionRecorder: start recording to [./export/TaichiSLAM.yml]
[Taichi] Starting on arch=cc
[Tina] Hint: MMB to orbit, Shift+MMB to pan, wheel to zoom
[Taichi] materializing...
[E 05/20/21 14:57:30.417] [struct_cc.cpp:generate_types@36] SNodeType=pointer not supported on C backend

Sparse data structure is improtant to my SLAM project (that's why I trying to write them in Taichi), and the IO speed of python too slow for a real-world aerial robot.

Does Taichi has the plan to support SNodeType=pointer on C backend or I have another way to call Taichi kernels in C/C++?

To Reproduce

git clone https://github.com/xuhao1/TaichiSLAM
python examples/TaichiSLAM_demo.py --record
# My data structure
  B = ti.root.pointer(ti.ijk, (block_num_xy, block_num_xy, block_num_z))
  B = B.dense(ti.ijk, (block_size, block_size, block_size))
  B.place(self.occupy, self.W_TSDF,self.TSDF, self.ESDF)
....
# Record to yml
if args.record:
    ti.core.start_recording('./export/TaichiSLAM.yml')
    ti.init(arch=ti.cc)
else:
    if args.cuda:
        ti.init(arch=ti.cuda)
    else:
        ti.init(arch=ti.cpu)

Log/Screenshots

(base) ➜  TaichiSLAM git:(main) ✗ python examples/TaichiSLAM_demo.py --record          
[Taichi] mode=release
[Taichi] preparing sandbox at /tmp/taichi-8bhk1fyu
[Taichi] version 0.7.19, llvm 10.0.0, commit 19356808, linux, python 3.8.5
Failed to load Python extension for LZ4 support. LZ4 compression will not be available.
[Tina] version 0.1.1
[Tina] Taichi properties hacked

Res [1024x768] GPU False RVIZ False scale of map [100, 10] grid 0.05 
[I 05/20/21 15:05:34.745] [action_recorder.cpp:start_recording@27] ActionRecorder: start recording to [./export/TaichiSLAM.yml]
[Taichi] Starting on arch=cc
The map voxel is:[2048x2048x256] all 1.02e+03M grid scale [0.049x0.049x0.039] map scale:[100mx100mx10m] tree depth [11, 8]
[Tina] Hint: MMB to orbit, Shift+MMB to pan, wheel to zoom
No data input, using random generate maps
[Taichi] materializing...
[E 05/20/21 15:05:36.469] [struct_cc.cpp:generate_types@36] SNodeType=pointer not supported on C backend

***********************************
* Taichi Compiler Stack Traceback *
***********************************
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so: taichi::Logger::error(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, bool)
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so: taichi::lang::cccp::CCLayoutGen::generate_types(taichi::lang::SNode*)
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so: taichi::lang::cccp::CCLayoutGen::generate_children(taichi::lang::SNode*)
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so: taichi::lang::cccp::CCLayoutGen::generate_types(taichi::lang::SNode*)
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so: taichi::lang::cccp::CCLayoutGen::compile()
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so: taichi::lang::cccp::CCProgram::compile_layout(taichi::lang::SNode*)
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so: taichi::lang::Program::materialize_layout()
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so: taichi::lang::layout(std::function<void ()> const&)
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so(+0x40a1e9) [0x7fe4818431e9]
/home/xuhao/anaconda3/lib/python3.8/site-packages/taichi/core/../lib/taichi_core.so(+0x35fda4) [0x7fe481798da4]
python(PyCFunction_Call+0x56) [0x56471736cf76]
python(_PyObject_MakeTpCall+0x22f) [0x56471732a85f]
python(_PyEval_EvalFrameDefault+0x4596) [0x5647173b1f56]
python(_PyEval_EvalCodeWithName+0x7df) [0x564717377f9f]
python(_PyFunction_Vectorcall+0x1e3) [0x564717378943]
python(+0x10075e) [0x5647172ed75e]
python(_PyEval_EvalCodeWithName+0x7df) [0x564717377f9f]
python(+0x18bd20) [0x564717378d20]
python(+0x10011a) [0x5647172ed11a]
python(_PyEval_EvalCodeWithName+0x2d2) [0x564717377a92]
python(_PyFunction_Vectorcall+0x1e3) [0x564717378943]
python(_PyObject_FastCallDict+0x24b) [0x5647173794cb]
python(_PyObject_Call_Prepend+0x63) [0x564717379733]
python(+0x18c83a) [0x56471737983a]
python(PyObject_Call+0x70) [0x56471732a200]
python(_PyEval_EvalFrameDefault+0x1fdb) [0x5647173af99b]
python(_PyEval_EvalCodeWithName+0x2d2) [0x564717377a92]
python(_PyObject_FastCallDict+0x20c) [0x56471737948c]
python(_PyObject_Call_Prepend+0x63) [0x564717379733]
python(+0x18c83a) [0x56471737983a]
python(_PyObject_MakeTpCall+0x22f) [0x56471732a85f]
python(_PyEval_EvalFrameDefault+0x4596) [0x5647173b1f56]
python(_PyEval_EvalCodeWithName+0x2d2) [0x564717377a92]
python(PyEval_EvalCodeEx+0x44) [0x564717378754]
python(PyEval_EvalCode+0x1c) [0x564717406edc]
python(+0x219f84) [0x564717406f84]
python(+0x24c1f4) [0x5647174391f4]
python(PyRun_FileExFlags+0xa1) [0x5647173016e1]
python(PyRun_SimpleFileExFlags+0x3b4) [0x564717301ac6]
python(+0x11598b) [0x56471730298b]
python(Py_BytesMain+0x39) [0x56471743bd19]
/lib/x86_64-linux-gnu/libc.so.6: __libc_start_main
python(+0x1dee93) [0x5647173cbe93]

Internal error occurred. Check out this page for possible solutions:
https://taichi.readthedocs.io/en/stable/install.html#troubleshooting
k-ye commented 3 years ago

Thanks! Supporting pointer on the CC backend is not on our near roadmap yet. We can try to figure sth out regarding how users can better utilize the LLVM backend's compiled results.

xuhao1 commented 3 years ago

@k-ye Thanks for your reply! Directly using LLVM's compiled result is more reasonable than use C as a intermediate representation . Do you have any roadmap on this?

k-ye commented 3 years ago

Hey @xuhao1 !

Basically, you can dump the .ll file using print_kernel_llvm_ir and print_kernel_llvm_ir_optimized (see also https://taichi.readthedocs.io/en/latest/contributor_guide.html#tips-on-the-taichi-compiler-development), then convert the ll to an actual library via LLVM CLI tools.

Inside the .ll file, you should be able to find functions like cpu_parallel_range_for and cpu_parallel_range_for_task (if you choose the ti.cpu backend), among various other arcane stuff... . Basically, cpu_parallel_range_for is the entry point of the CPU-backend's thread loop, and the body parameter is the LLVM IR generated from your @ti.kernel.

https://github.com/taichi-dev/taichi/blob/a07b5a5b20a3c2687025a167a422655c4a3b3724/taichi/runtime/llvm/runtime.cpp#L1278-L1287

Sorry I know this is not a great answer and lacks lots of details. But if you start playing around with a simple taichi kernel, it might become intuitive to figure out what's going on... We are trying to implement something that can dump your taichi kernels into a file and load it separately, please stay tuned! :-)

xuhao1 commented 3 years ago

@k-ye Well I am taking a look of this method and trying to understand the LLVM layer of taichi. Also hoping we can see your new easy-to-use dumping tool soon~ Thanks for your excellent work on taichi again.