diff --git a/ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/all_reduce_async_op.hpp b/ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/all_reduce_async_op.hpp index a826ee761b3..dae83a6bbc4 100644 --- a/ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/all_reduce_async_op.hpp +++ b/ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/all_reduce_async_op.hpp @@ -46,7 +46,7 @@ struct AllReduceAsync { MemoryConfig output_mem_config, ccl::Topology topology, GlobalSemaphore semaphore, - std::optional& sub_device_id, + std::optional& sub_device_id, bool enable_persistent_fabric_mode) : forward_device(forward_device), backward_device(backward_device), diff --git a/ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/all_reduce_async_program_minimal_variants.cpp b/ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/all_reduce_async_program_minimal_variants.cpp index 54372f82bc3..119a7b49832 100644 --- a/ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/all_reduce_async_program_minimal_variants.cpp +++ b/ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/all_reduce_async_program_minimal_variants.cpp @@ -121,7 +121,8 @@ tt::tt_metal::operation::ProgramWithCallbacks all_reduce_async_minimal_multi_cor tt::tt_metal::CircularBufferConfig cb_src0_config = tt::tt_metal::CircularBufferConfig(cb_num_pages * l1_scratch_cb_page_size_bytes, {{src0_cb_index, df}}) .set_page_size(src0_cb_index, l1_scratch_cb_page_size_bytes); - CBHandle cb_src0_workers = CreateCircularBuffer(program, sender_worker_core_range, cb_src0_config); + tt::tt_metal::CBHandle cb_src0_workers = + tt::tt_metal::CreateCircularBuffer(program, sender_worker_core_range, cb_src0_config); // Set aside a buffer we can use for storing packet headers in (particularly for atomic incs) const auto reserved_packet_header_CB_index = tt::CBIndex::c_3; static constexpr auto num_packet_headers_storable = 8; @@ -132,7 +133,7 @@ tt::tt_metal::operation::ProgramWithCallbacks all_reduce_async_minimal_multi_cor {{reserved_packet_header_CB_index, tt::DataFormat::RawUInt32}}) .set_page_size(reserved_packet_header_CB_index, packet_header_size_bytes); auto reserved_packet_header_CB_handle = - CreateCircularBuffer(program, sender_worker_core_range, cb_reserved_packet_header_config); + tt::tt_metal::CreateCircularBuffer(program, sender_worker_core_range, cb_reserved_packet_header_config); // Reduction kernel setup auto all_cores = output_tensor_cores.merge(sender_worker_core_range); @@ -282,8 +283,8 @@ tt::tt_metal::operation::ProgramWithCallbacks all_reduce_async_minimal_multi_cor tt::tt_metal::SetRuntimeArgs(program, reduction_kernel_id, output_tensor_cores, reduction_kernel_rt_args); // KERNEL CREATION - tt::tt_metal::NOC reader_noc = NOC::NOC_1; - tt::tt_metal::NOC writer_noc = NOC::NOC_0; + tt::tt_metal::NOC reader_noc = tt::tt_metal::NOC::NOC_1; + tt::tt_metal::NOC writer_noc = tt::tt_metal::NOC::NOC_0; // Reader std::vector reader_compile_args = { ring_index, // my_chip_id @@ -296,7 +297,7 @@ tt::tt_metal::operation::ProgramWithCallbacks all_reduce_async_minimal_multi_cor "ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/kernels/dataflow/" "worker_reader.cpp", sender_worker_core_range, - DataMovementConfig{ + tt::tt_metal::DataMovementConfig{ .processor = DataMovementProcessor::RISCV_1, .noc = reader_noc, .compile_args = reader_compile_args}); // Writer @@ -316,7 +317,7 @@ tt::tt_metal::operation::ProgramWithCallbacks all_reduce_async_minimal_multi_cor "ttnn/cpp/ttnn/operations/experimental/ccl/all_reduce_async/device/kernels/dataflow/" "worker_writer.cpp", sender_worker_core_range, - DataMovementConfig{ + tt::tt_metal::DataMovementConfig{ .processor = DataMovementProcessor::RISCV_0, .noc = writer_noc, .compile_args = writer_compile_args}); // Kernel Runtime Args @@ -388,7 +389,7 @@ tt::tt_metal::operation::ProgramWithCallbacks all_reduce_async_minimal_multi_cor if (mcast_range_contains_self) { num_mcast_cores -= 1; } - if (writer_noc == NOC::NOC_1) { + if (writer_noc == tt::tt_metal::NOC::NOC_1) { std::swap(start_core, end_core); } mcast_start_x.push_back(start_core.x);