From a935481a144b537b7ccc728744d529d1d7fe3c09 Mon Sep 17 00:00:00 2001 From: Samarth Agarwal Date: Mon, 30 Sep 2024 17:10:02 +0000 Subject: [PATCH] #11592: Modified calls to CreateSemaphore to use the return values --- .../dispatch/test_dispatcher.cpp | 12 ++++--- .../fast_dispatch_kernels/test_write_host.cpp | 13 +++++--- tt_metal/impl/device/device.cpp | 31 +++++++++---------- 3 files changed, 30 insertions(+), 26 deletions(-) diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_dispatcher.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_dispatcher.cpp index 966bb7b909b1..ae72de232139 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_dispatcher.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_dispatcher.cpp @@ -441,11 +441,13 @@ int main(int argc, char **argv) { gen_cmds(device, cmds, all_workers_g, device_data, dispatch_buffer_page_size_g); llrt::write_hex_vec_to_core(device->id(), phys_spoof_prefetch_core, cmds, l1_buf_base); - constexpr uint32_t dispatch_cb_sem = 0; - constexpr uint32_t prefetch_sync_sem = 1; - tt_metal::CreateSemaphore(program, {spoof_prefetch_core}, dispatch_buffer_pages); - tt_metal::CreateSemaphore(program, {dispatch_core}, 0); - tt_metal::CreateSemaphore(program, {spoof_prefetch_core}, 0); + const uint32_t spoof_prefetch_core_sem_0_id = tt_metal::CreateSemaphore(program, {spoof_prefetch_core}, dispatch_buffer_pages); + const uint32_t dispatch_core_sem_id = tt_metal::CreateSemaphore(program, {dispatch_core}, 0); + TT_ASSERT(spoof_prefetch_core_sem_0_id == dispatch_core_sem_id); + const uint32_t dispatch_cb_sem = spoof_prefetch_core_sem_0_id; + + const uint32_t spoof_prefetch_core_sem_1_id = tt_metal::CreateSemaphore(program, {spoof_prefetch_core}, 0); + const uint32_t prefetch_sync_sem = spoof_prefetch_core_sem_1_id; std::vector dispatch_compile_args = {l1_buf_base, diff --git a/tests/tt_metal/tt_metal/unit_tests/fast_dispatch_kernels/test_write_host.cpp b/tests/tt_metal/tt_metal/unit_tests/fast_dispatch_kernels/test_write_host.cpp index 1a777c03db5b..b3325a5aac33 100644 --- a/tests/tt_metal/tt_metal/unit_tests/fast_dispatch_kernels/test_write_host.cpp +++ b/tests/tt_metal/tt_metal/unit_tests/fast_dispatch_kernels/test_write_host.cpp @@ -127,11 +127,14 @@ bool test_write_host(Device *device, uint32_t data_size, std::pairid(), phys_spoof_prefetch_core, dispatch_cmds, l1_buf_base); tt::Cluster::instance().l1_barrier(device->id()); - constexpr uint32_t dispatch_cb_sem = 0; - constexpr uint32_t prefetch_sync_sem = 1; - tt::tt_metal::CreateSemaphore(program, {spoof_prefetch_core}, dispatch_buffer_pages); - tt::tt_metal::CreateSemaphore(program, {dispatch_core}, 0); - tt::tt_metal::CreateSemaphore(program, {spoof_prefetch_core}, 0); + const uint32_t spoof_prefetch_core_sem_0_id = + tt::tt_metal::CreateSemaphore(program, {spoof_prefetch_core}, dispatch_buffer_pages); + const uint32_t dispatch_core_sem_id = tt::tt_metal::CreateSemaphore(program, {dispatch_core}, 0); + TT_ASSERT(spoof_prefetch_core_sem_0_id == dispatch_core_sem_id); + const uint32_t dispatch_cb_sem = spoof_prefetch_core_sem_0_id; + + const uint32_t spoof_prefetch_core_sem_1_id = tt::tt_metal::CreateSemaphore(program, {spoof_prefetch_core}, 0); + const uint32_t prefetch_sync_sem = spoof_prefetch_core_sem_1_id; std::vector dispatch_compile_args = { l1_buf_base, diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp index 05c6793edfa3..474425a976e1 100644 --- a/tt_metal/impl/device/device.cpp +++ b/tt_metal/impl/device/device.cpp @@ -1459,17 +1459,17 @@ void Device::compile_command_queue_programs() { std::string dispatch_kernel_path = "tt_metal/impl/dispatch/kernels/cq_dispatch.cpp"; // TODO: These are semaphore IDs, remove these when CreateSemaphore returns ID rather than address - constexpr uint32_t prefetch_sync_sem = 0; - constexpr uint32_t prefetch_downstream_cb_sem = 1; - constexpr uint32_t prefetch_sem = 1; - constexpr uint32_t dispatch_sem = 0; - constexpr uint32_t mux_sem = 0; - constexpr uint32_t demux_sem = 0; - - constexpr uint32_t prefetch_d_sync_sem = 0; - constexpr uint32_t prefetch_d_upstream_cb_sem = 1; - constexpr uint32_t prefetch_d_downstream_cb_sem = 2; - constexpr uint32_t dispatch_downstream_cb_sem = 1; + // constexpr uint32_t prefetch_sync_sem = 0; + // constexpr uint32_t prefetch_downstream_cb_sem = 1; + // constexpr uint32_t prefetch_sem = 1; + // constexpr uint32_t dispatch_sem = 0; + // constexpr uint32_t mux_sem = 0; + // constexpr uint32_t demux_sem = 0; + + // constexpr uint32_t prefetch_d_sync_sem = 0; + // constexpr uint32_t prefetch_d_upstream_cb_sem = 1; + // constexpr uint32_t prefetch_d_downstream_cb_sem = 2; + // constexpr uint32_t dispatch_downstream_cb_sem = 1; // TODO: this->hw_command_queues_[cq_id]->noc_index is also hardcoded to NOC_0 elsewhere, should have one definition and remove assertion constexpr NOC my_noc_index = NOC::NOC_0; @@ -1504,6 +1504,10 @@ void Device::compile_command_queue_programs() { uint32_t completion_queue_start_addr = issue_queue_start_addr + issue_queue_size; uint32_t completion_queue_size = this->sysmem_manager_->get_completion_queue_size(cq_id); + const uint32_t prefetch_sync_sem = tt::tt_metal::CreateSemaphore(*command_queue_program_ptr, prefetch_core, 0, dispatch_core_type); // prefetch_sync_sem + const uint32_t prefetch_sem = tt::tt_metal::CreateSemaphore(*command_queue_program_ptr, prefetch_core, dispatch_constants::get(dispatch_core_type).dispatch_buffer_pages(), dispatch_core_type); // prefetch_sem + const uint32_t dispatch_sem = tt::tt_metal::CreateSemaphore(*command_queue_program_ptr, dispatch_core, 0, dispatch_core_type); // dispatch_sem + std::vector prefetch_compile_args = { dispatch_constants::DISPATCH_BUFFER_BASE, dispatch_constants::DISPATCH_BUFFER_LOG_PAGE_SIZE, @@ -1545,9 +1549,6 @@ void Device::compile_command_queue_programs() { my_noc_index ); - tt::tt_metal::CreateSemaphore(*command_queue_program_ptr, prefetch_core, 0, dispatch_core_type); // prefetch_sync_sem - tt::tt_metal::CreateSemaphore(*command_queue_program_ptr, prefetch_core, dispatch_constants::get(dispatch_core_type).dispatch_buffer_pages(), dispatch_core_type); // prefetch_sem - std::vector dispatch_compile_args = { dispatch_constants::DISPATCH_BUFFER_BASE, dispatch_constants::DISPATCH_BUFFER_LOG_PAGE_SIZE, @@ -1587,8 +1588,6 @@ void Device::compile_command_queue_programs() { dispatch_upstream_noc_index, my_noc_index ); - - tt::tt_metal::CreateSemaphore(*command_queue_program_ptr, dispatch_core, 0, dispatch_core_type); // dispatch_sem } detail::CompileProgram(this, *command_queue_program_ptr, /*fd_bootloader_mode=*/true); this->command_queue_programs.push_back(std::move(command_queue_program_ptr));