[cc] Add struct compiler to the experimental C backend

.github/workflows/persubmit.yml

@@ -31,12 +31,13 @@ jobs:
           export TAICHI_REPO_DIR=`pwd`
           export PATH=$TAICHI_REPO_DIR/taichi-llvm/bin/:$PATH
           export CXX=clang++
-          export CI_SETUP_CMAKE_ARGS="-DTI_WITH_OPENGL:BOOL=OFF -DTI_WITH_CC:BOOL=${{matrix.with_cc}}"
           # GLFW dependencies:
           #export CI_SETUP_CMAKE_ARGS=-DTI_WITH_OPENGL:BOOL=ON
           #sudo apt-get install libx11-dev libxcursor-dev libxi-dev
           #sudo apt-get install libxrandr-dev libxinerama-dev libglvnd-dev
           python misc/ci_setup.py ci
+        env:
+          CI_SETUP_CMAKE_ARGS: -DTI_WITH_OPENGL:BOOL=OFF -DTI_WITH_CC:BOOL=${{ matrix.with_cc }}
 
       - name: Functionallity Test
         run: |

hello.py

@@ -0,0 +1,23 @@
+import taichi as ti
+
+ti.init(arch=ti.cc)
+ti.core.toggle_advanced_optimization(False)
+
+x = ti.var(ti.i32)
+y = ti.var(ti.f32)
+z = ti.var(ti.f64)
+
+blk0 = ti.root
+blk1 = blk0.dense(ti.i, 8)
+blk2 = blk1.dense(ti.i, 4)
+blk1.place(x)
+blk1.place(y)
+blk2.place(z)
+
+
+@ti.kernel
+def func():
+    print('Nice to meet you!', z[3])
+
+
+func()

taichi/backends/cc/cc_configuation.h

@@ -6,9 +6,11 @@ TLANG_NAMESPACE_BEGIN
 namespace cccp {
 
 struct CCConfiguation {
-  std::string compile_cmd;
+  std::string compile_cmd, link_cmd;
 
-  CCConfiguation() : compile_cmd("gcc -shared -fPIC -o '{}' '{}'") {
+  CCConfiguation()
+      : compile_cmd("gcc -c -o '{}' '{}'"),
+        link_cmd("gcc -shared -fPIC -o '{}' '{}'") {
   }
 };
 

taichi/backends/cc/cc_kernel.h

@@ -9,18 +9,26 @@ class CCProgram;
 
 class CCKernel {
  public:
-  CCKernel(std::string const &source, std::string const &name)
-      : source(source), name(name) {
+  CCKernel(CCProgram *program,
+           std::string const &source,
+           std::string const &name)
+      : program(program), source(source), name(name) {
   }
 
   void compile();
-  void launch(CCProgram *program, Context *ctx);
+  void launch(Context *ctx);
+  std::string get_object() {
+    return obj_path;
+  }
+
+ private:
+  CCProgram *program;
 
   std::string source;
   std::string name;
 
   std::string src_path;
-  std::string bin_path;
+  std::string obj_path;
 };
 
 }  // namespace cccp

taichi/backends/cc/cc_layout.h

@@ -10,7 +10,17 @@ class CCLayout {
   CCLayout() {
   }
 
+  std::string get_object() {
+    return obj_path;
+  }
+
+  void compile();
+
   std::string source;
+
+ private:
+  std::string src_path;
+  std::string obj_path;
 };
 
 }  // namespace cccp

taichi/backends/cc/cc_program.cpp

@@ -12,28 +12,65 @@ namespace cccp {
 CCConfiguation cfg;
 
 void CCKernel::compile() {
-  bin_path = fmt::format("{}/{}.so", runtime_tmp_dir, name);
+  obj_path = fmt::format("{}/{}.o", runtime_tmp_dir, name);
   src_path = fmt::format("{}/{}.c", runtime_tmp_dir, name);
 
   std::ofstream(src_path) << source;
-  TI_INFO("[cc] compiling kernel [{}]:\n{}\n", name, source);
-  execute(cfg.compile_cmd, bin_path, src_path);
+  TI_DEBUG("[cc] compiling [{}] -> [{}]:\n{}\n", name, src_path, source);
+  execute(cfg.compile_cmd, obj_path, src_path);
 }
 
-void CCKernel::launch(CCProgram *launcher, Context *ctx) {
-  using FuncEntryType = void();
-  DynamicLoader dll(bin_path);
-  TI_ASSERT_INFO(dll.loaded(), "[cc] could not load shared object: {}",
-                 bin_path);
-  auto main =
-      reinterpret_cast<FuncEntryType *>(dll.load_function(get_sym_name(name)));
-  TI_INFO("[cc] entering kernel [{}]", name);
-  (*main)();
-  TI_INFO("[cc] leaving kernel [{}]", name);
+void CCKernel::launch(Context *ctx) {
+  program->relink();
+  auto entry = program->load_kernel(name);
+  TI_TRACE("[cc] entering kernel [{}]", name);
+  (*entry)();
+  TI_TRACE("[cc] leaving kernel [{}]", name);
 }
 
-void CCProgram::launch(CCKernel *kernel, Context *ctx) {
-  kernel->launch(this, ctx);
+void CCLayout::compile() {
+  obj_path = fmt::format("{}/_root.o", runtime_tmp_dir);
+  src_path = fmt::format("{}/_root.c", runtime_tmp_dir);
+
+  std::ofstream(src_path) << source
+                          << "\n\nstruct S0root *_Ti_get_root() {\n\tstatic "
+                             "struct S0root r;\n\treturn &r;\n}\n";
+  TI_DEBUG("[cc] compiling root struct -> [{}]:\n{}\n", obj_path, source);
+  execute(cfg.compile_cmd, obj_path, src_path);
+}
+
+void CCProgram::relink() {
+  if (!need_relink)
+    return;
+
+  dll_path = fmt::format("{}/libti_program.so", runtime_tmp_dir);
+
+  std::vector<std::string> objects;
+  objects.push_back(layout->get_object());
+  for (auto const &ker : kernels) {
+    objects.push_back(ker->get_object());
+  }
+
+  TI_DEBUG("[cc] linking shared object [{}] with [{}]", dll_path,
+           fmt::join(objects, "] ["));
+  execute(cfg.link_cmd, dll_path, fmt::join(objects, "' '"));
+
+  TI_DEBUG("[cc] loading shared object: {}", dll_path);
+  dll = std::make_unique<DynamicLoader>(dll_path);
+  TI_ASSERT_INFO(dll->loaded(), "[cc] could not load shared object: {}",
+                 dll_path);
+
+  need_relink = false;
+}
+
+void CCProgram::add_kernel(std::unique_ptr<CCKernel> kernel) {
+  kernels.push_back(std::move(kernel));
+  need_relink = true;
+}
+
+CCFuncEntryType *CCProgram::load_kernel(std::string const &name) {
+  return reinterpret_cast<CCFuncEntryType *>(
+      dll->load_function(get_func_sym(name)));
 }
 
 CCProgram::CCProgram() {

taichi/backends/cc/cc_program.h

@@ -2,22 +2,34 @@
 
 #include "taichi/lang_util.h"
 
+TI_NAMESPACE_BEGIN
+class DynamicLoader;
+TI_NAMESPACE_END
+
 TLANG_NAMESPACE_BEGIN
+
 namespace cccp {
 
 class CCKernel;
 class CCLayout;
 
+using CCFuncEntryType = void();
+
 class CCProgram {
   // Launch C compiler to compile generated source code, and run them
  public:
   CCProgram();
   ~CCProgram();
 
-  void launch(CCKernel *kernel, Context *ctx);
+  void add_kernel(std::unique_ptr<CCKernel> kernel);
+  CCFuncEntryType *load_kernel(std::string const &name);
+  void relink();
 
   std::vector<std::unique_ptr<CCKernel>> kernels;
   std::unique_ptr<CCLayout> layout;
+  std::unique_ptr<DynamicLoader> dll;
+  std::string dll_path;
+  bool need_relink{true};
 };
 
 }  // namespace cccp

taichi/backends/cc/cc_utils.h

@@ -16,21 +16,27 @@ inline std::string cc_data_type_name(DataType dt) {
       return "int";
     case DataType::f32:
       return "float";
+    case DataType::f64:
+      return "double";
     default:
-      TI_NOT_IMPLEMENTED
+      TI_ERROR("Unsupported DataType={} on C backend", data_type_name(dt));
   }
 }
 
-inline std::string get_sym_name(std::string const &name) {
+inline std::string get_func_sym(std::string const &name) {
   return fmt::format("Ti_{}", name);
 }
 
+inline std::string get_data_sym(std::string const &name) {
+  return fmt::format("ti_{}", name);
+}
+
 template <typename... Args>
 inline int execute(std::string fmt, Args &&... args) {
   auto cmd = fmt::format(fmt, std::forward<Args>(args)...);
-  TI_INFO("Executing command: {}", cmd);
+  TI_TRACE("Executing command: {}", cmd);
   int ret = std::system(cmd.c_str());
-  TI_INFO("Command exit status: {}", ret);
+  TI_TRACE("Command exit status: {}", ret);
   return ret;
 }