feat(3rd-party): Add fmtlog support.

.devcontainer/Dockerfile

@@ -0,0 +1,9 @@
+# FROM ubuntu:22.04
+FROM mcr.microsoft.com/vscode/devcontainers/base:ubuntu-22.04
+ENV DEBIAN_FRONTEND=noninteractive 
+
+# Install dependencies.
+RUN apt update && apt-get install -y git make cmake build-essential python-is-python3 python-dev-is-python3 python3-pip libdw-dev openssh-client
+
+# Update pip and switch to Tsinghua source.
+RUN python -m pip install -i https://pypi.tuna.tsinghua.edu.cn/simple --upgrade pip && pip config set global.index-url https://pypi.tuna.tsinghua.edu.cn/simple

.devcontainer/devcontainer.json

@@ -0,0 +1,26 @@
+// For format details, see https://aka.ms/devcontainer.json. For config options, see the
+// README at: https://github.com/devcontainers/templates/tree/main/src/cpp
+{
+    "name": "C++",
+    "build": {
+        "dockerfile": "Dockerfile"
+    },
+    // Features to add to the dev container. More info: https://containers.dev/features.
+    // "features": {},
+    // Configure tool-specific properties.
+    "customizations": {
+        // Configure properties specific to VS Code.
+        "vscode": {
+            "settings": {},
+            "extensions": [
+                "streetsidesoftware.code-spell-checker"
+            ]
+        }
+    }
+    // Use 'forwardPorts' to make a list of ports inside the container available locally.
+    // "forwardPorts": [],
+    // Use 'postCreateCommand' to run commands after the container is created.
+    // "postCreateCommand": "gcc -v",
+    // Uncomment to connect as root instead. More info: https://aka.ms/dev-containers-non-root.
+    // "remoteUser": "root"
+}

CMakeLists.txt

@@ -6,6 +6,17 @@ option(USE_CUDA "Support CUDA GPU" OFF)
 option(BUILD_TEST "Build test code" ON)
 option(BUILD_ASAN "Build code whith ASAN" OFF)
 
+option(DEBUG "Debug mode" OFF)
+
+if(DEBUG)
+  message("Debug mode is enabled.")
+  add_definitions(-DDEBUG)
+  set(CMAKE_BUILD_TYPE Debug)
+else()
+  message("Debug mode is disabled.")
+  set(CMAKE_BUILD_TYPE Release)
+endif()
+
 set(CMAKE_CXX_STANDARD 20)
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
 set(CMAKE_CXX_EXTENSIONS OFF) 
@@ -19,12 +30,17 @@ if(BUILD_ASAN)
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${CMAKE_ASAN_FLAGS}")
 endif()
 
+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
+
 include_directories(include)
-# include_directories(3rd-party/fmt/include)
-# include_directories(3rd-party/result)
-# include_directories(3rd-party/fmtlog)
-# add_definitions(-D FMTLOG_HEADER_ONLY)
-# add_definitions(-D FMT_HEADER_ONLY)
+
+include_directories(3rd-party/fmt/include)
+add_subdirectory(3rd-party/fmt)
+add_definitions(-D FMT_HEADER_ONLY)
+
+include_directories(3rd-party/fmtlog)
+add_definitions(-D FMTLOG_HEADER_ONLY)
+
 
 if(BUILD_TEST)
     include_directories(3rd-party/googletest/googletest/include)
@@ -42,13 +58,8 @@ file(GLOB_RECURSE SOURCE
   src/graph/*.cc
 )
 
-set(CMAKE_POSITION_INDEPENDENT_CODE ON)
 add_library(tiledkernel SHARED ${SOURCE})
 
-include_directories(3rd-party/fmt/include)
-add_subdirectory(3rd-party/fmt)
-
-
 add_compile_options(-Wall)
 add_compile_options(-Werror)
 

Makefile

@@ -1,21 +1,28 @@
 CC := g++
-EXAMPLE ?= rf_graph
+EXAMPLE ?= shared_graph
 EXAMPLE_SRCS := $(wildcard examples/*.cc)
 EXAMPLES := $(patsubst examples/%.cc, %, $(EXAMPLE_SRCS))
 
 LD_FLAGS := -Lbuild/ -ltiledkernel -Wl,-rpath,build/
-INC_FLAGS := -Iinclude -I3rd-party/fmt/include
-MACRO_FLAGS := -DFMT_HEADER_ONLY
+INC_FLAGS := -Iinclude -I3rd-party/fmt/include -I3rd-party/fmtlog
+MACRO_FLAGS := -DFMT_HEADER_ONLY -DFMTLOG_HEADER_ONLY
+
+CMAKE_OPTIONS :=
 
 CUDA ?= ON
+DEBUG ?= OFF
+
+ifeq ($(DEBUG), ON)
+	CMAKE_OPTIONS += -DDEBUG=ON
+endif
 
 BUILD := build
 
 .PHONY: build example 
 
 build:
 	@mkdir -p build
-	@cd build && cmake .. && make
+	cd build && cmake $(CMAKE_OPTIONS) .. && make
 
 example: build
 	@$(CC) -std=c++17 examples/$(EXAMPLE).cc $(INC_FLAGS) $(LD_FLAGS) $(MACRO_FLAGS) -o build/$(EXAMPLE)

examples/rf_graph.cc

@@ -7,14 +7,15 @@
 #include "platform.hpp"
 #include "op.hpp"
 #include <iostream>
+#include <fmtlog.h>
 
 using namespace tiledkernel;
 using namespace tiledkernel::type;
 using namespace tiledkernel::graph;
 
 int main() {
     // Build a RF Gemm graph
-
+    fmtlog::setLogLevel(fmtlog::LogLevel::DBG);
     std::cout << "Run RF GEMM Graph Example:" << std::endl << std::endl;
     // Define buffers
     auto sA = std::make_shared<TiledBuffer>("sA", MemoryLevel::Shared,

examples/shared_graph.cc

@@ -3,16 +3,18 @@
 #include "graph/tilededge.hpp"
 #include "type/data_type.hpp"
 #include "tiledbuffer.hpp"
+#include "generator.hpp"
 #include "op.hpp"
 #include <iostream>
+#include <fmtlog.h>
 
 using namespace tiledkernel;
 using namespace tiledkernel::type;
 using namespace tiledkernel::graph;
 
 int main() {
+    fmtlog::setLogLevel(fmtlog::LogLevel::DBG);
     // Build a Shared Gemm graph
-
     std::cout << "Shared GEMM Graph Example:" << std::endl;
     // Define buffers
     auto sA = std::make_shared<TiledBuffer>("sA", MemoryLevel::Shared,
@@ -99,11 +101,37 @@ int main() {
         std::vector<EdgePtr>{acc_sC_edge},
         std::vector<EdgePtr>{rA_gemm_edge, rB_gemm_edge, gemm_acc_edge});
 
+    auto rf_sorted_nodes = rf_gemm_graph->topoSort();
+    std::cout << "ID\tName" << std::endl;
+    for (auto node : rf_sorted_nodes) {
+        std::cout << node->id << "\t" << node->name << std::endl;
+    }
+
     auto rf_gemm_node = std::make_shared<TiledNode>(
         NodeType::Task, MemoryLevel::RF, TiledNodeData{rf_gemm_graph},
         "rf_gemm_node", std::vector<EdgePtr>{sA_rA_edge, sB_rB_edge},
         std::vector<EdgePtr>{acc_sC_edge});
 
+    // Define Access Map
+    auto rA_gemm_access_map_i = std::make_shared<AccessMap>(
+        1, 1, std::vector<std::vector<int32_t>>{std::vector<int32_t>{1}},
+        std::vector<std::pair<int32_t, int32_t>>{std::make_pair(0, 10)},
+        std::vector<int32_t>{0});
+
+    auto rB_gemm_access_map_i = std::make_shared<AccessMap>(
+        1, 1, std::vector<std::vector<int32_t>>{std::vector<int32_t>{1}},
+        std::vector<std::pair<int32_t, int32_t>>{std::make_pair(0, 10)},
+        std::vector<int32_t>{0});
+
+    auto gemm_acc_access_map_i = std::make_shared<AccessMap>(
+        1, 1, std::vector<std::vector<int32_t>>{std::vector<int32_t>{0}},
+        std::vector<std::pair<int32_t, int32_t>>{std::make_pair(0, 10)},
+        std::vector<int32_t>{0});
+
+    rA_gemm_edge->setAccessMapI(rA_gemm_access_map_i);
+    rB_gemm_edge->setAccessMapI(rB_gemm_access_map_i);
+    gemm_acc_edge->setAccessMapI(gemm_acc_access_map_i);
+
     std::cout << "Connect Shared GEMM Graph:" << std::endl;
     auto shared_gemm_graph = std::make_shared<TiledGraph>(
         MemoryLevel::Shared, "shared_gemm_graph",
@@ -118,4 +146,10 @@ int main() {
     for (auto node : sorted_nodes) {
         std::cout << node->id << "\t" << node->name << std::endl;
     }
+
+    auto generator = std::make_shared<TiledGenerator>();
+
+    auto kernel = generator->emit(Platform::Cute, shared_gemm_graph);
+
+    std::cout << "Generate kernel:" << std::endl << kernel << std::endl;
 }

include/generator.hpp

@@ -19,8 +19,14 @@ namespace tiledkernel {
        private:
         TiledContext::Pointer ctx;
 
+        // Emit code for the given platform.
         std::string emit_cute(TiledGraph::Pointer graph);
+
+        // Emit code for the given memory hierarchy graph.
         std::string emit_rf_cute(TiledGraph::Pointer graph);
+        std::string emit_shared_cute(TiledGraph::Pointer graph);
+        std::string emit_global_cute(TiledGraph::Pointer graph);
+
         std::string emit_rf_cute_gemm(TiledNode::Pointer node);
 
         std::string generate_loop(IterVar::Pointer iter_var, std::string body);

include/graph/tiledgraph.hpp

@@ -19,6 +19,10 @@ namespace tiledkernel::graph {
 
         MemoryLevel getMemLevel() { return mem_level; }
 
+        void connect();
+
+        void clear();
+
         std::vector<std::shared_ptr<TiledNode>> topoSort();
 
         bool isNodeExist(std::shared_ptr<TiledNode> node);
@@ -33,7 +37,5 @@ namespace tiledkernel::graph {
         std::vector<std::shared_ptr<TiledEdge>> in_edges;
         std::vector<std::shared_ptr<TiledEdge>> out_edges;
         std::vector<std::shared_ptr<TiledEdge>> intra_edges;
-
-        void connect();
     };
 };  // namespace tiledkernel::graph

include/graph/tilednode.hpp

@@ -59,6 +59,8 @@ namespace tiledkernel::graph {
             return {};
         }
 
+        TiledNodeData getData() { return data; }
+
         std::vector<std::shared_ptr<TiledEdge>> getInEdges() {
             return in_edges;
         }

include/id.hpp

@@ -20,6 +20,8 @@ namespace tiledkernel {
 
         bool isValid() const { return id_ != 0; }
 
+        uint64_t value() const { return id_; }
+
         operator uint64_t() const { return id_; }
     };
 };  // namespace tiledkernel

src/generator.cc

@@ -3,7 +3,9 @@
 #include "generator.hpp"
 #include "error_handler.hpp"
 #include <fmt/core.h>
+#include <fmtlog.h>
 #include <sstream>
+#include <iostream>
 
 namespace tiledkernel {
 
@@ -23,9 +25,9 @@ namespace tiledkernel {
         std::string kernel = "";
         switch (graph->getMemLevel()) {
             case MemoryLevel::Global:
-                break;
+                kernel += emit_global_cute(graph);
             case MemoryLevel::Shared:
-                break;
+                kernel += emit_shared_cute(graph);
             case MemoryLevel::RF:
                 kernel += emit_rf_cute(graph);
         }
@@ -60,6 +62,44 @@ namespace tiledkernel {
         return kernel;
     }
 
+    std::string TiledGenerator::emit_shared_cute(TiledGraph::Pointer graph) {
+        std::string kernel;
+        auto sorted_nodes = graph->topoSort();
+
+        std::vector<TiledNode::Pointer> compute_nodes;
+        std::vector<TiledNode::Pointer> rf_nodes;
+
+        for (auto node : sorted_nodes) {
+            if (node->getMemLevel() == MemoryLevel::RF) {
+                rf_nodes.push_back(node);
+            }
+        }
+
+        kernel += "__syncthreads();\n";
+        for (auto node : rf_nodes) {
+            auto graph = std::get<TiledGraph::Pointer>(node->getData());
+            auto sorted_nodes = graph->topoSort();
+            logd("ID\tName");
+#ifdef DEBUG
+            fmt::println("[DEBUG] ID\tName");
+#endif
+            for (auto node : sorted_nodes) {
+                logd("{}\t{}", node->id.value(), node->name);
+#ifdef DEBUG
+                fmt::println("[DEBUG] {}\t{}", node->id.value(), node->name);
+#endif
+            }
+            kernel += emit_rf_cute(graph);
+        }
+        kernel += "__syncthreads();\n";
+
+        return kernel;
+    }
+
+    std::string TiledGenerator::emit_global_cute(TiledGraph::Pointer graph) {
+        return "";
+    }
+
     std::string TiledGenerator::emit_rf_cute_gemm(TiledNode::Pointer node) {
         auto predecessors = node->getPredecessors();
         auto successors = node->getSuccessors();

src/graph/tiledgraph.cc

@@ -58,6 +58,8 @@ namespace tiledkernel::graph {
         return false;
     }
 
+    void TiledGraph::clear() {}
+
     void TiledGraph::connect() {
         for (auto node : nodes) {
             for (auto edge : node->in_edges) {

tests/graph/test_toposort.cc

@@ -6,14 +6,15 @@
 #include "op.hpp"
 
 #include <gtest/gtest.h>
+#include <fmtlog.h>
 
 using namespace tiledkernel;
 using namespace tiledkernel::type;
 using namespace tiledkernel::graph;
 
 TEST(Graph, tiledgraph_topo_rf_gemm_graph) {
+    fmtlog::setLogLevel(fmtlog::LogLevel::INF);
     // Build a RF Gemm graph
-
     // Define buffers
     auto sA = std::make_shared<TiledBuffer>("sA", MemoryLevel::Shared,
                                             DataType::Float32);