Merge branch 'main' into dvartanians/fix_yolov4_faster_webdemo

tenstorrent · Feb 19, 2025 · f34a56d · f34a56d
2 parents feacb38 + 686a4f0
commit f34a56d
Show file tree

Hide file tree

Showing 117 changed files with 1,767 additions and 1,765 deletions.
diff --git a/.github/workflows/build-and-test-wheels.yaml b/.github/workflows/build-and-test-wheels.yaml
@@ -15,6 +15,8 @@ jobs:
     if: ${{ github.event_name == 'workflow_dispatch' && inputs.from-precompiled }}
     uses: ./.github/workflows/build-artifact.yaml
     secrets: inherit
+    with:
+      build-wheel: true
   test-wheels:
     needs: build-artifact
     if: ${{ always() }}

diff --git a/.github/workflows/build-and-unit-tests.yaml b/.github/workflows/build-and-unit-tests.yaml
@@ -58,7 +58,7 @@ jobs:
           {name: api, cmd: "./build/test/tt_metal/unit_tests_api_${{ inputs.arch }}"},
           {name: debug_tools, cmd: "./build/test/tt_metal/unit_tests_debug_tools_${{ inputs.arch }}"},
           {name: device, cmd: "./build/test/tt_metal/unit_tests_device"},
-          {name: dispatch, cmd: "./build/test/tt_metal/unit_tests_dispatch_${{ inputs.arch }}"},
+          {name: dispatch, cmd: "./build/test/tt_metal/unit_tests_dispatch"},
           {name: eth, cmd: "./build/test/tt_metal/unit_tests_eth_${{ inputs.arch }}"},
           {name: llk, cmd: "./build/test/tt_metal/unit_tests_llk"},
           {name: stl, cmd: "./build/test/tt_metal/unit_tests_stl"},

diff --git a/.github/workflows/cpp-post-commit.yaml b/.github/workflows/cpp-post-commit.yaml
@@ -58,13 +58,13 @@ jobs:
           {name: api, cmd: "./build/test/tt_metal/unit_tests_api_${{ inputs.arch }}"},
           {name: debug_tools, cmd: "./build/test/tt_metal/unit_tests_debug_tools_${{ inputs.arch }}"},
           {name: device, cmd: "./build/test/tt_metal/unit_tests_device"},
-          {name: dispatch, cmd: "./build/test/tt_metal/unit_tests_dispatch_${{ inputs.arch }}"},
+          {name: dispatch, cmd: "./build/test/tt_metal/unit_tests_dispatch"},
           {name: eth, cmd: "./build/test/tt_metal/unit_tests_eth_${{ inputs.arch }}"},
           {name: llk, cmd: "./build/test/tt_metal/unit_tests_llk"},
           {name: stl, cmd: "./build/test/tt_metal/unit_tests_stl"},
           {name: distributed, cmd: "./build/test/tt_metal/distributed/distributed_unit_tests_${{ inputs.arch }}"},
           {name: lightmetal, cmd: "./build/test/tt_metal/unit_tests_lightmetal"},
-          {name: dispatch multicmd queue, cmd: "TT_METAL_GTEST_NUM_HW_CQS=2 ./build/test/tt_metal/unit_tests_dispatch_${{ inputs.arch }} --gtest_filter=MultiCommandQueue*Fixture.*"},
+          {name: dispatch multicmd queue, cmd: "TT_METAL_GTEST_NUM_HW_CQS=2 ./build/test/tt_metal/unit_tests_dispatch --gtest_filter=MultiCommandQueue*Fixture.*"},
           {name: ttnn cpp unit tests, cmd: ./build/test/ttnn/unit_tests_ttnn},
           {name: ttnn ccl cpp unit tests, cmd: ./build/test/ttnn/unit_tests_ttnn_ccl},
           {name: ttnn tensor cpp unit tests, cmd: ./build/test/ttnn/unit_tests_ttnn_tensor},

diff --git a/.github/workflows/package-and-release.yaml b/.github/workflows/package-and-release.yaml
@@ -13,10 +13,20 @@ permissions:
 jobs:
   build-artifact:
     needs: create-tag
+    strategy:
+      matrix:
+        config:
+          - version: "20.04"
+            publish-artifact: true
+          - version: "22.04"
+            publish-artifact: false
     uses: ./.github/workflows/build-artifact.yaml
-    secrets: inherit
     with:
+      version: ${{ matrix.config.version }}
+      distro: ubuntu
+      publish-artifact: ${{ matrix.config.publish-artifact }}
       build-wheel: true
+    secrets: inherit
   build-artifact-profiler:
     uses: ./.github/workflows/build-artifact.yaml
     with:
@@ -123,14 +133,7 @@ jobs:
           path: RELEASE_NOTES.txt
   # Candidate for breaking up
   create-and-upload-draft-release:
-    needs: [
-      create-tag,
-      create-release-notes,
-      build-artifact,
-    ]
-    strategy:
-      matrix:
-        os: [ubuntu-20.04]
+    needs: [create-tag, create-release-notes, build-artifact]
     # May accidentally create two releases without restricting to 1 job
     concurrency: create_upload_draft_release
     runs-on: ubuntu-latest
@@ -143,10 +146,14 @@ jobs:
         uses: qmonnet/git-archive-all-action@791fb850881cf58b1d1fcc9b06c01940080bba0a
         with:
           output-files: tt-metalium.tar.gz
-      - name: Download eager Python packages
+      - name: Download eager 20.04 Python packages
+        uses: actions/download-artifact@v4
+        with:
+          name: eager-dist-ubuntu-20.04-any
+      - name: Download eager 22.04 Python packages
         uses: actions/download-artifact@v4
         with:
-          name: eager-dist-${{ matrix.os }}-any
+          name: eager-dist-ubuntu-22.04-any
       - name: Create VERSION
         run: echo ${{ needs.create-tag.outputs.version }} > VERSION
       - name : Download release notes

diff --git a/METALIUM_GUIDE.md b/METALIUM_GUIDE.md
@@ -125,7 +125,7 @@ kernel:
 ```
 namespace NAMESPACE {
 void MAIN {
-  mm_init();
+  mm_init(tt::CBIndex::c_0, tt::CBIndex::c_1, tt::CBIndex::c_16);
   acquire_dst();
 
   cb_wait_front(tt::CBIndex::c_0, /* number of tiles */ 1);
@@ -297,7 +297,7 @@ with `tile_regs_..()` functions like:
 ```
 namespace NAMESPACE {
 void MAIN {
-  mm_init();
+  mm_init(tt::CBIndex::c_0, tt::CBIndex::c_1, tt::CBIndex::c_16);
 
   cb_wait_front(tt::CBIndex::c_0, /* number of tiles */ 1);
   cb_wait_front(tt::CBIndex::c_1, /* number of tiles */ 1);

diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/add_tiles.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/add_tiles.rst
@@ -3,5 +3,5 @@ add_tiles
 
 
 .. doxygenfunction:: add_tiles_init_nof()
-.. doxygenfunction:: add_tiles_init(uint32_t icb0 = 0, uint32_t icb1 = 1, bool acc_to_dest = false)
+.. doxygenfunction:: add_tiles_init(uint32_t icb0, uint32_t icb1, bool acc_to_dest = false)
 .. doxygenfunction:: add_tiles(uint32_t icb0, uint32_t icb1, uint32_t itile0, uint32_t itile1, uint32_t idst)
diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/add_tiles_bcast.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/add_tiles_bcast.rst
@@ -1,6 +1,6 @@
 add_tiles_bcast
 ===============
 
-.. doxygenfunction:: add_bcast_cols_init_short(uint32_t icb0 = 0, uint32_t icb1 = 1)
-.. doxygenfunction:: add_bcast_rows_init_short(uint32_t icb0 = 0, uint32_t icb1 = 1)
+.. doxygenfunction:: add_bcast_cols_init_short(uint32_t icb0, uint32_t icb1)
+.. doxygenfunction:: add_bcast_rows_init_short(uint32_t icb0, uint32_t icb1)
 .. doxygenfunction:: add_tiles_bcast(uint32_t icb0, uint32_t icb1, uint32_t itile0, uint32_t itile1, uint32_t idst)
diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/matmul_block.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/matmul_block.rst
@@ -1,7 +1,7 @@
 matmul_block
 ============
 
-.. doxygenfunction:: mm_block_init(uint32_t in0_cb_id = 0, uint32_t in1_cb_id = 1, uint32_t out_cb_id = 16, const uint32_t transpose=0, uint32_t ct_dim = 1, uint32_t rt_dim = 1, uint32_t kt_dim = 1)
-.. doxygenfunction:: mm_block_init_short(uint32_t in0_cb_id = 0, uint32_t in1_cb_id = 1, const uint32_t transpose=0, uint32_t ct_dim = 1, uint32_t rt_dim = 1, uint32_t kt_dim = 1)
-.. doxygenfunction:: mm_block_init_short_with_dt(uint32_t in0_cb_id = 0, uint32_t in1_cb_id = 1, uint32_t old_in1_cb_id=2, const uint32_t transpose=0, uint32_t ct_dim = 1, uint32_t rt_dim = 1, uint32_t kt_dim = 1)
+.. doxygenfunction:: mm_block_init(uint32_t in0_cb_id, uint32_t in1_cb_id, uint32_t out_cb_id, const uint32_t transpose=0, uint32_t ct_dim = 1, uint32_t rt_dim = 1, uint32_t kt_dim = 1)
+.. doxygenfunction:: mm_block_init_short(uint32_t in0_cb_id, uint32_t in1_cb_id, const uint32_t transpose=0, uint32_t ct_dim = 1, uint32_t rt_dim = 1, uint32_t kt_dim = 1)
+.. doxygenfunction:: mm_block_init_short_with_dt(uint32_t in0_cb_id, uint32_t in1_cb_id, uint32_t old_in1_cb_id, const uint32_t transpose=0, uint32_t ct_dim = 1, uint32_t rt_dim = 1, uint32_t kt_dim = 1)
 .. doxygenfunction:: matmul_block(uint32_t in0_cb_id, uint32_t in1_cb_id, uint32_t in0_tile_index, uint32_t in1_tile_index, uint32_t idst, const uint32_t transpose, uint32_t ct_dim, uint32_t rt_dim, uint32_t kt_dim)
diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/matmul_tiles.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/matmul_tiles.rst
@@ -1,7 +1,7 @@
 matmul_tiles
 ============
 
-.. doxygenfunction:: mm_init(uint32_t in0_cb_id = 0, uint32_t in1_cb_id = 1, uint32_t out_cb_id = 16, const uint32_t transpose=0)
-.. doxygenfunction:: mm_init_short_with_dt(uint32_t in0_cb_id = 0, uint32_t in1_cb_id = 1, uint32_t c_in_old_srca = 2, const uint32_t transpose=0)
-.. doxygenfunction:: mm_init_short(uint32_t in0_cb_id = 0, uint32_t in1_cb_id = 1, const uint32_t transpose=0)
+.. doxygenfunction:: mm_init(uint32_t in0_cb_id, uint32_t in1_cb_id, uint32_t out_cb_id, const uint32_t transpose=0)
+.. doxygenfunction:: mm_init_short_with_dt(uint32_t in0_cb_id, uint32_t in1_cb_id, uint32_t c_in_old_srca, const uint32_t transpose=0)
+.. doxygenfunction:: mm_init_short(uint32_t in0_cb_id, uint32_t in1_cb_id, const uint32_t transpose=0)
 .. doxygenfunction:: matmul_tiles(uint32_t in0_cb_id, uint32_t in1_cb_id, uint32_t in0_tile_index, uint32_t in1_tile_index, uint32_t idst, const uint32_t transpose)
diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/move_copy_tile.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/move_copy_tile.rst
@@ -3,5 +3,5 @@ move_copy_tile
 
 
 .. doxygenfunction:: copy_tile_to_dst_init_short_with_dt(uint32_t old_cbid, uint32_t new_cbid, uint32_t transpose = 0)
-.. doxygenfunction:: copy_tile_to_dst_init_short(uint32_t cbid = 0, uint32_t transpose = 0)
+.. doxygenfunction:: copy_tile_to_dst_init_short(uint32_t cbid, uint32_t transpose = 0)
 .. doxygenfunction:: copy_tile_init(uint32_t cbid)
diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/mul_tiles.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/mul_tiles.rst
@@ -2,5 +2,5 @@ mul_tiles
 =========
 
 .. doxygenfunction:: mul_tiles_init_f()
-.. doxygenfunction:: mul_tiles_init(uint32_t icb0 = 0, uint32_t icb1 = 1)
+.. doxygenfunction:: mul_tiles_init(uint32_t icb0, uint32_t icb1)
 .. doxygenfunction:: mul_tiles(uint32_t icb0, uint32_t icb1, uint32_t itile0, uint32_t itile1, uint32_t idst)
diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/mul_tiles_bcast.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/mul_tiles_bcast.rst
@@ -1,8 +1,8 @@
 mul_tiles_bcast
 ===============
 
-.. doxygenfunction:: mul_bcast_cols_init_short(uint32_t icb0 = 0, uint32_t icb1 = 1)
-.. doxygenfunction:: mul_bcast_rows_init_short(uint32_t icb0 = 0, uint32_t icb1 = 1)
+.. doxygenfunction:: mul_bcast_cols_init_short(uint32_t icb0, uint32_t icb1)
+.. doxygenfunction:: mul_bcast_rows_init_short(uint32_t icb0, uint32_t icb1)
 .. doxygenfunction:: mul_tiles_bcast(uint32_t icb0, uint32_t icb1, uint32_t itile0, uint32_t itile1, uint32_t idst)
-.. doxygenfunction:: mul_tiles_bcast_scalar_init_short(uint32_t icb0 = 0, uint32_t icb1 = 1)
+.. doxygenfunction:: mul_tiles_bcast_scalar_init_short(uint32_t icb0, uint32_t icb1)
 .. doxygenfunction:: mul_tiles_bcast_scalar(uint32_t icb0, uint32_t icb1, uint32_t itile0, uint32_t itile1, uint32_t idst)
diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/sub_tiles.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/sub_tiles.rst
@@ -2,5 +2,5 @@ sub_tiles
 =========
 
 .. doxygenfunction:: sub_tiles_init_nof()
-.. doxygenfunction:: sub_tiles_init(uint32_t icb0 = 0, uint32_t icb1 = 1, bool acc_to_dest = false)
+.. doxygenfunction:: sub_tiles_init(uint32_t icb0, uint32_t icb1, bool acc_to_dest = false)
 .. doxygenfunction:: sub_tiles( uint32_t icb0, uint32_t icb1, uint32_t itile0, uint32_t itile1, uint32_t idst)
diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/sub_tiles_bcast.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/sub_tiles_bcast.rst
@@ -2,5 +2,5 @@ sub_tiles_bcast
 ===============
 
 
-.. doxygenfunction:: sub_bcast_cols_init_short(uint32_t icb0 = 0, uint32_t icb1 = 1)
+.. doxygenfunction:: sub_bcast_cols_init_short(uint32_t icb0, uint32_t icb1)
 .. doxygenfunction:: sub_tiles_bcast(uint32_t icb0, uint32_t icb1, uint32_t itile0, uint32_t itile1, uint32_t idst)
diff --git a/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/transpose_wh_tile.rst b/docs/source/tt-metalium/tt_metal/apis/kernel_apis/compute/transpose_wh_tile.rst
@@ -1,5 +1,5 @@
 transpose_wh_tile
 =================
 
-.. doxygenfunction:: transpose_wh_init(uint32_t icb, uint32_t ocb = 16)
+.. doxygenfunction:: transpose_wh_init(uint32_t icb, uint32_t ocb)
 .. doxygenfunction:: transpose_wh_tile(uint32_t icb, uint32_t itile, uint32_t idst)
diff --git a/models/demos/t3000/falcon40b/tests/test_falcon_model.py b/models/demos/t3000/falcon40b/tests/test_falcon_model.py
@@ -6,7 +6,7 @@
 import pytest
 from loguru import logger
 import ttnn
-from ttnn import ShardTensorToMesh, ConcatMeshToTensor, ListMeshToTensor
+from ttnn import ShardTensorToMesh, ConcatMeshToTensor
 from models.demos.t3000.falcon40b.reference.hf_modeling_falcon import (
     FalconForCausalLM,
 )
@@ -196,7 +196,7 @@ def run_test_FalconModel_inference(
                 use_cache=use_cache,
             )
             # output of model is replicated
-            tensors = ttnn.to_torch(tt_out, device=mesh_device, mesh_composer=ListMeshToTensor(mesh_device))
+            tensors = [ttnn.to_torch(shard) for shard in ttnn.get_device_tensors(tt_out.cpu())]
             tt_outs.append(tensors[0].squeeze(1))
 
         tt_out = torch.vstack(tt_outs)
@@ -213,7 +213,7 @@ def run_test_FalconModel_inference(
             use_cache=use_cache,
         )
         # Output of model is replicated
-        tensors = ttnn.to_torch(tt_out, device=mesh_device, mesh_composer=ListMeshToTensor(mesh_device))
+        tensors = [ttnn.to_torch(shard) for shard in ttnn.get_device_tensors(tt_out.cpu())]
         tt_out = tensors[0].squeeze(1).transpose(0, 1)
 
     # check outputs ----------------------------------------------------------------------

diff --git a/models/demos/t3000/llama2_70b/tests/test_llama_generation.py b/models/demos/t3000/llama2_70b/tests/test_llama_generation.py
@@ -6,7 +6,7 @@
 import torch
 from torch import nn
 import ttnn
-from ttnn import ShardTensorToMesh, ReplicateTensorToMesh, ConcatMeshToTensor, ListMeshToTensor
+from ttnn import ShardTensorToMesh, ReplicateTensorToMesh, ConcatMeshToTensor
 
 
 import scipy

diff --git a/models/demos/tg/llama3_70b/tests/test_llama_attention_galaxy.py b/models/demos/tg/llama3_70b/tests/test_llama_attention_galaxy.py
@@ -337,7 +337,7 @@ def run_test_LlamaAttention_inference(
             attn_mask,
             mode=mode,
         )
-        # tt_out = ttnn.to_torch(tt_out, mesh_composer=ListMeshToTensor(mesh_device))[0]
+        # tt_out = [ttnn.to_torch(shard) for shard in ttnn.get_device_tensors(tt_out.cpu())]
 
         tt_out = ttnn.to_torch(
             tt_out, mesh_composer=ConcatMesh2DToTensor(mesh_device, dims=(3, 1), cluster_shape=cluster_shape)

diff --git a/tests/scripts/run_cpp_unit_tests.sh b/tests/scripts/run_cpp_unit_tests.sh
@@ -18,7 +18,7 @@ if [[ ! -z "$TT_METAL_SLOW_DISPATCH_MODE" ]]; then
 else
     # Enable this on BH after #14613
     if [[ "$ARCH_NAME" == "wormhole_b0" ]]; then
-        TT_METAL_GTEST_ETH_DISPATCH=1 ./build/test/tt_metal/unit_tests_dispatch_${ARCH_NAME}
+        TT_METAL_GTEST_ETH_DISPATCH=1 ./build/test/tt_metal/unit_tests_dispatch
     fi
     env python3 tests/scripts/run_tt_eager.py --dispatch-mode fast
     env python3 tests/scripts/run_tt_metal.py --dispatch-mode fast

diff --git a/tests/scripts/run_tests.sh b/tests/scripts/run_tests.sh
@@ -79,7 +79,7 @@ run_frequent_api_pipeline_tests() {
     local dispatch_mode=$3
 
     if [[ $dispatch_mode == "slow" ]]; then
-        TT_METAL_SLOW_DISPATCH_MODE=1 ./build/test/tt_metal/unit_tests_dispatch_${ARCH_NAME} --gtest_filter=DispatchStress.TensixRunManyTimes
+        TT_METAL_SLOW_DISPATCH_MODE=1 ./build/test/tt_metal/unit_tests_dispatch --gtest_filter=DispatchStress.TensixRunManyTimes
         echo "Running Python API unit tests in SD for frequent..."
         ./tests/scripts/run_python_api_unit_tests.sh
     fi

diff --git a/tests/scripts/t3000/run_t3000_unit_tests.sh b/tests/scripts/t3000/run_t3000_unit_tests.sh
@@ -18,8 +18,8 @@ run_t3000_ttmetal_tests() {
   TT_METAL_SLOW_DISPATCH_MODE=1 ./build/test/tt_metal/unit_tests_eth_${ARCH_NAME} --gtest_filter="DeviceFixture.ActiveEthKernelsSendInterleavedBufferAllConnectedChips" ; fail+=$?
   TT_METAL_SLOW_DISPATCH_MODE=1 ./build/test/tt_metal/unit_tests_eth_${ARCH_NAME} --gtest_filter="DeviceFixture.ActiveEthKernelsDirectRingGatherAllChips" ; fail+=$?
   TT_METAL_SLOW_DISPATCH_MODE=1 ./build/test/tt_metal/unit_tests_eth_${ARCH_NAME} --gtest_filter="DeviceFixture.ActiveEthKernelsInterleavedRingGatherAllChips" ; fail+=$?
-  TT_METAL_ENABLE_REMOTE_CHIP=1 ./build/test/tt_metal/unit_tests_dispatch_${ARCH_NAME} --gtest_filter="CommandQueueSingleCard*Fixture.*" ; fail+=$?
-  ./build/test/tt_metal/unit_tests_dispatch_${ARCH_NAME} --gtest_filter="CommandQueueMultiDevice*Fixture.*" ; fail+=$?
+  TT_METAL_ENABLE_REMOTE_CHIP=1 ./build/test/tt_metal/unit_tests_dispatch --gtest_filter="CommandQueueSingleCard*Fixture.*" ; fail+=$?
+  ./build/test/tt_metal/unit_tests_dispatch --gtest_filter="CommandQueueMultiDevice*Fixture.*" ; fail+=$?
   ./build/test/tt_metal/unit_tests_debug_tools_${ARCH_NAME} --gtest_filter="DPrintFixture.*:WatcherFixture.*" ; fail+=$?
 
   # Programming examples

diff --git a/tests/scripts/tg/run_tg_unit_tests.sh b/tests/scripts/tg/run_tg_unit_tests.sh
@@ -105,11 +105,11 @@ run_tg_prefetcher_tests() {
 run_tg_tests() {
   if [[ "$1" == "unit" ]]; then
     echo "LOG_METAL: running run_tg_unit_tests"
-    TT_METAL_ENABLE_REMOTE_CHIP=1 ./build/test/tt_metal/unit_tests_dispatch_${ARCH_NAME} --gtest_filter="CommandQueueSingleCard*Fixture.*"
+    TT_METAL_ENABLE_REMOTE_CHIP=1 ./build/test/tt_metal/unit_tests_dispatch --gtest_filter="CommandQueueSingleCard*Fixture.*"
     ./build/test/ttnn/galaxy_unit_tests_ttnn
     TT_METAL_SLOW_DISPATCH_MODE=1 ./build/test/tt_metal/unit_tests_device --gtest_filter="GalaxyFixture.*:TGFixture.*"
     ./build/test/tt_metal/unit_tests_device --gtest_filter="GalaxyFixture.*:TGFixture.*"
-    TT_METAL_GTEST_NUM_HW_CQS=2 ./build/test/tt_metal/unit_tests_dispatch_${ARCH_NAME} --gtest_filter="MultiCommandQueueMultiDevice*Fixture.*"
+    TT_METAL_GTEST_NUM_HW_CQS=2 ./build/test/tt_metal/unit_tests_dispatch --gtest_filter="MultiCommandQueueMultiDevice*Fixture.*"
 
   elif [[ "$1" == "fabric" ]]; then
     echo "LOG_FABRIC: running run_tg_fabric_tests"

diff --git a/tests/scripts/tgg/run_tgg_unit_tests.sh b/tests/scripts/tgg/run_tgg_unit_tests.sh
@@ -5,7 +5,7 @@ run_tgg_tests() {
 
   echo "LOG_METAL: running run_tgg_unit_tests"
 
-  TT_METAL_ENABLE_REMOTE_CHIP=1 ./build/test/tt_metal/unit_tests_dispatch_${ARCH_NAME} --gtest_filter="CommandQueueSingleCard*Fixture.*"
+  TT_METAL_ENABLE_REMOTE_CHIP=1 ./build/test/tt_metal/unit_tests_dispatch --gtest_filter="CommandQueueSingleCard*Fixture.*"
   ./build/test/ttnn/galaxy_unit_tests_ttnn
   TT_METAL_SLOW_DISPATCH_MODE=1 ./build/test/tt_metal/unit_tests_device --gtest_filter="GalaxyFixture.*:TGGFixture.*"
   ./build/test/tt_metal/unit_tests_device --gtest_filter="GalaxyFixture.*:TGGFixture.*"

diff --git a/tests/tt_metal/distributed/test_mesh_coord.cpp b/tests/tt_metal/distributed/test_mesh_coord.cpp
@@ -4,14 +4,15 @@
 
 #include <gtest/gtest.h>
 #include <gmock/gmock.h>
+#include <unordered_set>
 
 #include "mesh_coord.hpp"
 
 namespace tt::tt_metal::distributed {
 namespace {
 
 using ::testing::ElementsAre;
-
+using ::testing::UnorderedElementsAre;
 TEST(SimpleMeshShapeTest, Construction) {
     SimpleMeshShape shape_1d(3);
     EXPECT_EQ(shape_1d.dims(), 1);
@@ -100,6 +101,21 @@ TEST(MeshCoordinateTest, Comparison) {
     EXPECT_NE(coord1, MeshCoordinate(1, 2, 1));
 }
 
+TEST(MeshCoordinateTest, UnorderedSet) {
+    std::unordered_set<MeshCoordinate> set;
+    set.insert(MeshCoordinate(0, 0, 0));
+    set.insert(MeshCoordinate(0, 0, 1));
+    set.insert(MeshCoordinate(0, 0, 2));
+
+    EXPECT_FALSE(set.insert(MeshCoordinate(0, 0, 2)).second);
+    EXPECT_THAT(
+        set,
+        UnorderedElementsAre(
+            MeshCoordinate(0, 0, 0),  //
+            MeshCoordinate(0, 0, 1),
+            MeshCoordinate(0, 0, 2)));
+}
+
 TEST(MeshCoordinateRangeTest, FromShape) {
     SimpleMeshShape shape(2, 3);
     MeshCoordinateRange range(shape);
@@ -232,6 +248,7 @@ TEST(MeshContainerTest, ElementAccessRowMajor) {
             MeshCoordinate(1, 1),
             MeshCoordinate(1, 2)));
     EXPECT_THAT(values, ElementsAre(0, 1, 2, 3, 4, 5));
+    EXPECT_THAT(container.values(), ElementsAre(0, 1, 2, 3, 4, 5));
 }
 
 TEST(MeshContainerTest, ConstContainer) {
@@ -254,6 +271,7 @@ TEST(MeshContainerTest, ConstContainer) {
             MeshCoordinate(1, 1),
             MeshCoordinate(1, 2)));
     EXPECT_THAT(values, ElementsAre(0, 0, 0, 0, 0, 0));
+    EXPECT_THAT(container.values(), ElementsAre(0, 0, 0, 0, 0, 0));
 }
 
 TEST(MeshContainerTest, MutateThroughProxy) {
@@ -276,6 +294,7 @@ TEST(MeshContainerTest, MutateThroughProxy) {
         values.push_back(value);
     }
     EXPECT_THAT(values, ElementsAre(0, 1, 2, 3, 4, 5));
+    EXPECT_THAT(container.values(), ElementsAre(0, 1, 2, 3, 4, 5));
 }
 
 TEST(MeshContainerTest, OutOfBounds) {