diff --git a/tests/tt_metal/tt_metal/api/CMakeLists.txt b/tests/tt_metal/tt_metal/api/CMakeLists.txt index fc95afeb92e6..864e4b911e9f 100644 --- 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 new file mode 100644 index 000000000000..8fccfe95aae3 --- /dev/null +++ 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 + +#include + +#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 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 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 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); + } + } +}