print_kernel_llvm_ir_optimized; reproduce Root_lookup_element not inl…

…ined
jamusven · Jan 1, 2020 · a2b591d · a2b591d
1 parent 8cb6d92
commit a2b591d
Show file tree

Hide file tree

Showing 5 changed files with 11 additions and 1 deletion.
diff --git a/examples/mpm_snow.py b/examples/mpm_snow.py
@@ -22,6 +22,8 @@
 grid_m = ti.var(dt=ti.f32, shape=(n_grid, n_grid))
 ti.cfg.arch = ti.cuda # Run on a GPU if equipped
 ti.cfg.enable_profiler = True
+ti.cfg.print_kernel_llvm_ir = True
+ti.cfg.print_kernel_llvm_ir_optimized = True
 
 @ti.kernel
 def substep():
@@ -93,6 +95,7 @@ def substep():
   t = time.time()
   for s in range(100):
     grid_v.fill([0, 0])
+    exit(0)
     grid_m.fill(0)
     substep()
   print('{:.4f} ms'.format((time.time() - t) * 10))

diff --git a/taichi/backends/codegen_llvm_ptx.cpp b/taichi/backends/codegen_llvm_ptx.cpp
@@ -69,6 +69,9 @@ class CodeGenLLVMGPU : public CodeGenLLVM {
     }
 
     auto ptx = compile_module_to_ptx(module);
+    if (get_current_program().config.print_kernel_llvm_ir_optimized) {
+      TC_P(ptx);
+    }
     auto cuda_module = cuda_context->compile(ptx);
 
     for (auto &task : offloaded_local) {

diff --git a/taichi/backends/codegen_llvm_x86.cpp b/taichi/backends/codegen_llvm_x86.cpp
@@ -99,7 +99,8 @@ void global_optimize_module_x86_64(std::unique_ptr<llvm::Module> &module) {
   module_pass_manager.run(*module);
   t = Time::get_time() - t;
   // TC_INFO("Global optimization time: {} ms", t * 1000);
-  // module->print(llvm::errs(), nullptr);
+  if (get_current_program().config.print_kernel_llvm_ir_optimized)
+    module->print(llvm::errs(), nullptr);
 }
 
 TLANG_NAMESPACE_END
diff --git a/taichi/python_bindings.cpp b/taichi/python_bindings.cpp
@@ -54,6 +54,8 @@ void export_lang(py::module &m) {
                      &CompileConfig::print_struct_llvm_ir)
       .def_readwrite("print_kernel_llvm_ir",
                      &CompileConfig::print_kernel_llvm_ir)
+      .def_readwrite("print_kernel_llvm_ir_optimized",
+                     &CompileConfig::print_kernel_llvm_ir_optimized)
       .def_readwrite("simplify_before_lower_access",
                      &CompileConfig::simplify_before_lower_access)
       .def_readwrite("simplify_after_lower_access",

diff --git a/taichi/tlang_util.h b/taichi/tlang_util.h
@@ -338,6 +338,7 @@ struct CompileConfig {
   bool use_llvm;
   bool print_struct_llvm_ir;
   bool print_kernel_llvm_ir;
+  bool print_kernel_llvm_ir_optimized;
   bool verbose_kernel_launches;
   bool enable_profiler;
   bool verbose;