TT-Metal VCS Simulator

CMakeLists.txt

@@ -69,7 +69,7 @@ option(WITH_PYTHON_BINDINGS "Enables build of python bindings" ON)
 message(STATUS "Build Python bindings: ${WITH_PYTHON_BINDINGS}")
 
 option(ENABLE_CODE_TIMERS "Enable code timers" OFF)
-option(TT_METAL_VERSIM_DISABLED "Disable TT_METAL_VERSIM" ON)
+
 option(ENABLE_TRACY "Enable Tracy Profiling" OFF)
 
 
@@ -121,7 +121,13 @@ endif()
 ############################################################################################################################
 add_library(stdlib INTERFACE)
 if(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
-    target_link_libraries(stdlib INTERFACE c++ c++abi)
+    find_library(LIBC++ c++)
+    find_library(LIBC++ABI c++abi)
+    if(NOT LIBC++ OR NOT LIBC++ABI)
+        message(FATAL_ERROR "libc++ or libc++abi not found. Make sure you have libc++ and libc++abi installed and in your PATH")
+    endif()
+
+    target_link_libraries(stdlib INTERFACE ${LIBC++} ${LIBC++ABI})
     target_compile_options(stdlib INTERFACE -stdlib=libc++)
 else()
     target_link_libraries(stdlib INTERFACE stdc++)
@@ -147,9 +153,6 @@ add_library(compiler_flags INTERFACE)
 target_link_libraries(compiler_flags INTERFACE compiler_warnings stdlib)
 target_compile_options(compiler_flags INTERFACE -mavx2 -fPIC -DFMT_HEADER_ONLY -fvisibility-inlines-hidden -fno-lto)
 
-if(TT_METAL_VERSIM_DISABLED)
-    target_compile_options(compiler_flags INTERFACE -DTT_METAL_VERSIM_DISABLED)
-endif()
 if(ENABLE_CODE_TIMERS)
     target_compile_options(compiler_flags INTERFACE -DTT_ENABLE_CODE_TIMERS)
 endif()
@@ -160,7 +163,7 @@ endif()
 
 add_library(metal_header_directories INTERFACE)
 target_include_directories(metal_header_directories INTERFACE ${PROJECT_SOURCE_DIR}/tt_metal/hw/inc)
-target_include_directories(metal_header_directories SYSTEM INTERFACE ${reflect_SOURCE_DIR})
+target_include_directories(metal_header_directories SYSTEM INTERFACE ${reflect_SOURCE_DIR} ${flatbuffers_SOURCE_DIR}/include ${nanomsg_SOURCE_DIR}/include)
 foreach(lib ${BoostPackages})
     target_include_directories(metal_header_directories INTERFACE ${Boost${lib}_SOURCE_DIR}/include)
 endforeach()

cmake/dependencies.cmake

@@ -56,3 +56,38 @@ CPMAddPackage(
   GITHUB_REPOSITORY boost-ext/reflect
   GIT_TAG v1.1.1
 )
+
+############################################################################################################################
+# Packages needed for tt-metal simulator
+#   NNG for IPC/TCP communication
+#   Google Flatbuffers for serialization
+#   libuv for process mgmt
+############################################################################################################################
+CPMAddPackage(
+  NAME nanomsg
+  GITHUB_REPOSITORY nanomsg/nng
+  GIT_TAG v1.8.0
+  OPTIONS
+      "BUILD_SHARED_LIBS ON"
+      "NNG_TESTS OFF"
+      "NNG_TOOLS OFF"
+)
+CPMAddPackage(
+  NAME flatbuffers
+  GITHUB_REPOSITORY google/flatbuffers
+  GIT_TAG v24.3.25
+  OPTIONS
+      "FLATBUFFERS_BUILD_FLATC OFF"
+      "FLATBUFFERS_BUILD_TESTS OFF"
+      "FLATBUFFERS_INSTALL OFF"
+      "FLATBUFFERS_BUILD_FLATLIB OFF"
+      "FLATBUFFERS_SKIP_MONSTER_EXTRA ON"
+      "FLATBUFFERS_STRICT_MODE ON"
+)
+CPMAddPackage(
+  NAME libuv
+  GITHUB_REPOSITORY libuv/libuv
+  GIT_TAG v1.48.0
+  OPTIONS
+      "LIBUV_BUILD_TESTS OFF"
+)

cmake/umd_device.cmake

@@ -1,3 +1,21 @@
+
+add_library(simulation STATIC
+    ${UMD_HOME}/device/simulation/tt_simulation_device.cpp
+    ${UMD_HOME}/device/simulation/tt_simulation_host.cpp
+)
+target_link_libraries(simulation 
+    PUBLIC compiler_flags 
+    PRIVATE nng uv
+)
+target_include_directories(simulation PUBLIC
+    ${UMD_HOME}
+    ${UMD_HOME}/device
+    ${UMD_HOME}/third_party/fmt/include
+    ${flatbuffers_SOURCE_DIR}/include
+    ${nanomsg_SOURCE_DIR}/include
+    ${libuv_SOURCE_DIR}/include
+)
+
 set(UMD_SRC
     ${UMD_HOME}/device/architecture_implementation.cpp
     ${UMD_HOME}/device/blackhole_implementation.cpp
@@ -10,10 +28,12 @@ set(UMD_SRC
     ${UMD_HOME}/device/tt_silicon_driver.cpp
     ${UMD_HOME}/device/tt_silicon_driver_common.cpp
     ${UMD_HOME}/device/tt_soc_descriptor.cpp
-    ${UMD_HOME}/device/tt_versim_stub.cpp
     ${UMD_HOME}/device/wormhole_implementation.cpp
 )
 
 add_library(umd_device STATIC ${UMD_SRC})
-target_include_directories(umd_device PRIVATE ${UMD_HOME} ${UMD_HOME}/third_party/fmt/include)
-target_link_libraries(umd_device PRIVATE yaml-cpp::yaml-cpp Boost::interprocess rt compiler_flags)
+target_include_directories(umd_device PRIVATE 
+    ${UMD_HOME}
+    ${UMD_HOME}/device
+    ${UMD_HOME}/third_party/fmt/include)
+target_link_libraries(umd_device PRIVATE yaml-cpp::yaml-cpp Boost::interprocess rt simulation)

tests/tt_metal/tt_metal/perf_microbenchmark/common/util.hpp

@@ -66,11 +66,7 @@ inline uint64_t get_t0_to_any_riscfw_end_cycle(tt::tt_metal::Device *device, con
 }
 
 inline int get_tt_npu_clock(tt::tt_metal::Device *device) {
-    int ai_clk = 0;
-#ifdef TT_METAL_VERSIM_DISABLED
-    ai_clk = tt::Cluster::instance().get_device_aiclk(device->id());
-#endif
-    return ai_clk;
+    return tt::Cluster::instance().get_device_aiclk(device->id());
 }
 
 template <typename T>

tt_metal/common/base.hpp

@@ -42,10 +42,8 @@ namespace tt
 */
 enum class TargetDevice : uint8_t
 {
-    Model = 0,
-    Versim = 1,
-    Silicon = 2,
-    Golden = 3,
+    Silicon = 0,
+    Simulator = 1,
     Invalid = 0xFF,
 };
 

tt_metal/common/core_descriptor.hpp

@@ -35,19 +35,16 @@ inline std::string get_core_descriptor_file(const tt::ARCH &arch) {
         wh_arch = "tt_metal/core_descriptors/";
         wh_arch += getenv("WH_ARCH_YAML");
     }
-    bool targeting_versim = false;
-#ifndef TT_METAL_VERSIM_DISABLED
-    targeting_versim = true;
-#endif
-
-    if (targeting_versim) {
+
+    bool targeting_sim = std::getenv("TT_METAL_SIMULATOR_EN") != nullptr; 
+    if (targeting_sim) {
         switch (arch) {
             case tt::ARCH::Invalid: throw std::runtime_error("Invalid arch not supported"); // will be overwritten in tt_global_state constructor
             case tt::ARCH::JAWBRIDGE: throw std::runtime_error("JAWBRIDGE arch not supported");
-            case tt::ARCH::GRAYSKULL: return tt_metal_home + "tt_metal/core_descriptors/grayskull_versim_1x1_arch.yaml";
-            case tt::ARCH::WORMHOLE: throw std::runtime_error("WORMHOLE arch not supported");
-            case tt::ARCH::WORMHOLE_B0: return tt_metal_home + "tt_metal/core_descriptors/wormhole_b0_versim_1x1_arch.yaml";
-            case tt::ARCH::BLACKHOLE: return tt_metal_home + "tt_metal/core_descriptors/blackhole_versim_1x1_arch.yaml";
+            case tt::ARCH::GRAYSKULL: throw std::runtime_error("GRAYSKULL arch not supported for simulator");
+            case tt::ARCH::WORMHOLE: throw std::runtime_error("WORMHOLE arch not supported for simulator");
+            case tt::ARCH::WORMHOLE_B0: throw std::runtime_error("WORMHOLE_B0 arch not supported for simulator");
+            case tt::ARCH::BLACKHOLE: return tt_metal_home + "tt_metal/core_descriptors/blackhole_simulation_1x2_arch.yaml";
             default: throw std::runtime_error("Unsupported device arch");
         };
     } else {

tt_metal/common/test_common.hpp

@@ -23,7 +23,6 @@
 
 inline std::string get_soc_description_file(const tt::ARCH &arch, tt::TargetDevice target_device, string output_dir = "") {
     // Ability to skip this runtime opt, since trimmed SOC desc limits which DRAM channels are available.
-    bool use_full_soc_desc = getenv("TT_METAL_VERSIM_FORCE_FULL_SOC_DESC");
     string tt_metal_home;
     if (getenv("TT_METAL_HOME")) {
         tt_metal_home = getenv("TT_METAL_HOME");
@@ -33,14 +32,14 @@ inline std::string get_soc_description_file(const tt::ARCH &arch, tt::TargetDevi
     if (tt_metal_home.back() != '/') {
         tt_metal_home += "/";
     }
-    if (target_device == tt::TargetDevice::Versim && !use_full_soc_desc) {
+    if (target_device == tt::TargetDevice::Simulator){
         switch (arch) {
-            case tt::ARCH::Invalid: throw std::runtime_error("Invalid arch not supported"); // will be overwritten in tt_global_state constructor
+            case tt::ARCH::Invalid: throw std::runtime_error("Invalid arch not supported");
             case tt::ARCH::JAWBRIDGE: throw std::runtime_error("JAWBRIDGE arch not supported");
-            case tt::ARCH::GRAYSKULL: return tt_metal_home + "tt_metal/soc_descriptors/grayskull_versim_1x1_arch.yaml";
+            case tt::ARCH::GRAYSKULL: throw std::runtime_error("GRAYSKULL arch not supported");
             case tt::ARCH::WORMHOLE: throw std::runtime_error("WORMHOLE arch not supported");
-            case tt::ARCH::WORMHOLE_B0: return tt_metal_home + "tt_metal/soc_descriptors/wormhole_b0_versim_1x1_arch.yaml";
-            case tt::ARCH::BLACKHOLE: return tt_metal_home + "tt_metal/soc_descriptors/blackhole_versim_1x1_arch.yaml";
+            case tt::ARCH::WORMHOLE_B0: throw std::runtime_error("WORMHOLE_B0 arch not supported");
+            case tt::ARCH::BLACKHOLE: return tt_metal_home + "tt_metal/soc_descriptors/blackhole_simulation_1x2_arch.yaml";
             default: throw std::runtime_error("Unsupported device arch");
         };
     } else {

tt_metal/core_descriptors/blackhole_simulation_1x2_arch.yaml

@@ -0,0 +1,34 @@
+# Anything using [[#, #]] is logical coordinates (Can be relative)
+# relative index: 0 means first row, -1 means last row of functional grid...
+
+# product name:
+#   num of HW command queues:
+#     core descriptor config
+
+blackhole:
+  1:
+    compute_with_storage_grid_range: # Logical only start and end [x, y]
+      start: [0, 1]
+      end: [1, 1]
+
+    storage_cores: # Relative to grid of tensix cores
+      []
+
+    dispatch_cores:
+      []
+
+    dispatch_core_type:
+      "tensix"
+  2:
+    compute_with_storage_grid_range: # Logical only start and end [x, y]
+      start: [0, 1]
+      end: [1, 1]
+
+    storage_cores: # Relative to grid of tensix cores
+      []
+
+    dispatch_cores:
+      []
+
+    dispatch_core_type:
+      "tensix"

tt_metal/hw/inc/grayskull/tensix.h

@@ -11,7 +11,6 @@
 #include <tensix_types.h>
 #include "cfg_defines.h"
 #include "dev_mem_map.h"
-#include "dev_mem_map_versim.h"
 
 // Convenience and type defines
 typedef std::uint32_t uint;

tt_metal/hw/inc/wormhole/tensix.h

@@ -11,7 +11,6 @@
 #include <tensix_types.h>
 #include "cfg_defines.h"
 #include "dev_mem_map.h"
-#include "dev_mem_map_versim.h"
 
 // Convenience and type defines
 typedef std::uint32_t uint;

tt_metal/impl/device/device.cpp

@@ -147,10 +147,8 @@ void Device::initialize_cluster() {
     if (llrt::OptionsG.get_clear_l1()) {
         this->clear_l1_state();
     }
-#ifdef TT_METAL_VERSIM_DISABLED
     int ai_clk = tt::Cluster::instance().get_device_aiclk(this->id_);
     log_info(tt::LogMetal, "AI CLK for device {} is:   {} MHz", this->id_, ai_clk);
-#endif
 }
 
 void Device::initialize_allocator(size_t l1_small_size, size_t trace_region_size, const std::vector<uint32_t> &l1_bank_remap) {

tt_metal/llrt/llrt.hpp

@@ -13,7 +13,7 @@
 // clang-format off
 #include "llrt/tt_cluster.hpp"
 #include "tensix.h"
-#include "tt_metal/third_party/umd/device/device_api.h"
+#include "tt_metal/third_party/umd/device/device_api_metal.h"
 #include "tt_metal/third_party/umd/device/tt_xy_pair.h"
 #include "llrt_common/tiles.hpp"
 #include "llrt/tt_memory.h"

tt_metal/llrt/tlb_config.hpp

@@ -4,7 +4,7 @@
 
 #pragma once
 
-#include "third_party/umd/device/device_api.h"
+#include "third_party/umd/device/device_api_metal.h"
 #include "tt_metal/common/tt_backend_api_types.hpp"
 #include "tt_metal/common/metal_soc_descriptor.h"