Remove more namespace pollution caused by `using namespace tt::tt_met…

…al` in header file (#16342)

tt_metal/impl/dispatch/kernel_config/demux.cpp

@@ -8,6 +8,8 @@
 #include "tt_metal/host_api.hpp"
 #include "tt_metal/detail/tt_metal.hpp"
 
+using namespace tt::tt_metal;
+
 void DemuxKernel::GenerateStaticConfigs() {
     uint16_t channel =
         tt::Cluster::instance().get_assigned_channel_for_device(servicing_device_id_);  // TODO: this can be mmio

tt_metal/impl/dispatch/kernel_config/dispatch.cpp

@@ -10,6 +10,8 @@
 #include "tt_metal/host_api.hpp"
 #include "tt_metal/detail/tt_metal.hpp"
 
+using namespace tt::tt_metal;
+
 void DispatchKernel::GenerateStaticConfigs() {
     uint16_t channel = tt::Cluster::instance().get_assigned_channel_for_device(device_->id());
     uint8_t cq_id_ = this->cq_id_;

tt_metal/impl/dispatch/kernel_config/dispatch_s.cpp

@@ -8,6 +8,8 @@
 #include "tt_metal/host_api.hpp"
 #include "tt_metal/detail/tt_metal.hpp"
 
+using namespace tt::tt_metal;
+
 void DispatchSKernel::GenerateStaticConfigs() {
     uint16_t channel = tt::Cluster::instance().get_assigned_channel_for_device(device_->id());
     uint8_t cq_id_ = this->cq_id_;

tt_metal/impl/dispatch/kernel_config/eth_router.cpp

@@ -8,6 +8,8 @@
 #include "tt_metal/host_api.hpp"
 #include "tt_metal/detail/tt_metal.hpp"
 
+using namespace tt::tt_metal;
+
 void EthRouterKernel::GenerateStaticConfigs() {
     auto& my_dispatch_constants = dispatch_constants::get(GetCoreType());
     if (as_mux_) {

tt_metal/impl/dispatch/kernel_config/eth_tunneler.cpp

@@ -9,6 +9,8 @@
 #include "tt_metal/host_api.hpp"
 #include "tt_metal/detail/tt_metal.hpp"
 
+using namespace tt::tt_metal;
+
 void EthTunnelerKernel::GenerateStaticConfigs() {
     chip_id_t downstream_device_id = FDKernel::GetDownstreamDeviceId(device_id_);
     // For MMIO devices, the above function just gets one of the possible downstream devices, we've populated this

tt_metal/impl/dispatch/kernel_config/fd_kernel.cpp

@@ -15,6 +15,8 @@
 #include "eth_router.hpp"
 #include "eth_tunneler.hpp"
 
+using namespace tt::tt_metal;
+
 // Helper function to get upstream device in the tunnel from current device, not valid for mmio
 chip_id_t FDKernel::GetUpstreamDeviceId(chip_id_t device_id) {
     chip_id_t mmio_device_id = tt::Cluster::instance().get_associated_mmio_device(device_id);

tt_metal/impl/dispatch/kernel_config/fd_kernel.hpp

@@ -72,13 +72,15 @@ class FDKernel {
         chip_id_t servicing_device_id,
         uint8_t cq_id,
         noc_selection_t noc_selection,
-        DispatchWorkerType type);
+        tt::tt_metal::DispatchWorkerType type);
 
     // Register another kernel as upstream/downstream of this one
     void AddUpstreamKernel(FDKernel* upstream) { upstream_kernels_.push_back(upstream); }
     void AddDownstreamKernel(FDKernel* downstream) { downstream_kernels_.push_back(downstream); }
 
-    virtual CoreType GetCoreType() { return dispatch_core_manager::instance().get_dispatch_core_type(device_->id()); }
+    virtual CoreType GetCoreType() {
+        return tt::tt_metal::dispatch_core_manager::instance().get_dispatch_core_type(device_->id());
+    }
     tt_cxy_pair GetLogicalCore() { return logical_core_; }
     tt_cxy_pair GetVirtualCore() {
         return tt::Cluster::instance().get_virtual_coordinate_from_logical_coordinates(logical_core_, GetCoreType());
@@ -88,7 +90,7 @@ class FDKernel {
     // Get the port index for which a given kernel is upstream/downstream of this one
     int GetUpstreamPort(FDKernel* other) { return GetPort(other, this->upstream_kernels_); }
     int GetDownstreamPort(FDKernel* other) { return GetPort(other, this->downstream_kernels_); }
-    void AddDeviceAndProgram(Device* device, Program* program) {
+    void AddDeviceAndProgram(tt::tt_metal::Device* device, tt::tt_metal::Program* program) {
         device_ = device;
         program_ = program;
     };
@@ -116,8 +118,8 @@ class FDKernel {
     static chip_id_t GetDownstreamDeviceId(chip_id_t device_id);
     static uint32_t GetTunnelStop(chip_id_t device_id);
 
-    Device* device_ = nullptr;  // Set at configuration time by AddDeviceAndProgram()
-    Program* program_ = nullptr;
+    tt::tt_metal::Device* device_ = nullptr;  // Set at configuration time by AddDeviceAndProgram()
+    tt::tt_metal::Program* program_ = nullptr;
     tt_cxy_pair logical_core_;
     chip_id_t device_id_;
     chip_id_t servicing_device_id_;  // Remote chip that this PREFETCH_H/DISPATCH_H is servicing

tt_metal/impl/dispatch/kernel_config/mux.cpp

@@ -9,6 +9,8 @@
 #include "tt_metal/host_api.hpp"
 #include "tt_metal/detail/tt_metal.hpp"
 
+using namespace tt::tt_metal;
+
 void MuxKernel::GenerateStaticConfigs() {
     uint16_t channel = tt::Cluster::instance().get_assigned_channel_for_device(device_->id());
     logical_core_ = dispatch_core_manager::instance().mux_d_core(device_->id(), channel, this->cq_id_);

tt_metal/impl/dispatch/kernel_config/prefetch.cpp

@@ -9,6 +9,8 @@
 #include "tt_metal/host_api.hpp"
 #include "tt_metal/detail/tt_metal.hpp"
 
+using namespace tt::tt_metal;
+
 void PrefetchKernel::GenerateStaticConfigs() {
     uint16_t channel = tt::Cluster::instance().get_assigned_channel_for_device(device_->id());
     uint8_t cq_id_ = this->cq_id_;

tt_metal/impl/dispatch/topology.cpp

@@ -18,6 +18,8 @@
 #define DISPATCH_MAX_UPSTREAM 4
 #define DISPATCH_MAX_DOWNSTREAM 4
 
+using namespace tt::tt_metal;
+
 typedef struct {
     int id;
     chip_id_t device_id;                          // Device that this kernel is located on

tt_metal/impl/dispatch/topology.hpp

@@ -9,7 +9,7 @@
 void populate_fd_kernels(const std::set<chip_id_t>& device_ids, uint32_t num_hw_cqs);
 
 // Fill out all settings for FD kernels on the given device, and add them to a Program and return it.
-std::unique_ptr<Program> create_and_compile_cq_program(Device* device);
+std::unique_ptr<tt::tt_metal::Program> create_and_compile_cq_program(tt::tt_metal::Device* device);
 
 // Performa additional configuration (writing to specific L1 addresses, etc.) for FD kernels on this device.
-void configure_dispatch_cores(Device* device);
+void configure_dispatch_cores(tt::tt_metal::Device* device);

ttnn/cpp/ttnn/operations/conv/conv2d/device/conv2d_op.cpp

@@ -23,10 +23,12 @@
 
 #include "ttnn/operations/sliding_window/sliding_window.hpp"
 #include "ttnn/tensor/tensor_utils.hpp"
+
 using namespace tt::constants;
+using namespace tt::tt_metal;
+
 namespace optimized_conv_op_utils {
 using namespace tt;
-using namespace tt::tt_metal;
 
 std::pair<std::vector<uint32_t>, std::vector<uint32_t>> compute_opt_conv_activation_as_mm_shape(
     const tt::tt_metal::LegacyShape& conv_activation_shape,

ttnn/cpp/ttnn/operations/core/to_dtype/to_dtype_op.hpp

@@ -74,10 +74,14 @@ inline Tensor convert_to_cpp_supported_dtype(const Tensor& input_tensor) {
         [&](auto&& buffer) -> Tensor {
             using T = std::decay_t<decltype(buffer)>;
             if constexpr (std::is_same_v<T, tt::tt_metal::OwnedBuffer>) {
-                return Tensor{OwnedStorage{buffer}, input_tensor.get_shape(), input_dtype, input_tensor.get_layout()};
+                return Tensor{
+                    tt::tt_metal::OwnedStorage{buffer},
+                    input_tensor.get_shape(),
+                    input_dtype,
+                    input_tensor.get_layout()};
             } else if constexpr (std::is_same_v<T, tt::tt_metal::BorrowedBuffer>) {
                 return Tensor{
-                    BorrowedStorage{buffer, []() {}, []() {}},
+                    tt::tt_metal::BorrowedStorage{buffer, []() {}, []() {}},
                     input_tensor.get_shape(),
                     input_dtype,
                     input_tensor.get_layout()};
@@ -110,7 +114,7 @@ inline std::vector<NewT> cast(const tt::tt_metal::borrowed_buffer::Buffer<OldT>&
 template <typename T>
 Tensor create_owned_tensor(std::vector<T>&& data, const Shape& shape, DataType data_type, Layout layout) {
     auto buffer = tt::tt_metal::owned_buffer::create(std::move(data));
-    auto storage = OwnedStorage{std::move(buffer)};
+    auto storage = tt::tt_metal::OwnedStorage{std::move(buffer)};
     return Tensor(std::move(storage), shape, data_type, layout);
 }
 
@@ -146,15 +150,19 @@ inline Tensor create_tensor_from_buffer(
             auto data = cast<float, T>(input_buffer);
             auto buffer = tt::tt_metal::owned_buffer::create<float>(std::move(data));
             auto tensor =
-                Tensor(OwnedStorage{std::move(buffer)}, shape, DataType::FLOAT32, Layout::ROW_MAJOR).to(Layout::TILE);
+                Tensor(tt::tt_metal::OwnedStorage{std::move(buffer)}, shape, DataType::FLOAT32, Layout::ROW_MAJOR)
+                    .to(Layout::TILE);
             auto output_float_data = tt::tt_metal::owned_buffer::get_as<float>(tensor).get();
             auto output_packed_data =
                 dtype == DataType::BFLOAT8_B
                     ? pack_fp32_vec_as_bfp8_tiles(output_float_data, /*row_major_input=*/false, /*is_exp_a=*/false)
                     : pack_fp32_vec_as_bfp4_tiles(output_float_data, /*row_major_input=*/false, /*is_exp_a=*/false);
             auto output_buffer = tt::tt_metal::owned_buffer::create<uint32_t>(std::move(output_packed_data));
             return Tensor(
-                OwnedStorage{std::move(output_buffer)}, shape, dtype, Layout::TILE);  // has to be in tile layout
+                tt::tt_metal::OwnedStorage{std::move(output_buffer)},
+                shape,
+                dtype,
+                Layout::TILE);  // has to be in tile layout
         }
         default: {
             TT_THROW("Unsupported DataType: {}", dtype);

ttnn/cpp/ttnn/operations/data_movement/pad/device/pad_op.cpp

@@ -6,6 +6,7 @@
 #include "pad_op.hpp"
 #include "pad_program_factory.hpp"
 
+using namespace tt::tt_metal;
 namespace ttnn::operations::data_movement {
 
 void Pad::validate_with_output_tensors(

ttnn/cpp/ttnn/operations/data_movement/permute/device/permute_device_operation.hpp

@@ -54,9 +54,9 @@ struct PermuteDeviceOperation {
     struct MultiCoreBlockedGeneric {
         // Shared variables are the variables that are shared between the create and override_runtime_arguments methods
         struct shared_variables_t {
-            KernelHandle unary_reader_kernel_id;
-            KernelHandle unary_writer_kernel_id;
-            KernelHandle compute_kernel_id;
+            tt::tt_metal::KernelHandle unary_reader_kernel_id;
+            tt::tt_metal::KernelHandle unary_writer_kernel_id;
+            tt::tt_metal::KernelHandle compute_kernel_id;
             CoreRangeSet core_range;
         };
         using cached_program_t = ttnn::device_operation::CachedProgram<shared_variables_t>;

ttnn/cpp/ttnn/operations/data_movement/reshape_view/device/reshape_rm_op.cpp

@@ -7,6 +7,8 @@
 
 #include <cstdint>
 
+using namespace tt::tt_metal;
+
 namespace ttnn {
 
 void RM_RESHAPE_STRUCT::validate(const std::vector<Tensor>& input_tensors) const {

ttnn/cpp/ttnn/operations/data_movement/reshape_view/device/reshape_rm_op.hpp

@@ -11,21 +11,20 @@ namespace ttnn {
 
 struct RM_RESHAPE_STRUCT {
     const ttnn::Shape output_shape;
-    MemoryConfig output_mem_config;
-
+    tt::tt_metal::MemoryConfig output_mem_config;
 
     //Required functions to all tensor op functions
     void update_structure (const Tensor& input_tensor);
     void validate(const std::vector<Tensor> &input_tensors) const;
     std::vector<SimpleShape> compute_output_shapes(const std::vector<Tensor> &input_tensors) const;
     std::vector<Tensor> create_output_tensors(const std::vector<Tensor> &input_tensors) const;
-    operation::ProgramWithCallbacks create_program(
-        const std::vector<Tensor> &input_tensors, std::vector<Tensor> &output_tensors) const;
+    tt::tt_metal::operation::ProgramWithCallbacks create_program(
+        const std::vector<Tensor>& input_tensors, std::vector<Tensor>& output_tensors) const;
 };
 
 
 }// namespace ttnn
 namespace ttnn::operations::data_movement::rm_reshape{
 
-operation::ProgramWithCallbacks rm_reshape_preparer(const Tensor& input, const Tensor& output);
+tt::tt_metal::operation::ProgramWithCallbacks rm_reshape_preparer(const Tensor& input, const Tensor& output);
 }

ttnn/cpp/ttnn/operations/data_movement/sharded/sharded_to_interleaved/device/sharded_to_interleaved_program_factory.hpp

@@ -8,7 +8,7 @@
 
 namespace ttnn::operations::data_movement::detail {
 
-operation::ProgramWithCallbacks sharded_to_interleaved_multi_core(
+tt::tt_metal::operation::ProgramWithCallbacks sharded_to_interleaved_multi_core(
     const Tensor& a,
     const Tensor& output,
     bool is_l1_aligned = false,

ttnn/cpp/ttnn/operations/data_movement/slice/device/slice_op.hpp

@@ -23,7 +23,7 @@ struct SliceDeviceOperation {
     void validate_with_output_tensors(
         const std::vector<Tensor>& input_tensors, const std::vector<std::optional<Tensor>>& output_tensors) const;
     std::vector<ttnn::TensorSpec> compute_output_specs(const std::vector<Tensor>& input_tensors) const;
-    operation::ProgramWithCallbacks create_program(
+    tt::tt_metal::operation::ProgramWithCallbacks create_program(
         const std::vector<Tensor>& input_tensors, std::vector<Tensor>& output_tensors) const;
 };
 

ttnn/cpp/ttnn/operations/data_movement/slice/device/slice_program_factory.hpp

@@ -7,7 +7,7 @@
 
 namespace ttnn::operations::data_movement::detail {
 
-operation::ProgramWithCallbacks slice_multi_core(
+tt::tt_metal::operation::ProgramWithCallbacks slice_multi_core(
     const Tensor& a,
     Tensor& output,
     const tt::tt_metal::LegacyShape& output_tensor_start,

ttnn/cpp/ttnn/operations/data_movement/split/device/split_op.hpp

@@ -17,7 +17,7 @@ struct SplitDeviceOperation {
     void validate(const std::vector<Tensor>& input_tensors) const;
     std::vector<tt::tt_metal::LegacyShape> compute_output_shapes(const std::vector<Tensor>& input_tensors) const;
     std::vector<Tensor> create_output_tensors(const std::vector<Tensor>& input_tensors) const;
-    operation::ProgramWithCallbacks create_program(
+    tt::tt_metal::operation::ProgramWithCallbacks create_program(
         const std::vector<Tensor>& input_tensors, std::vector<Tensor>& output_tensors) const;
 };
 

ttnn/cpp/ttnn/operations/data_movement/split/device/split_program_factory.hpp

@@ -6,6 +6,6 @@
 
 namespace ttnn::operations::data_movement::detail {
 
-operation::ProgramWithCallbacks split_last_dim_two_chunks_tiled(
+tt::tt_metal::operation::ProgramWithCallbacks split_last_dim_two_chunks_tiled(
     const Tensor& input_tensor, std::vector<Tensor>& output_tensors, const tt::tt_metal::MemoryConfig& mem_config);
 }

ttnn/cpp/ttnn/operations/data_movement/transpose/device/transpose_op.cpp

@@ -14,6 +14,7 @@
 #include "noc/noc_parameters.h"  // DRAM_ALIGNMENT
 
 using namespace tt::constants;
+using namespace tt::tt_metal;
 
 namespace ttnn::operations::data_movement {
 

ttnn/cpp/ttnn/operations/data_movement/transpose/device/transpose_op.hpp

@@ -20,7 +20,7 @@ struct Transpose {
 
     void validate(const std::vector<Tensor>& input_tensors) const;
     std::vector<ttnn::TensorSpec> compute_output_specs(const std::vector<Tensor>& input_tensors) const;
-    operation::ProgramWithCallbacks create_program(
+    tt::tt_metal::operation::ProgramWithCallbacks create_program(
         const std::vector<Tensor>& input_tensors, std::vector<Tensor>& output_tensors) const;
     TransposeOpParallelizationStrategy get_parallelization_strategy(const std::vector<Tensor>& input_tensors) const;
 };

ttnn/cpp/ttnn/operations/data_movement/transpose/device/transpose_program_factory.hpp

@@ -4,14 +4,14 @@
 
 namespace ttnn::operations::data_movement::detail {
 
-operation::ProgramWithCallbacks transpose_wh_multi_core(const Tensor& a, Tensor& output);
-operation::ProgramWithCallbacks transpose_wh_multi_core_sharded(const Tensor& a, Tensor& output);
-operation::ProgramWithCallbacks transpose_wh_multi_core_sharded_rm(const Tensor& a, Tensor& output);
-operation::ProgramWithCallbacks transpose_hc_multi_core(
+tt::tt_metal::operation::ProgramWithCallbacks transpose_wh_multi_core(const Tensor& a, Tensor& output);
+tt::tt_metal::operation::ProgramWithCallbacks transpose_wh_multi_core_sharded(const Tensor& a, Tensor& output);
+tt::tt_metal::operation::ProgramWithCallbacks transpose_wh_multi_core_sharded_rm(const Tensor& a, Tensor& output);
+tt::tt_metal::operation::ProgramWithCallbacks transpose_hc_multi_core(
     const Tensor& a, Tensor& output, const std::optional<float>& pad_value);
-operation::ProgramWithCallbacks transpose_hc_multi_core_tiled_interleaved(
+tt::tt_metal::operation::ProgramWithCallbacks transpose_hc_multi_core_tiled_interleaved(
     const Tensor& a, Tensor& output, const std::optional<float>& pad_value);
-operation::ProgramWithCallbacks transpose_hc_multi_core_sharded(const Tensor& a, Tensor& output);
-operation::ProgramWithCallbacks transpose_cn_multi_core(const Tensor& a, Tensor& output);
+tt::tt_metal::operation::ProgramWithCallbacks transpose_hc_multi_core_sharded(const Tensor& a, Tensor& output);
+tt::tt_metal::operation::ProgramWithCallbacks transpose_cn_multi_core(const Tensor& a, Tensor& output);
 
 }  // namespace ttnn::operations::data_movement::detail

ttnn/cpp/ttnn/operations/data_movement/untilize/device/untilize_op.hpp

@@ -18,7 +18,7 @@ uint32_t get_num_cores(CoreCoord grid_size, uint32_t nblocks);
 }  // namespace untilize_helpers
 
 struct Untilize {
-    const MemoryConfig output_mem_config;
+    const tt::tt_metal::MemoryConfig output_mem_config;
     const bool use_multicore;
     const bool use_pack_untilize;
     const bool fp32_dest_acc_en;
@@ -27,7 +27,7 @@ struct Untilize {
     std::vector<tt::tt_metal::LegacyShape> compute_output_shapes(const std::vector<Tensor>& input_tensors) const;
     std::vector<Tensor> create_output_tensors(
         const std::vector<Tensor>& input_tensors, const std::vector<std::optional<Tensor>>& output_tensors) const;
-    operation::ProgramWithCallbacks create_program(
+    tt::tt_metal::operation::ProgramWithCallbacks create_program(
         const std::vector<Tensor>& input_tensors, std::vector<Tensor>& output_tensors) const;
 };
 

ttnn/cpp/ttnn/operations/data_movement/untilize/device/untilize_program_factory.hpp

@@ -8,10 +8,10 @@
 
 namespace ttnn::operations::data_movement::detail {
 
-operation::ProgramWithCallbacks untilize_multi_core(
+tt::tt_metal::operation::ProgramWithCallbacks untilize_multi_core(
     const Tensor& a, Tensor& output, bool use_pack_untilize, bool fp32_dest_acc_en);
 
-operation::ProgramWithCallbacks untilize_single_core(
+tt::tt_metal::operation::ProgramWithCallbacks untilize_single_core(
     const Tensor& a, Tensor& output, bool use_pack_untilize, bool fp32_dest_acc_en);
 
 }  // namespace ttnn::operations::data_movement::detail

ttnn/cpp/ttnn/operations/data_movement/untilize_with_halo_v2/device/untilize_with_halo_v2_op.hpp

@@ -16,15 +16,15 @@ struct UntilizeWithHaloV2 {
     const uint32_t pad_val_;
     const uint32_t ncores_nhw_;
     const uint32_t max_out_nsticks_per_core_;
-    const MemoryConfig out_mem_config_;
+    const tt::tt_metal::MemoryConfig out_mem_config_;
     const bool remote_read_;
     const bool transpose_mcast_;
 
     void validate(const std::vector<Tensor>& input_tensors) const;
     std::vector<tt::tt_metal::LegacyShape> compute_output_shapes(const std::vector<Tensor>& input_tensors) const;
     std::vector<Tensor> create_output_tensors(
         const std::vector<Tensor>& input_tensors, const std::vector<std::optional<Tensor>>& output_tensors) const;
-    operation::ProgramWithCallbacks create_program(
+    tt::tt_metal::operation::ProgramWithCallbacks create_program(
         const std::vector<Tensor>& input_tensors, std::vector<Tensor>& output_tensors) const;
 };
 

ttnn/cpp/ttnn/operations/data_movement/untilize_with_halo_v2/device/untilize_with_halo_v2_program_factory.hpp

@@ -8,7 +8,7 @@
 
 namespace ttnn::operations::data_movement::detail {
 
-operation::ProgramWithCallbacks untilize_with_halo_multi_core_v2(
+tt::tt_metal::operation::ProgramWithCallbacks untilize_with_halo_multi_core_v2(
     Program& program,
     const Tensor& input_tensor,
     const uint32_t pad_val,

ttnn/cpp/ttnn/operations/data_movement/untilize_with_unpadding/device/untilize_with_unpadding_op.cpp

@@ -8,6 +8,8 @@
 #include "ttnn/run_operation.hpp"
 #include "untilize_with_unpadding_program_factory.hpp"
 
+using namespace tt::tt_metal;
+
 namespace ttnn::operations::data_movement {
 
 void UntilizeWithUnpadding::validate(const std::vector<Tensor>& input_tensors) const {