#0: Add global circular buffer tests

tenstorrent · Dec 2, 2024 · 218b0f9 · 218b0f9
1 parent 6cc1d32
commit 218b0f9
Show file tree

Hide file tree

Showing 2 changed files with 51 additions and 0 deletions.
diff --git a/tests/tt_metal/tt_metal/api/CMakeLists.txt b/tests/tt_metal/tt_metal/api/CMakeLists.txt
@@ -19,6 +19,7 @@ set(UNIT_TESTS_API_SRC
     ${CMAKE_CURRENT_SOURCE_DIR}/test_direct.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_dram_to_l1_multicast.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_dram.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/test_global_circular_buffers.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_global_semaphores.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_kernel_creation.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/test_noc.cpp

diff --git a/tests/tt_metal/tt_metal/api/test_global_circular_buffers.cpp b/tests/tt_metal/tt_metal/api/test_global_circular_buffers.cpp
@@ -0,0 +1,50 @@
+// SPDX-FileCopyrightText: © 2024 Tenstorrent Inc.
+//
+// SPDX-License-Identifier: Apache-2.0
+
+#include <gtest/gtest.h>
+
+#include <vector>
+
+#include "device_fixture.hpp"
+#include "tt_metal/common/core_coord.hpp"
+#include "tt_metal/detail/tt_metal.hpp"
+#include "tt_metal/host_api.hpp"
+#include "tt_metal/impl/buffers/global_circular_buffer.hpp"
+
+TEST_F(DispatchFixture, CreateGlobalCircularBuffers) {
+    CoreRangeSet cores(CoreRange({1, 1}, {1, 1}));
+    CoreRangeSet cores2(CoreRange({1, 1}, {2, 2}));
+    CoreRangeSet cores3(CoreRange({3, 3}, {3, 3}));
+
+    for (auto device : devices_) {
+        {
+            std::unordered_map<CoreCoord, CoreRangeSet> sender_receiver_core_mapping;
+            sender_receiver_core_mapping[CoreCoord(0, 0)] = cores;
+            auto global_cb = tt::tt_metal::experimental::CreateGlobalCircularBuffer(
+                device, sender_receiver_core_mapping, 3200, tt::tt_metal::BufferType::L1);
+            auto buffer_address = global_cb->buffer_address();
+            auto config_address = global_cb->config_address();
+        }
+        {
+            std::unordered_map<CoreCoord, CoreRangeSet> sender_receiver_core_mapping;
+            sender_receiver_core_mapping[CoreCoord(0, 0)] = cores;
+            sender_receiver_core_mapping[CoreCoord(1, 1)] = cores3;
+            // sender receiver cores overlap
+            EXPECT_THROW(
+                tt::tt_metal::experimental::CreateGlobalCircularBuffer(
+                    device, sender_receiver_core_mapping, 3200, tt::tt_metal::BufferType::L1),
+                std::exception);
+        }
+        {
+            std::unordered_map<CoreCoord, CoreRangeSet> sender_receiver_core_mapping;
+            sender_receiver_core_mapping[CoreCoord(0, 0)] = cores;
+            sender_receiver_core_mapping[CoreCoord(0, 1)] = cores2;
+            // receiver cores overlap
+            EXPECT_THROW(
+                tt::tt_metal::experimental::CreateGlobalCircularBuffer(
+                    device, sender_receiver_core_mapping, 3200, tt::tt_metal::BufferType::L1),
+                std::exception);
+        }
+    }
+}