Skip to content

Commit

Permalink
#11592: Modified calls to CreateSemaphore to use the return values
Browse files Browse the repository at this point in the history
  • Loading branch information
@sagarwalTT
sagarwalTT committed Sep 30, 2024
1 parent fbbfd60 commit a935481
Show file tree
Hide file tree
Showing 3 changed files with 30 additions and 26 deletions.
12 changes: 7 additions & 5 deletions tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_dispatcher.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<uint32_t> dispatch_compile_args =
{l1_buf_base,
Expand Down
13 changes: 8 additions & 5 deletions tests/tt_metal/tt_metal/unit_tests/fast_dispatch_kernels/test_write_host.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -127,11 +127,14 @@ bool test_write_host(Device *device, uint32_t data_size, std::pair<uint32_t, uin
tt::llrt::write_hex_vec_to_core(device->id(), 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<uint32_t> dispatch_compile_args = {
l1_buf_base,
Expand Down
31 changes: 15 additions & 16 deletions tt_metal/impl/device/device.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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<uint32_t> prefetch_compile_args = {
dispatch_constants::DISPATCH_BUFFER_BASE,
dispatch_constants::DISPATCH_BUFFER_LOG_PAGE_SIZE,
Expand Down Expand Up @@ -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<uint32_t> dispatch_compile_args = {
dispatch_constants::DISPATCH_BUFFER_BASE,
dispatch_constants::DISPATCH_BUFFER_LOG_PAGE_SIZE,
Expand Down Expand Up @@ -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));
Expand Down

0 comments on commit a935481

Please sign in to comment.