Changing how we setup simulator. (#16375)

Previously, `tt_SimulationDevice` needed to be setup with two arguments:
1. Constructor argument with soc descriptor path.
2. Environment variable that points to simulator run script.

Change in `tt_SimulationDevice` now requires only one argument, path to
directory that contains all necessary files for simulator to run. That
directory contains `soc_descriptor.yaml`, `run.sh` script and simulator
application.

This PR also removes need for two environment variables
(`TT_METAL_SIMULATOR_EN` and `TT_REMOTE_EXE`) and replaces them with
single `TT_METAL_SIMULATOR` variable that contains path to simulator
build directory.

In the future, when tt-metal is built only once for all architectures,
this would mean that one would setup `TT_METAL_SIMULATOR` environment
variable that point to simulator directory and would run it against
architecture that is compiled in that directory.

Related PRs:
- [UMD
simulators](https://yyz-gitlab.local.tenstorrent.com/tenstorrent/tt-metal/-/merge_requests/15)
- [UMD](tenstorrent/tt-umd#449)

tests/tt_metal/test_utils/env_vars.hpp

@@ -7,6 +7,7 @@
 
 #include "umd/device/device_api_metal.h"
 #include "umd/device/tt_cluster_descriptor.h"
+#include "umd/device/tt_simulation_device.h"
 
 #include <string>
 
@@ -40,8 +41,9 @@ inline std::string get_env_arch_name() {
 
 inline std::string get_umd_arch_name() {
 
-    if(std::getenv("TT_METAL_SIMULATOR_EN")) {
-        return get_env_arch_name();
+    if(std::getenv("TT_METAL_SIMULATOR")) {
+        tt_SimulationDeviceInit init(std::getenv("TT_METAL_SIMULATOR"));
+        return tt::arch_to_str(init.get_arch_name());
     }
 
     auto cluster_desc = tt_ClusterDescriptor::create();

tt_metal/api/tt-metalium/core_descriptor.hpp

@@ -36,13 +36,13 @@ inline std::string get_core_descriptor_file(
     }
     core_desc_dir += "tt_metal/core_descriptors/";
 
-    bool targeting_sim = std::getenv("TT_METAL_SIMULATOR_EN") != nullptr;
+    bool targeting_sim = std::getenv("TT_METAL_SIMULATOR") != 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::GRAYSKULL: throw std::runtime_error("GRAYSKULL arch not supported for simulator");
+            case tt::ARCH::GRAYSKULL: return core_desc_dir + "grayskull_versim_1x1_arch.yaml";
             case tt::ARCH::WORMHOLE_B0: return core_desc_dir + "wormhole_b0_versim_1x1_arch.yaml";
             case tt::ARCH::BLACKHOLE: return core_desc_dir + "blackhole_simulation_1x2_arch.yaml";
             default: throw std::runtime_error("Unsupported device arch");

tt_metal/api/tt-metalium/get_platform_architecture.hpp

@@ -10,6 +10,7 @@
 #include "assert.hpp"
 #include "umd/device/pci_device.hpp"
 #include "umd/device/tt_soc_descriptor.h"
+#include "umd/device/tt_simulation_device.h"
 
 namespace tt::tt_metal {
 
@@ -18,21 +19,17 @@ namespace tt::tt_metal {
  *
  * This function determines the platform architecture by inspecting the environment
  * variables or available physical devices. If the environment variable
- * `TT_METAL_SIMULATOR_EN` is set, the architecture is retrieved from the
- * `ARCH_NAME` environment variable. Otherwise, the architecture is deduced
- * by detecting available physical devices.
+ * `TT_METAL_SIMULATOR` is set, the architecture is retrieved from simulator.
+ * Otherwise, the architecture is deduced by detecting available physical devices.
  *
  * @return tt::ARCH The detected platform architecture. Returns tt::ARCH::Invalid
  *                  if no valid architecture could be detected.
  *
  * @note
- * - If the system is in simulation mode (`TT_METAL_SIMULATOR_EN` is set),
- *   the `ARCH_NAME` environment variable must be defined.
  * - A fatal error occurs if multiple devices are detected with conflicting
  *   architectures.
  *
  * @exception std::runtime_error Throws a fatal error if:
- * - `ARCH_NAME` is not set when `TT_METAL_SIMULATOR_EN` is enabled.
  * - Multiple devices with inconsistent architectures are detected.
  *
  * Example usage:
@@ -52,12 +49,10 @@ namespace tt::tt_metal {
  */
 inline tt::ARCH get_platform_architecture() {
     auto arch = tt::ARCH::Invalid;
-    if (std::getenv("TT_METAL_SIMULATOR_EN")) {
-        auto arch_env = std::getenv("ARCH_NAME");
-        TT_FATAL(arch_env, "ARCH_NAME env var needed for VCS");
-        arch = tt::get_arch_from_string(arch_env);
+    if (std::getenv("TT_METAL_SIMULATOR")) {
+        tt_SimulationDeviceInit init(std::getenv("TT_METAL_SIMULATOR"));
+        arch = init.get_arch_name();
     } else {
-
         // Issue tt_umd#361: tt_ClusterDescriptor::create() won't work here.
         // This map holds PCI info for each mmio chip.
         auto devices_info = PCIDevice::enumerate_devices_info();

tt_metal/common/core_descriptor.cpp

@@ -126,7 +126,7 @@ const core_descriptor_t& get_core_descriptor_config(
         }
         dispatch_cores.push_back(coord);
     }
-    TT_ASSERT(dispatch_cores.size() || std::getenv("TT_METAL_SIMULATOR_EN"), "Dispatch cores size must be positive");
+    TT_ASSERT(dispatch_cores.size() || std::getenv("TT_METAL_SIMULATOR"), "Dispatch cores size must be positive");
 
     std::vector<CoreCoord> logical_compute_cores;
     logical_compute_cores.reserve(compute_cores.size());

tt_metal/core_descriptors/grayskull_versim_1x1_arch.yaml

@@ -17,3 +17,19 @@ E150:
 
       dispatch_cores:
         []
+
+      dispatch_core_type:
+        "tensix"
+    2:
+      compute_with_storage_grid_range:
+        start: [0, 0]
+        end: [0, 0]
+
+      storage_cores:
+        []
+
+      dispatch_cores:
+        []
+
+      dispatch_core_type:
+        "tensix"

tt_metal/llrt/hal.cpp

@@ -22,6 +22,8 @@ Hal::Hal() : arch_(get_platform_architecture()) {
 
         case tt::ARCH::BLACKHOLE: initialize_bh(); break;
 
+        case tt::ARCH::QUASAR: TT_THROW("HAL doesn't support Quasar"); break;
+
         case tt::ARCH::Invalid: /*TT_THROW("Unsupported arch for HAL")*/; break;
     }
 }

tt_metal/llrt/llrt.cpp

@@ -219,7 +219,9 @@ void wait_until_cores_done(
     // poll the cores until the set of not done cores is empty
     int loop_count = 1;
     auto start = std::chrono::high_resolution_clock::now();
-    if (std::getenv("TT_METAL_SIMULATOR_EN")) timeout_ms = 0;
+    bool is_simulator = std::getenv("TT_METAL_SIMULATOR") != nullptr;
+
+    if (is_simulator) timeout_ms = 0;
     while (!not_done_phys_cores.empty()) {
         if (timeout_ms > 0) {
             auto now = std::chrono::high_resolution_clock::now();
@@ -253,6 +255,11 @@ void wait_until_cores_done(
             }
         }
         loop_count++;
+
+        // Continuously polling cores on simulator can cause it to run much slower than real hardware.
+        if (is_simulator)
+            std::this_thread::sleep_for(std::chrono::milliseconds(100));
+
         // Continuously polling cores here can cause other host-driven noc transactions (dprint, watcher) to drastically
         // slow down for remote devices. So when debugging with these features, add a small delay to allow other
         // host-driven transactions through.

tt_metal/llrt/tt_cluster.cpp

@@ -120,7 +120,7 @@ Cluster::Cluster() {
 
 void Cluster::detect_arch_and_target() {
 
-    this->target_type_ = (std::getenv("TT_METAL_SIMULATOR_EN")) ? TargetDevice::Simulator : TargetDevice::Silicon;
+    this->target_type_ = (std::getenv("TT_METAL_SIMULATOR")) ? TargetDevice::Simulator : TargetDevice::Silicon;
 
     this->arch_ = tt_metal::get_platform_architecture();
 
@@ -237,10 +237,9 @@ const std::unordered_map<CoreCoord, int32_t>& Cluster::get_virtual_routing_to_pr
 }
 
 void Cluster::open_driver(const bool &skip_driver_allocs) {
-    const std::string sdesc_path = get_soc_description_file(this->arch_, this->target_type_);
-
     std::unique_ptr<tt_device> device_driver;
     if (this->target_type_ == TargetDevice::Silicon) {
+        const std::string sdesc_path = get_soc_description_file(this->arch_, this->target_type_);
         std::unordered_set<chip_id_t> all_chips = this->cluster_desc_->get_all_chips();
         std::set<chip_id_t> all_chips_set(all_chips.begin(), all_chips.end());
         // This is the target/desired number of mem channels per arch/device.
@@ -268,7 +267,8 @@ void Cluster::open_driver(const bool &skip_driver_allocs) {
         // that is later expected to be populated by unrelated APIs
         // TT_FATAL(device_driver->get_target_mmio_device_ids().size() == 1, "Only one target mmio device id allowed.");
     } else if (this->target_type_ == TargetDevice::Simulator) {
-        device_driver = std::make_unique<tt_SimulationDevice>(sdesc_path);
+        auto simulator_directory = std::getenv("TT_METAL_SIMULATOR");
+        device_driver = std::make_unique<tt_SimulationDevice>(simulator_directory);
     }
 
     barrier_address_params barrier_params;