diff --git a/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_op.cpp b/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_op.cpp index 008e460c018..882db9c4842 100644 --- a/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_op.cpp +++ b/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_op.cpp @@ -41,11 +41,11 @@ void MorehClipGradNormStep1::validate( const std::vector &input_tensors, const std::vector> &optional_input_tensors) const { for (const auto &input : input_tensors) { - check_tensor(input, "moreh_clip_grad_norm_step1", "input"); + ttnn::operations::check_tensor(input, "moreh_clip_grad_norm_step1", "input"); } const auto &tmp_pow_sum = optional_input_tensors.at(0).value(); - check_tensor(tmp_pow_sum, "moreh_clip_grad_norm_step1", "tmp_pow_sum"); + ttnn::operations::check_tensor(tmp_pow_sum, "moreh_clip_grad_norm_step1", "tmp_pow_sum"); }; std::vector MorehClipGradNormStep1::compute_output_shapes(const std::vector &) const { return {}; } @@ -99,10 +99,10 @@ void moreh_clip_grad_norm_step1(const std::vector &inputs, float norm_ty void MorehClipGradNormStep2::validate(const std::vector &input_tensors) const { const auto &tmp_pow_sum = input_tensors.at(0); - check_tensor(tmp_pow_sum, "moreh_clip_grad_norm_step2", "tmp_pow_sum"); + ttnn::operations::check_tensor(tmp_pow_sum, "moreh_clip_grad_norm_step2", "tmp_pow_sum"); const auto &total_norm = input_tensors.at(1); - check_tensor(total_norm, "moreh_clip_grad_norm_step2", "total_norm"); + ttnn::operations::check_tensor(total_norm, "moreh_clip_grad_norm_step2", "total_norm"); } std::vector MorehClipGradNormStep2::compute_output_shapes(const std::vector &) const { return {}; } @@ -139,11 +139,11 @@ void MorehClipGradNormStep3::validate( const std::vector &input_tensors, const std::vector> &optional_input_tensors) const { for (const auto &input : input_tensors) { - check_tensor(input, "moreh_clip_grad_norm_step3", "input"); + ttnn::operations::check_tensor(input, "moreh_clip_grad_norm_step3", "input"); } const auto &clip_coef_clamped = optional_input_tensors.at(0).value(); - check_tensor(clip_coef_clamped, "moreh_clip_grad_norm_step3", "clip_coef_clamped"); + ttnn::operations::check_tensor(clip_coef_clamped, "moreh_clip_grad_norm_step3", "clip_coef_clamped"); } std::vector MorehClipGradNormStep3::compute_output_shapes(const std::vector &) const { return {}; } diff --git a/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step1/moreh_clip_grad_norm_step1.cpp b/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step1/moreh_clip_grad_norm_step1.cpp index 29b9f51b546..3ae17e72f25 100644 --- a/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step1/moreh_clip_grad_norm_step1.cpp +++ b/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step1/moreh_clip_grad_norm_step1.cpp @@ -84,7 +84,7 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step1_impl( const auto cb_data_format = tt_metal::datatype_to_dataformat_converter(tmp_pow_sum.get_dtype()); - CreateCircularBuffer( + ttnn::operations::CreateCircularBuffer( program, core_group_1, cb_data_format, @@ -112,8 +112,8 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step1_impl( "ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step1/kernels/" "writer_moreh_clip_grad_norm_step1.cpp"; - const auto reader_kernels_id = CreateReadKernel(program, reader_kernel_file, core_group_1); - const auto writer_kernels_id = CreateWriteKernel(program, writer_kernel_file, core_group_1); + const auto reader_kernels_id = ttnn::operations::CreateReadKernel(program, reader_kernel_file, core_group_1); + const auto writer_kernels_id = ttnn::operations::CreateWriteKernel(program, writer_kernel_file, core_group_1); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -127,7 +127,7 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step1_impl( "moreh_clip_grad_norm_step1_kernel.cpp"; const auto compute_kernels_id = - CreateComputeKernel(program, compute_kernel_file, {core_group_1, num_inputs_per_core_group_1}, compute_defines); + ttnn::operations::CreateComputeKernel(program, compute_kernel_file, {core_group_1, num_inputs_per_core_group_1}, compute_defines); //////////////////////////////////////////////////////////////////////////// // RuntimeArgs SetUp @@ -146,7 +146,7 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step1_impl( // reader const std::array reader_runtime_args{ input_addr, - static_cast(is_dram(input)), + static_cast(ttnn::operations::is_dram(input)), num_tiles, *reinterpret_cast(&decimal), origin_h, @@ -155,7 +155,7 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step1_impl( // writer const std::array writer_runtime_args{ - output_addr, static_cast(is_dram(tmp_pow_sum)), tile_offset}; + output_addr, static_cast(ttnn::operations::is_dram(tmp_pow_sum)), tile_offset}; SetRuntimeArgs(program, writer_kernels_id, core, writer_runtime_args); // compute diff --git a/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step2/moreh_clip_grad_norm_step2.cpp b/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step2/moreh_clip_grad_norm_step2.cpp index b65a1a7ea58..5b7826cbb6e 100644 --- a/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step2/moreh_clip_grad_norm_step2.cpp +++ b/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step2/moreh_clip_grad_norm_step2.cpp @@ -58,7 +58,7 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step2_impl( const auto cb_data_format = tt_metal::datatype_to_dataformat_converter(total_norm.get_dtype()); - CreateCircularBuffer( + ttnn::operations::CreateCircularBuffer( program, single_core, cb_data_format, @@ -82,8 +82,8 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step2_impl( "ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step2/kernels/" "writer_moreh_clip_grad_norm_step2.cpp"; - const auto reader_kernels_id = CreateReadKernel(program, reader_kernel_file, single_core); - const auto writer_kernels_id = CreateWriteKernel(program, writer_kernel_file, single_core); + const auto reader_kernels_id = ttnn::operations::CreateReadKernel(program, reader_kernel_file, single_core); + const auto writer_kernels_id = ttnn::operations::CreateWriteKernel(program, writer_kernel_file, single_core); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -92,7 +92,7 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step2_impl( "ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step2/kernels/" "moreh_clip_grad_norm_step2_kernel.cpp"; - const auto compute_kernels_id = CreateComputeKernel(program, compute_kernel_file, {single_core, num_tiles}); + const auto compute_kernels_id = ttnn::operations::CreateComputeKernel(program, compute_kernel_file, {single_core, num_tiles}); //////////////////////////////////////////////////////////////////////////// // RuntimeArgs SetUp @@ -102,11 +102,11 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step2_impl( // reader const std::array reader_runtime_args{ - input_addr, static_cast(is_dram(tmp_pow_sum)), num_tiles, *reinterpret_cast(&decimal)}; + input_addr, static_cast(ttnn::operations::is_dram(tmp_pow_sum)), num_tiles, *reinterpret_cast(&decimal)}; SetRuntimeArgs(program, reader_kernels_id, single_core, reader_runtime_args); // writer - const std::array writer_runtime_args{output_addr, static_cast(is_dram(total_norm))}; + const std::array writer_runtime_args{output_addr, static_cast(ttnn::operations::is_dram(total_norm))}; SetRuntimeArgs(program, writer_kernels_id, single_core, writer_runtime_args); // compute diff --git a/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step3/moreh_clip_grad_norm_step3.cpp b/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step3/moreh_clip_grad_norm_step3.cpp index f62b9eace7a..0d5696ea59a 100644 --- a/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step3/moreh_clip_grad_norm_step3.cpp +++ b/ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step3/moreh_clip_grad_norm_step3.cpp @@ -62,7 +62,7 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step3_impl( const auto cb_data_format = tt_metal::datatype_to_dataformat_converter(inputs.at(0).get_dtype()); - CreateCircularBuffer( + ttnn::operations::CreateCircularBuffer( program, core_group_1, cb_data_format, @@ -82,8 +82,8 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step3_impl( "ttnn/cpp/ttnn/deprecated/tt_dnn/op_library/moreh_clip_grad_norm/moreh_clip_grad_norm_step3/kernels/" "writer_moreh_clip_grad_norm_step3.cpp"; - const auto reader_kernels_id = CreateReadKernel(program, reader_kernel_file, core_group_1); - const auto writer_kernels_id = CreateWriteKernel(program, writer_kernel_file, core_group_1); + const auto reader_kernels_id = ttnn::operations::CreateReadKernel(program, reader_kernel_file, core_group_1); + const auto writer_kernels_id = ttnn::operations::CreateWriteKernel(program, writer_kernel_file, core_group_1); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -93,7 +93,7 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step3_impl( "moreh_clip_grad_norm_step3_kernel.cpp"; const auto compute_kernels_id = - CreateComputeKernel(program, compute_kernel_file, {core_group_1, num_inputs_per_core_group_1}); + ttnn::operations::CreateComputeKernel(program, compute_kernel_file, {core_group_1, num_inputs_per_core_group_1}); //////////////////////////////////////////////////////////////////////////// // RuntimeArgs SetUp @@ -109,14 +109,14 @@ operation::ProgramWithCallbacks moreh_clip_grad_norm_step3_impl( // reader const std::array reader_runtime_args{ input_addr, - static_cast(is_dram(input)), + static_cast(ttnn::operations::is_dram(input)), clip_coef_clamped_addr, - static_cast(is_dram(clip_coef_clamped)), + static_cast(ttnn::operations::is_dram(clip_coef_clamped)), num_tiles}; SetRuntimeArgs(program, reader_kernels_id, core, reader_runtime_args); // writer - const std::array writer_runtime_args{input_addr, static_cast(is_dram(input)), num_tiles}; + const std::array writer_runtime_args{input_addr, static_cast(ttnn::operations::is_dram(input)), num_tiles}; SetRuntimeArgs(program, writer_kernels_id, core, writer_runtime_args); // compute diff --git a/ttnn/cpp/ttnn/operations/experimental/reduction/fast_reduce_nc/device/fast_reduce_nc_device_operation.cpp b/ttnn/cpp/ttnn/operations/experimental/reduction/fast_reduce_nc/device/fast_reduce_nc_device_operation.cpp index e4ddeaf7d9e..c0680538354 100644 --- a/ttnn/cpp/ttnn/operations/experimental/reduction/fast_reduce_nc/device/fast_reduce_nc_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/experimental/reduction/fast_reduce_nc/device/fast_reduce_nc_device_operation.cpp @@ -50,8 +50,8 @@ void FastReduceNCDeviceOperation::validate_with_output_tensors( auto& output = output_tensors.at(0); // validate tensor - tt::operations::primary::check_tensor(input, "FastReduceNC", "input", {DataType::BFLOAT16, DataType::BFLOAT8_B}); - tt::operations::primary::check_tensor(output, "FastReduceNC", "output", {DataType::BFLOAT16, DataType::BFLOAT8_B}); + check_tensor(input, "FastReduceNC", "input", {DataType::BFLOAT16, DataType::BFLOAT8_B}); + check_tensor(output, "FastReduceNC", "output", {DataType::BFLOAT16, DataType::BFLOAT8_B}); // validate input dim const auto input_rank = input.get_logical_shape().rank(); diff --git a/ttnn/cpp/ttnn/operations/full/device/full_program_factory.cpp b/ttnn/cpp/ttnn/operations/full/device/full_program_factory.cpp index 910e8dfe8fa..940642ba8bd 100644 --- a/ttnn/cpp/ttnn/operations/full/device/full_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/full/device/full_program_factory.cpp @@ -24,7 +24,7 @@ FullOperation::ProgramFactory::cached_program_t FullOperation::ProgramFactory::c auto grid = tensor_args.any.device()->compute_with_storage_grid_size(); auto num_tiles = output.volume() / TILE_HW; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - split_work_to_cores(grid, num_tiles); + tt::tt_metal::split_work_to_cores(grid, num_tiles); tt::DataFormat data_format = tt::tt_metal::datatype_to_dataformat_converter(dtype); uint32_t single_tile_size = tt::tt_metal::detail::TileSize(data_format); @@ -40,7 +40,7 @@ FullOperation::ProgramFactory::cached_program_t FullOperation::ProgramFactory::c // Create circular buffer auto cb_index = tt::CB::c_intermed0; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -57,7 +57,7 @@ FullOperation::ProgramFactory::cached_program_t FullOperation::ProgramFactory::c default: break; } - auto writer_id = tt::operations::primary::CreateWriteKernel( + auto writer_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/full/device/kernels/writer_full.cpp", all_cores, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/moreh_adam_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/moreh_adam_device_operation.cpp index 3cc32ff7ed1..420ed10cb3f 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/moreh_adam_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/moreh_adam_device_operation.cpp @@ -18,31 +18,31 @@ void MorehAdamOperation::validate_inputs( auto& exp_avg_in = tensor_args.exp_avg_in; auto& exp_avg_sq_in = tensor_args.exp_avg_sq_in; - tt::operations::primary::check_tensor(params_in, "moreh_adam", "params_in"); - tt::operations::primary::check_tensor(grad, "moreh_adam", "grad"); - tt::operations::primary::check_tensor(exp_avg_in, "moreh_adam", "exp_avg_in"); - tt::operations::primary::check_tensor(exp_avg_sq_in, "moreh_adam", "exp_avg_sq_in"); + check_tensor(params_in, "moreh_adam", "params_in"); + check_tensor(grad, "moreh_adam", "grad"); + check_tensor(exp_avg_in, "moreh_adam", "exp_avg_in"); + check_tensor(exp_avg_sq_in, "moreh_adam", "exp_avg_sq_in"); if (tensor_args.max_exp_avg_sq_in) { - tt::operations::primary::check_tensor(*tensor_args.max_exp_avg_sq_in, "moreh_adam", "max_exp_avg_sq_in"); + check_tensor(*tensor_args.max_exp_avg_sq_in, "moreh_adam", "max_exp_avg_sq_in"); } const auto& params_out = tensor_args.output_tensors.at(0); if (params_out.has_value()) { - tt::operations::primary::check_tensor(params_out.value(), "moreh_adam", "params_out"); + check_tensor(params_out.value(), "moreh_adam", "params_out"); } if (tensor_args.output_tensors.at(1).has_value()) { - tt::operations::primary::check_tensor(tensor_args.output_tensors.at(1).value(), "moreh_adam", "exp_avg_out"); + check_tensor(tensor_args.output_tensors.at(1).value(), "moreh_adam", "exp_avg_out"); } if (tensor_args.output_tensors.at(2).has_value()) { - tt::operations::primary::check_tensor(tensor_args.output_tensors.at(2).value(), "moreh_adam", "exp_avg_sq_out"); + check_tensor(tensor_args.output_tensors.at(2).value(), "moreh_adam", "exp_avg_sq_out"); } if (tensor_args.output_tensors.at(3).has_value()) { - tt::operations::primary::check_tensor( + check_tensor( tensor_args.output_tensors.at(3).value(), "moreh_adam", "max_exp_avg_sq_out"); } } diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/moreh_adam_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/moreh_adam_program_factory.cpp index 29527eae2ae..a9c8be4a79c 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/moreh_adam_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/moreh_adam_program_factory.cpp @@ -61,7 +61,7 @@ MorehAdamOperation::ProgramFactory::cached_program_t MorehAdamOperation::Program //////////////////////////////////////////////////////////////////////////// auto data_format = tt::tt_metal::datatype_to_dataformat_converter(param_in.get_dtype()); auto intermed_cb_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -94,17 +94,17 @@ MorehAdamOperation::ProgramFactory::cached_program_t MorehAdamOperation::Program //////////////////////////////////////////////////////////////////////////// const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(param_in)), - static_cast(tt::operations::primary::is_dram(grad)), - static_cast(tt::operations::primary::is_dram(exp_avg_in)), - static_cast(tt::operations::primary::is_dram(exp_avg_sq_in)), - static_cast(tt::operations::primary::is_dram(max_exp_avg_sq_in))}; + static_cast(is_dram(param_in)), + static_cast(is_dram(grad)), + static_cast(is_dram(exp_avg_in)), + static_cast(is_dram(exp_avg_sq_in)), + static_cast(is_dram(max_exp_avg_sq_in))}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(param_out)), - static_cast(tt::operations::primary::is_dram(exp_avg_out)), - static_cast(tt::operations::primary::is_dram(exp_avg_sq_out)), - static_cast(tt::operations::primary::is_dram(max_exp_avg_sq_out.value()))}; + static_cast(is_dram(param_out)), + static_cast(is_dram(exp_avg_out)), + static_cast(is_dram(exp_avg_sq_out)), + static_cast(is_dram(max_exp_avg_sq_out.value()))}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/kernels/" @@ -120,9 +120,9 @@ MorehAdamOperation::ProgramFactory::cached_program_t MorehAdamOperation::Program if (fp32_dest_acc_en) { data_movement_defines["FP32_DEST_ACC_EN"] = "1"; } - const auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + const auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, data_movement_defines); - const auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + const auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, data_movement_defines); //////////////////////////////////////////////////////////////////////////// @@ -143,7 +143,7 @@ MorehAdamOperation::ProgramFactory::cached_program_t MorehAdamOperation::Program "ttnn/cpp/ttnn/operations/moreh/moreh_adam/device/kernels/" "moreh_adam.cpp"; - auto compute_kernel_1_id = tt ::operations::primary::CreateComputeKernel( + auto compute_kernel_1_id = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_tiles_per_core_group_1, compute_args_group_1}, @@ -155,7 +155,7 @@ MorehAdamOperation::ProgramFactory::cached_program_t MorehAdamOperation::Program if (!core_group_2.ranges().empty()) { const std::vector compute_args_group_2{num_tiles_per_core_group_2}; - compute_kernel_2_id = tt::operations::primary::CreateComputeKernel( + compute_kernel_2_id = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_tiles_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/moreh_adamw_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/moreh_adamw_device_operation.cpp index 99dc832ec5f..71cbb5b4bb7 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/moreh_adamw_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/moreh_adamw_device_operation.cpp @@ -18,17 +18,17 @@ MorehAdamWDeviceOperation::program_factory_t MorehAdamWDeviceOperation::select_p void MorehAdamWDeviceOperation::validate_inputs( const operation_attributes_t& attributes, const tensor_args_t& tensor_args) { - tt::operations::primary::check_tensor( + check_tensor( tensor_args.param_in, "moreh_adamw", "param_in", {DataType::BFLOAT16, DataType::BFLOAT8_B}); - tt::operations::primary::check_tensor( + check_tensor( tensor_args.grad, "moreh_adamw", "grad", {DataType::BFLOAT16, DataType::BFLOAT8_B}); - tt::operations::primary::check_tensor( + check_tensor( tensor_args.exp_avg_in, "moreh_adamw", "exp_avg_in", {DataType::BFLOAT16, DataType::BFLOAT8_B}); - tt::operations::primary::check_tensor( + check_tensor( tensor_args.exp_avg_sq_in, "moreh_adamw", "exp_avg_sq_in", {DataType::BFLOAT16, DataType::BFLOAT8_B}); if (tensor_args.max_exp_avg_sq_in.has_value()) { - tt::operations::primary::check_tensor( + check_tensor( tensor_args.max_exp_avg_sq_in.value(), "moreh_adamw", "max_exp_avg_sq_in", @@ -36,22 +36,22 @@ void MorehAdamWDeviceOperation::validate_inputs( } if (tensor_args.param_out.has_value()) { - tt::operations::primary::check_tensor( + check_tensor( tensor_args.param_out.value(), "moreh_adamw", "param_out", {DataType::BFLOAT16, DataType::BFLOAT8_B}); } if (tensor_args.exp_avg_out.has_value()) { - tt::operations::primary::check_tensor( + check_tensor( tensor_args.exp_avg_out.value(), "moreh_adamw", "exp_avg_out", {DataType::BFLOAT16, DataType::BFLOAT8_B}); } if (tensor_args.exp_avg_sq_out.has_value()) { - tt::operations::primary::check_tensor( + check_tensor( tensor_args.exp_avg_sq_out.value(), "moreh_adamw", "exp_avg_sq_out", {DataType::BFLOAT16, DataType::BFLOAT8_B}); } if (tensor_args.max_exp_avg_sq_out.has_value()) { - tt::operations::primary::check_tensor( + check_tensor( tensor_args.max_exp_avg_sq_out.value(), "moreh_adamw", "max_exp_avg_sq_out", diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/multi_core_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/multi_core_program_factory.cpp index 7034e7e2b72..6e663eeea22 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/multi_core_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/multi_core_program_factory.cpp @@ -52,7 +52,7 @@ MorehAdamWDeviceOperation::MultiCore::cached_program_t MorehAdamWDeviceOperation const auto num_cores_y = grid.y; auto [num_cores, all_cores, core_group_1, core_group_2, num_units_per_core_group_1, num_units_per_core_group_2] = - split_work_to_cores(grid, num_units); + tt_metal::split_work_to_cores(grid, num_units); auto arch = param_in.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -64,7 +64,7 @@ MorehAdamWDeviceOperation::MultiCore::cached_program_t MorehAdamWDeviceOperation auto data_format = tt_metal::datatype_to_dataformat_converter(param_in.get_dtype()); auto intermed_cb_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -96,19 +96,19 @@ MorehAdamWDeviceOperation::MultiCore::cached_program_t MorehAdamWDeviceOperation // DataMovementKernel SetUp //////////////////////////////////////////////////////////////////////////// const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(param_in)), - static_cast(tt::operations::primary::is_dram(grad)), - static_cast(tt::operations::primary::is_dram(exp_avg_in)), - static_cast(tt::operations::primary::is_dram(exp_avg_sq_in)), + static_cast(is_dram(param_in)), + static_cast(is_dram(grad)), + static_cast(is_dram(exp_avg_in)), + static_cast(is_dram(exp_avg_sq_in)), static_cast( - max_exp_avg_sq_in.has_value() ? tt::operations::primary::is_dram(max_exp_avg_sq_in.value()) : false)}; + max_exp_avg_sq_in.has_value() ? is_dram(max_exp_avg_sq_in.value()) : false)}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(param_out)), - static_cast(tt::operations::primary::is_dram(exp_avg_out)), - static_cast(tt::operations::primary::is_dram(exp_avg_sq_out)), + static_cast(is_dram(param_out)), + static_cast(is_dram(exp_avg_out)), + static_cast(is_dram(exp_avg_sq_out)), static_cast( - max_exp_avg_sq_out.has_value() ? tt::operations::primary::is_dram(max_exp_avg_sq_out.value()) : false)}; + max_exp_avg_sq_out.has_value() ? is_dram(max_exp_avg_sq_out.value()) : false)}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/kernels/" @@ -128,15 +128,15 @@ MorehAdamWDeviceOperation::MultiCore::cached_program_t MorehAdamWDeviceOperation compute_defines["FP32_DEST_ACC_EN"] = "1"; } - const auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + const auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, data_movement_defines); - const auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + const auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, data_movement_defines); const std::vector compute_args_group_1{num_units_per_core_group_1}; const std::vector compute_args_group_2{num_units_per_core_group_2}; - auto compute_kernel_ids = tt::operations::primary::CreateComputeKernel( + auto compute_kernel_ids = CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_adamw/device/kernels/moreh_adamw.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_arange/device/moreh_arange_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_arange/device/moreh_arange_program_factory.cpp index 30a9fec34cc..8ff9536e22f 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_arange/device/moreh_arange_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_arange/device/moreh_arange_program_factory.cpp @@ -26,7 +26,7 @@ MorehArangeOperation::ProgramFactory::cached_program_t MorehArangeOperation::Pro Program program = Program(); // Create circular buffer - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, tt::tt_metal::datatype_to_dataformat_converter(dtype), @@ -44,7 +44,7 @@ MorehArangeOperation::ProgramFactory::cached_program_t MorehArangeOperation::Pro } uint32_t dst_is_dram = output.buffer()->buffer_type() == tt::tt_metal::BufferType::DRAM ? 1 : 0; - auto kernel_id = tt::operations::primary::CreateWriteKernel( + auto kernel_id = CreateWriteKernel( program, operation_attributes.untilize_out ? "ttnn/cpp/ttnn/operations/moreh/moreh_arange/device/kernels/writer_moreh_arange_rm.cpp" diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_cumsum/device/moreh_cumsum_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_cumsum/device/moreh_cumsum_program_factory.cpp index 74ad47ba25c..11961069c30 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_cumsum/device/moreh_cumsum_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_cumsum/device/moreh_cumsum_program_factory.cpp @@ -77,7 +77,7 @@ MorehCumsumDeviceOperation::ProgramFactory::cached_program_t MorehCumsumDeviceOp //////////////////////////////////////////////////////////////////////////// // CircularBuffer Setup //////////////////////////////////////////////////////////////////////////// - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -98,9 +98,9 @@ MorehCumsumDeviceOperation::ProgramFactory::cached_program_t MorehCumsumDeviceOp const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_cumsum/device/kernels/writer_moreh_cumsum_nc.cpp"; const auto reader_kernel_id = - tt::operations::primary::CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); + CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -108,7 +108,7 @@ MorehCumsumDeviceOperation::ProgramFactory::cached_program_t MorehCumsumDeviceOp const std::vector compute_args_group_1{num_cols_per_core_group_1}; std::map compute_defines; const auto compute_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_cumsum/device/kernels/moreh_cumsum_nc.cpp"; - const auto compute_kernel_1_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_1_id = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, @@ -120,7 +120,7 @@ MorehCumsumDeviceOperation::ProgramFactory::cached_program_t MorehCumsumDeviceOp std::optional compute_kernel_2_id = std::nullopt; if (!core_group_2.ranges().empty()) { const std::vector compute_args_group_2{num_cols_per_core_group_2}; - compute_kernel_2_id = tt::operations::primary::CreateComputeKernel( + compute_kernel_2_id = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, @@ -154,7 +154,7 @@ MorehCumsumDeviceOperation::ProgramFactory::cached_program_t MorehCumsumDeviceOp num_tiles_per_core, input_tile_offset, tile_offset, - static_cast(tt::operations::primary::is_dram(input)), + static_cast(is_dram(input)), HtWt, CHtWt, static_cast(dim), @@ -169,7 +169,7 @@ MorehCumsumDeviceOperation::ProgramFactory::cached_program_t MorehCumsumDeviceOp num_tiles_per_core, input_tile_offset, tile_offset, - static_cast(tt::operations::primary::is_dram(output)), + static_cast(is_dram(output)), HtWt, CHtWt, static_cast(dim), diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/moreh_dot_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/moreh_dot_device_operation.cpp index 3536fa55129..d8c67f67898 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/moreh_dot_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/moreh_dot_device_operation.cpp @@ -18,8 +18,8 @@ void MorehDotOperation::validate(const operation_attributes_t& operation_attribu const auto& input_a = tensor_args.input_a; const auto& input_b = tensor_args.input_b; - TT_FATAL(tt::operations::primary::is_1d_tensor(input_a), "Invalid input tensor dimensions."); - TT_FATAL(tt::operations::primary::is_1d_tensor(input_b), "Invalid input tensor dimensions."); + TT_FATAL(is_1d_tensor(input_a), "Invalid input tensor dimensions."); + TT_FATAL(is_1d_tensor(input_b), "Invalid input tensor dimensions."); const auto& a_shape_wo_padding = input_a.get_legacy_shape().without_padding(); const auto& b_shape_wo_padding = input_b.get_legacy_shape().without_padding(); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/moreh_dot_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/moreh_dot_program_factory.cpp index cad1e01e515..78a7ee2664a 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/moreh_dot_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/moreh_dot_program_factory.cpp @@ -53,7 +53,7 @@ MorehDotOperation::SingleCore::cached_program_t MorehDotOperation::SingleCore::c CoreCoord core = {0, 0}; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, std::set{CoreRange(core, core)}, cb_data_format, @@ -67,19 +67,19 @@ MorehDotOperation::SingleCore::cached_program_t MorehDotOperation::SingleCore::c }); std::vector reader_compile_time_args = { - (std::uint32_t)tt::operations::primary::is_dram(src0_buffer), - (std::uint32_t)tt::operations::primary::is_dram(src1_buffer), + (std::uint32_t)is_dram(src0_buffer), + (std::uint32_t)is_dram(src1_buffer), *reinterpret_cast(&scaler)}; std::vector writer_compile_time_args = { - (std::uint32_t)CB::c_out0, (std::uint32_t)tt::operations::primary::is_dram(dst_buffer)}; + (std::uint32_t)CB::c_out0, (std::uint32_t)is_dram(dst_buffer)}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/kernels/reader_moreh_dot.cpp"; const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/kernels/writer_moreh_dot.cpp"; const auto reader_kernel_id = - tt::operations::primary::CreateReadKernel(program, reader_kernel_file, core, reader_compile_time_args); + CreateReadKernel(program, reader_kernel_file, core, reader_compile_time_args); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, core, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, core, writer_compile_time_args); std::vector compute_kernel_args = {}; std::map compute_defines; @@ -88,7 +88,7 @@ MorehDotOperation::SingleCore::cached_program_t MorehDotOperation::SingleCore::c const uint32_t core_num = 1; const auto compute_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_dot/device/kernels/moreh_dot.cpp"; - const auto compute_kernel_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_id = CreateComputeKernel( program, compute_kernel_file, {core, core_num, compute_kernel_args}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/moreh_dot_backward_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/moreh_dot_backward_device_operation.cpp index e088c1aeaa1..84e363bf3f9 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/moreh_dot_backward_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/moreh_dot_backward_device_operation.cpp @@ -29,10 +29,10 @@ void validate_tensors( const auto& input = tensor_args.input; const auto& other = tensor_args.other; - TT_FATAL(tt::operations::primary::is_scalar(output_grad), "Invalid value type"); - TT_FATAL(tt::operations::primary::is_1d_tensor(input), "Invalid input tensor dimensions."); - TT_FATAL(tt::operations::primary::is_1d_tensor(other), "Invalid input tensor dimensions."); - TT_FATAL(tt::operations::primary::is_same_shape(input, other), "Tensor A and B should have the same shape."); + TT_FATAL(is_scalar(output_grad), "Invalid value type"); + TT_FATAL(is_1d_tensor(input), "Invalid input tensor dimensions."); + TT_FATAL(is_1d_tensor(other), "Invalid input tensor dimensions."); + TT_FATAL(is_same_shape(input, other), "Tensor A and B should have the same shape."); TT_FATAL( input.get_dtype() == DataType::BFLOAT16 || input.get_dtype() == DataType::BFLOAT8_B, "Unsupported data format"); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/moreh_dot_backward_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/moreh_dot_backward_program_factory.cpp index 7bcd39d21bf..9c79145f8a8 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/moreh_dot_backward_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/moreh_dot_backward_program_factory.cpp @@ -48,7 +48,7 @@ MorehDotBackwardOperation::SingleCore::cached_program_t MorehDotBackwardOperatio const uint32_t out0_t = 2; const uint32_t out1_t = 2; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, std::set{CoreRange(core, core)}, cb_data_format, @@ -63,9 +63,9 @@ MorehDotBackwardOperation::SingleCore::cached_program_t MorehDotBackwardOperatio bool has_other_grad = other_grad.has_value(); std::vector reader_compile_time_args = { - (std::uint32_t)tt::operations::primary::is_dram(src0_buffer), - (std::uint32_t)tt::operations::primary::is_dram(src1_buffer), - (std::uint32_t)tt::operations::primary::is_dram(src2_buffer)}; + (std::uint32_t)is_dram(src0_buffer), + (std::uint32_t)is_dram(src1_buffer), + (std::uint32_t)is_dram(src2_buffer)}; bool dst0_is_dram = false; bool dst1_is_dram = false; @@ -76,7 +76,7 @@ MorehDotBackwardOperation::SingleCore::cached_program_t MorehDotBackwardOperatio const auto& input_grad_tensor = input_grad.value(); auto* dst0_buffer = input_grad_tensor.buffer(); TT_ASSERT(dst0_buffer != nullptr, "input_grad buffer should be allocated on device!"); - dst0_is_dram = tt::operations::primary::is_dram(dst0_buffer); + dst0_is_dram = is_dram(dst0_buffer); dst0_address = dst0_buffer->address(); } @@ -84,7 +84,7 @@ MorehDotBackwardOperation::SingleCore::cached_program_t MorehDotBackwardOperatio const auto& other_grad_tensor = other_grad.value(); auto* dst1_buffer = other_grad_tensor.buffer(); TT_ASSERT(dst1_buffer != nullptr, "other_grad buffer should be allocated on device!"); - dst1_is_dram = tt::operations::primary::is_dram(dst1_buffer); + dst1_is_dram = is_dram(dst1_buffer); dst1_address = dst1_buffer->address(); } @@ -101,16 +101,16 @@ MorehDotBackwardOperation::SingleCore::cached_program_t MorehDotBackwardOperatio "ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/kernels/writer_moreh_dot_backward.cpp"; const auto reader_kernel_id = - tt::operations::primary::CreateReadKernel(program, reader_kernel_file, core, reader_compile_time_args); + CreateReadKernel(program, reader_kernel_file, core, reader_compile_time_args); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, core, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, core, writer_compile_time_args); std::vector compute_kernel_args = {}; std::map compute_defines; const auto compute_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_dot_backward/device/kernels/moreh_dot_backward.cpp"; - const auto compute_kernel_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_id = CreateComputeKernel( program, compute_kernel_file, {core, core_num, compute_kernel_args}, compute_defines); SetRuntimeArgs( diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_fold/device/fold_program_factory_rm.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_fold/device/fold_program_factory_rm.cpp index 39f869d22f7..453bcded4b1 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_fold/device/fold_program_factory_rm.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_fold/device/fold_program_factory_rm.cpp @@ -119,9 +119,9 @@ MorehFoldOperation::ProgramFactory::cached_program_t MorehFoldOperation::Program const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_fold/device/kernels/reader_fold_rm.cpp"; const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_fold/device/kernels/writer_fold_rm.cpp"; - const auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + const auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); - const auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + const auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, writer_defines); //////////////////////////////////////////////////////////////////////////// diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_getitem/device/moreh_getitem_rm_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_getitem/device/moreh_getitem_rm_factory.cpp index c0f7e7ae8f4..266896792e8 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_getitem/device/moreh_getitem_rm_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_getitem/device/moreh_getitem_rm_factory.cpp @@ -23,7 +23,6 @@ MorehGetItemOperation::MorehGetItemRmFactory::cached_program_t MorehGetItemOpera tensor_return_value_t &output_tensor) { using namespace tt; using namespace tt::tt_metal; - using namespace tt::operations::primary; using namespace CMAKE_UNIQUE_NAMESPACE; auto input = tensor_args.input; @@ -88,7 +87,7 @@ MorehGetItemOperation::MorehGetItemRmFactory::cached_program_t MorehGetItemOpera uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_units_per_core_group_1, num_units_per_core_group_2] = - split_work_to_cores(core_range, num_units); + split_work_to_cores_wt_core_range(core_range, num_units); Program program = Program(); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_getitem/device/moreh_getitem_tilized_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_getitem/device/moreh_getitem_tilized_factory.cpp index e3c81d815c7..b9313d8e3e7 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_getitem/device/moreh_getitem_tilized_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_getitem/device/moreh_getitem_tilized_factory.cpp @@ -27,7 +27,6 @@ MorehGetItemOperation::MorehGetItemTilizedFactory::create( tensor_return_value_t &output_tensor) { using namespace tt; using namespace tt::tt_metal; - using namespace tt::operations::primary; using namespace CMAKE_UNIQUE_NAMESPACE; auto input = tensor_args.input; @@ -116,7 +115,7 @@ MorehGetItemOperation::MorehGetItemTilizedFactory::create( auto [num_cores, all_cores, core_group_1, core_group_2, num_units_per_core_group_1, num_units_per_core_group_2] = - split_work_to_cores(core_range, num_units); + split_work_to_cores_wt_core_range(core_range, num_units); Program program = Program(); @@ -349,7 +348,7 @@ MorehGetItemOperation::MorehGetItemTilizedFactory::create( auto [num_cores, all_cores, core_group_1, core_group_2, num_units_per_core_group_1, num_units_per_core_group_2] = - split_work_to_cores(core_range, num_units); + split_work_to_cores_wt_core_range(core_range, num_units); Program program = Program(); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm/device/moreh_group_norm_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm/device/moreh_group_norm_device_operation.cpp index 011a8150dd5..33e5bae29ac 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm/device/moreh_group_norm_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm/device/moreh_group_norm_device_operation.cpp @@ -21,8 +21,6 @@ void MorehGroupNormOperation::validate_tensors( auto num_groups = operation_attributes.num_groups; - using namespace tt::operations::primary; - check_tensor(input, "moreh_group_norm", "input"); check_tensor(output, "moreh_group_norm", "output"); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm/device/moreh_group_norm_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm/device/moreh_group_norm_program_factory.cpp index 30be36d8998..dee89cc0014 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm/device/moreh_group_norm_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm/device/moreh_group_norm_program_factory.cpp @@ -24,7 +24,6 @@ MorehGroupNormOperation::MorehGroupNormFactory::cached_program_t MorehGroupNormO tensor_return_value_t &outputs) { using namespace tt; using namespace tt::constants; - using namespace tt::operations::primary; const auto &input = tensor_args.input; auto gamma = tensor_args.gamma; diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/gamma_beta_grad/moreh_group_norm_backward_gamma_beta_grad_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/gamma_beta_grad/moreh_group_norm_backward_gamma_beta_grad_device_operation.cpp index 741a021d6c1..9b9b32661c1 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/gamma_beta_grad/moreh_group_norm_backward_gamma_beta_grad_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/gamma_beta_grad/moreh_group_norm_backward_gamma_beta_grad_device_operation.cpp @@ -24,8 +24,6 @@ void MorehGroupNormBackwardGammaBetaGradOperation::validate_tensors( auto num_groups = operation_attributes.num_groups; - using namespace tt::operations::primary; - check_tensor(output_grad, "moreh_group_norm_backward_gamma_beta_grad", "output_grad"); check_tensor(input, "moreh_group_norm_backward_gamma_beta_grad", "input"); check_tensor(mean, "moreh_group_norm_backward_gamma_beta_grad", "mean"); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/gamma_beta_grad/moreh_group_norm_backward_gamma_beta_grad_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/gamma_beta_grad/moreh_group_norm_backward_gamma_beta_grad_factory.cpp index 77e4d85ea9f..05a9e054090 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/gamma_beta_grad/moreh_group_norm_backward_gamma_beta_grad_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/gamma_beta_grad/moreh_group_norm_backward_gamma_beta_grad_factory.cpp @@ -14,7 +14,6 @@ MorehGroupNormBackwardGammaBetaGradOperation::MorehGroupNormBackwardGammaBetaGra tensor_return_value_t &outputs) { using namespace tt; using namespace tt::constants; - using namespace tt::operations::primary; const auto &output_grad = tensor_args.output_grad; const auto &input = tensor_args.input; diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/input_grad/moreh_group_norm_backward_input_grad_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/input_grad/moreh_group_norm_backward_input_grad_device_operation.cpp index 07612dfb1bc..bdcce3512e5 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/input_grad/moreh_group_norm_backward_input_grad_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/input_grad/moreh_group_norm_backward_input_grad_device_operation.cpp @@ -20,8 +20,6 @@ void MorehGroupNormBackwardInputGradOperation::validate_tensors( auto num_groups = operation_attributes.num_groups; - using namespace tt::operations::primary; - check_tensor(output_grad, "moreh_group_norm_backward_input_grad", "output_grad"); check_tensor(input, "moreh_group_norm_backward_input_grad", "input"); check_tensor(mean, "moreh_group_norm_backward_input_grad", "mean"); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/input_grad/moreh_group_norm_backward_input_grad_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/input_grad/moreh_group_norm_backward_input_grad_factory.cpp index d2eaa508ec0..5e0431048c8 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/input_grad/moreh_group_norm_backward_input_grad_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_group_norm_backward/device/input_grad/moreh_group_norm_backward_input_grad_factory.cpp @@ -16,7 +16,6 @@ MorehGroupNormBackwardInputGradOperation::MorehGroupNormBackwardInputGradFactory tensor_return_value_t& outputs) { using namespace tt; using namespace tt::constants; - using namespace tt::operations::primary; const auto& output_grad = tensor_args.output_grad; const auto& input = tensor_args.input; diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_helper_functions.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_helper_functions.cpp index f5322fce5eb..3ea46496d11 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_helper_functions.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_helper_functions.cpp @@ -4,15 +4,17 @@ #include "moreh_helper_functions.hpp" -#include "common/constants.hpp" #include -#include "tt_metal/detail/util.hpp" + +#include "common/constants.hpp" #include "tt_metal/common/work_split.hpp" +#include "tt_metal/detail/util.hpp" -namespace tt { +namespace ttnn { namespace operations { -namespace primary { +using namespace tt; +using namespace tt::tt_metal; using namespace constants; std::tuple add_core_offset( @@ -50,7 +52,7 @@ std::tuple add_core_offset( return std::make_tuple(new_all_cores, new_core_group_1, new_core_group_2); } -std::tuple split_work_to_cores( +std::tuple split_work_to_cores_wt_core_range( CoreRange core_range, uint32_t units_to_divide) { uint32_t core_w = core_range.end_coord.x - core_range.start_coord.x + 1; uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; @@ -99,7 +101,7 @@ std::tuple= rank - 2); -} +bool is_hw_dim(uint32_t dim, uint32_t rank) { return (dim >= rank - 2); } uint32_t compute_inner(tt::tt_metal::LegacyShape shape, uint32_t dim) { uint32_t num_inner = 1; @@ -440,7 +433,7 @@ ttnn::SmallVector get_dim( return dims; } -std::tuple extract_spatial_dims(const ttnn::SimpleShape& shape) { +std::tuple extract_spatial_dims(const ttnn::SimpleShape &shape) { const auto rank = shape.rank(); TT_FATAL(rank >= 2, "Shape must have at least two dims."); @@ -452,10 +445,11 @@ std::tuple extract_spatial_dims(const ttnn::Simple other_dims_product *= shape[i]; } - return { W, H, other_dims_product}; + return {W, H, other_dims_product}; } -std::tuple extract_and_scale_spatial_dims(const ttnn::SimpleShape& shape, uint32_t dim) { +std::tuple extract_and_scale_spatial_dims( + const ttnn::SimpleShape &shape, uint32_t dim) { const auto rank = shape.rank(); TT_FATAL(rank >= 2, "Shape must have at least two dims."); @@ -471,9 +465,8 @@ std::tuple extract_and_scale_spatial_dim uint32_t inner_tile_size = inner_dims_product * Ht * Wt; uint32_t reduce_tile_size = reduce_dim * inner_tile_size; - return { Wt, Ht, inner_tile_size, reduce_tile_size}; + return {Wt, Ht, inner_tile_size, reduce_tile_size}; } -} // namespace primary } // namespace operations -} // namespace tt +} // namespace ttnn diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_helper_functions.hpp b/ttnn/cpp/ttnn/operations/moreh/moreh_helper_functions.hpp index 689bc85e137..f341431a467 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_helper_functions.hpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_helper_functions.hpp @@ -12,9 +12,8 @@ #include "tt_metal/host_api.hpp" #include "ttnn/tensor/tensor.hpp" -namespace tt { +namespace ttnn { namespace operations { -namespace primary { using namespace tt::tt_metal; @@ -60,7 +59,7 @@ inline bool is_same_shape(const Tensor &tensor_a, const Tensor &tensor_b) { std::tuple add_core_offset( CoreRangeSet all_cores, CoreRangeSet core_group_1, CoreRangeSet core_group_2, uint32_t offset_x, uint32_t offset_y); -std::tuple split_work_to_cores( +std::tuple split_work_to_cores_wt_core_range( CoreRange core_range, uint32_t units_to_divide); [[maybe_unused]] KernelHandle CreateReadKernel( @@ -308,6 +307,5 @@ std::tuple extract_spatial_dims(const ttnn::Simple std::tuple extract_and_scale_spatial_dims(const ttnn::SimpleShape& shape, uint32_t dim); -} // namespace primary } // namespace operations } // namespace tt diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/moreh_layer_norm_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/moreh_layer_norm_device_operation.cpp index 7f1a68c57d8..735f5106eee 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/moreh_layer_norm_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/moreh_layer_norm_device_operation.cpp @@ -99,7 +99,7 @@ MorehLayerNormOperation::shape_return_value_t MorehLayerNormOperation::compute_o for (uint32_t dim = 0; dim < output_rank; dim++) { auto input_shape_without_padding_size = input_shape_without_padding[dim]; - if (tt::operations::primary::is_hw_dim(dim, output_rank)) { + if (is_hw_dim(dim, output_rank)) { output_size_vec.push_back(round_up_to_mul32(input_shape_without_padding_size)); auto padding_back = output_size_vec[dim] - input_shape_without_padding_size; dimensions_pads.push_back(Padding::PadDimension{.front = 0, .back = padding_back}); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/moreh_layer_norm_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/moreh_layer_norm_program_factory.cpp index d5c3e384f7c..a1062d4a275 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/moreh_layer_norm_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/moreh_layer_norm_program_factory.cpp @@ -71,8 +71,8 @@ MorehLayerNormOperation::ProgramFactory::cached_program_t MorehLayerNormOperatio const bool is_lastdim_layer_norm = normalized_dims == 1; const bool is_groupnorm = false; - auto num_inner = tt::operations::primary::compute_inner(input_shape, normalized_dims); - auto num_outer = tt::operations::primary::compute_outer(input_shape, normalized_dims); + auto num_inner = compute_inner(input_shape, normalized_dims); + auto num_outer = compute_outer(input_shape, normalized_dims); const auto gamma_has_value = gamma.has_value(); const auto beta_has_value = beta.has_value(); @@ -168,7 +168,7 @@ MorehLayerNormOperation::ProgramFactory::cached_program_t MorehLayerNormOperatio log_info(tt::LogTest, "Small moreh_layer_norm algorithm is selected."); } - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -197,15 +197,15 @@ MorehLayerNormOperation::ProgramFactory::cached_program_t MorehLayerNormOperatio // DataMovementKernel SetUp //////////////////////////////////////////////////////////////////////////// const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input)), - static_cast(tt::operations::primary::is_dram(gamma)), - static_cast(tt::operations::primary::is_dram(beta)), + static_cast(is_dram(input)), + static_cast(is_dram(gamma)), + static_cast(is_dram(beta)), block_size}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(output)), - static_cast(tt::operations::primary::is_dram(mean_as_tensor)), - static_cast(tt::operations::primary::is_dram(rstd_as_tensor)), + static_cast(is_dram(output)), + static_cast(is_dram(mean_as_tensor)), + static_cast(is_dram(rstd_as_tensor)), static_cast(mean_has_value), static_cast(rstd_has_value), block_size}; @@ -242,10 +242,10 @@ MorehLayerNormOperation::ProgramFactory::cached_program_t MorehLayerNormOperatio const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/kernels/writer_moreh_layer_norm.cpp"; - const auto reader_kernels_id = tt::operations::primary::CreateReadKernel( + const auto reader_kernels_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); const auto writer_kernels_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); const std::vector compute_args_group_1{ num_rows_per_core_group_1, @@ -265,7 +265,7 @@ MorehLayerNormOperation::ProgramFactory::cached_program_t MorehLayerNormOperatio ? "ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/kernels/moreh_layer_norm_large_kernel.cpp" : "ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm/device/kernels/moreh_layer_norm_small_kernel.cpp"; - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_rows_per_core_group_1, compute_args_group_1}, @@ -288,7 +288,7 @@ MorehLayerNormOperation::ProgramFactory::cached_program_t MorehLayerNormOperatio static_cast(is_lastdim_layer_norm), static_cast(is_groupnorm)}; - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_rows_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/moreh_layer_norm_backward_gamma_beta_grad_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/moreh_layer_norm_backward_gamma_beta_grad_program_factory.cpp index 6b686835649..565705ce44a 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/moreh_layer_norm_backward_gamma_beta_grad_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/moreh_layer_norm_backward_gamma_beta_grad_program_factory.cpp @@ -69,8 +69,8 @@ MorehLayerNormBackwardGammaBetaGradOperation::ProgramFactory::create( auto mean_rstd_height = mean_rstd_shape_without_padding[-2]; auto mean_rstd_width = mean_rstd_shape_without_padding[-1]; - auto num_inner = tt::operations::primary::compute_inner(output_grad_shape, normalized_dims); - auto num_outer = tt::operations::primary::compute_outer(output_grad_shape, normalized_dims); + auto num_inner = compute_inner(output_grad_shape, normalized_dims); + auto num_outer = compute_outer(output_grad_shape, normalized_dims); const bool gamma_grad_has_value = gamma_grad.has_value(); const bool beta_grad_has_value = beta_grad.has_value(); @@ -113,7 +113,7 @@ MorehLayerNormBackwardGammaBetaGradOperation::ProgramFactory::create( const auto cb_data_format = tt::tt_metal::datatype_to_dataformat_converter(output_grad.get_dtype()); auto intermed_cb_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : cb_data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -138,16 +138,16 @@ MorehLayerNormBackwardGammaBetaGradOperation::ProgramFactory::create( // DataMovementKernel SetUp //////////////////////////////////////////////////////////////////////////// const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(output_grad)), - static_cast(tt::operations::primary::is_dram(input)), - static_cast(tt::operations::primary::is_dram(mean)), - static_cast(tt::operations::primary::is_dram(rstd)), + static_cast(is_dram(output_grad)), + static_cast(is_dram(input)), + static_cast(is_dram(mean)), + static_cast(is_dram(rstd)), static_cast(gamma_grad_has_value), static_cast(do_mask_h)}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(gamma_grad)), - static_cast(tt::operations::primary::is_dram(beta_grad)), + static_cast(is_dram(gamma_grad)), + static_cast(is_dram(beta_grad)), static_cast(gamma_grad_has_value), static_cast(beta_grad_has_value)}; @@ -167,10 +167,10 @@ MorehLayerNormBackwardGammaBetaGradOperation::ProgramFactory::create( "ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/kernels/" "writer_moreh_layer_norm_backward_gamma_beta_grad.cpp"; - const auto reader_kernels_id = tt::operations::primary::CreateReadKernel( + const auto reader_kernels_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); const auto writer_kernels_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); const std::vector compute_args_group_1{ num_cols_per_core_group_1, @@ -187,7 +187,7 @@ MorehLayerNormBackwardGammaBetaGradOperation::ProgramFactory::create( "ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/kernels/" "moreh_layer_norm_backward_gamma_beta_grad_kernel.cpp"; - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, @@ -208,7 +208,7 @@ MorehLayerNormBackwardGammaBetaGradOperation::ProgramFactory::create( static_cast(is_lastdim_layer_norm), static_cast(is_groupnorm)}; - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/moreh_layer_norm_backward_input_grad_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/moreh_layer_norm_backward_input_grad_program_factory.cpp index e30eab1ea63..d423b24b60e 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/moreh_layer_norm_backward_input_grad_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/moreh_layer_norm_backward_input_grad_program_factory.cpp @@ -66,8 +66,8 @@ MorehLayerNormBackwardInputGradOperation::ProgramFactory::create( auto n = static_cast(normalized_numel); auto recip_n = 1.0f / n; - auto num_inner = tt::operations::primary::compute_inner(output_grad_shape, normalized_dims); - auto num_outer = tt::operations::primary::compute_outer(output_grad_shape, normalized_dims); + auto num_inner = compute_inner(output_grad_shape, normalized_dims); + auto num_outer = compute_outer(output_grad_shape, normalized_dims); const bool gamma_has_value = gamma.has_value(); @@ -129,7 +129,7 @@ MorehLayerNormBackwardInputGradOperation::ProgramFactory::create( log_info(tt::LogTest, "Small moreh_layer_norm_backward_input_grad algorithm is selected."); } - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -157,17 +157,17 @@ MorehLayerNormBackwardInputGradOperation::ProgramFactory::create( // DataMovementKernel SetUp //////////////////////////////////////////////////////////////////////////// const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(output_grad)), - static_cast(tt::operations::primary::is_dram(input)), - static_cast(tt::operations::primary::is_dram(mean)), - static_cast(tt::operations::primary::is_dram(rstd)), - static_cast(tt::operations::primary::is_dram(gamma)), + static_cast(is_dram(output_grad)), + static_cast(is_dram(input)), + static_cast(is_dram(mean)), + static_cast(is_dram(rstd)), + static_cast(is_dram(gamma)), static_cast(gamma_has_value), static_cast(do_mask_h), static_cast(do_mask_w)}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; std::map reader_defines{}; std::map compute_defines{}; @@ -191,10 +191,10 @@ MorehLayerNormBackwardInputGradOperation::ProgramFactory::create( "ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/kernels/" "writer_moreh_layer_norm_backward_input_grad.cpp"; - const auto reader_kernels_id = tt::operations::primary::CreateReadKernel( + const auto reader_kernels_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); const auto writer_kernels_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); const std::vector compute_args_group_1{ num_rows_per_core_group_1, @@ -211,7 +211,7 @@ MorehLayerNormBackwardInputGradOperation::ProgramFactory::create( : "ttnn/cpp/ttnn/operations/moreh/moreh_layer_norm_backward/device/kernels/" "moreh_layer_norm_backward_input_grad_small_kernel.cpp"; - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_rows_per_core_group_1, compute_args_group_1}, @@ -230,7 +230,7 @@ MorehLayerNormBackwardInputGradOperation::ProgramFactory::create( static_cast(is_lastdim_layer_norm), static_cast(is_groupnorm)}; - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_rows_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_device_operation.cpp index b150ce0ff4b..c64d02c4ef4 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_device_operation.cpp @@ -20,8 +20,8 @@ void MorehBiasAddBackwardOperation::validate_inputs( auto bias_grad_shape = bias_grad->get_shape(); auto bias_grad_tensor = bias_grad.value(); TT_FATAL( - tt::operations::primary::is_scalar(bias_grad_tensor) || - tt::operations::primary::is_1d_tensor(bias_grad_tensor), + is_scalar(bias_grad_tensor) || + is_1d_tensor(bias_grad_tensor), "bias_grad tensor should be 1d or scalar"); } } @@ -29,7 +29,7 @@ void MorehBiasAddBackwardOperation::validate_inputs( MorehBiasAddBackwardOperation::program_factory_t MorehBiasAddBackwardOperation::select_program_factory( const operation_attributes_t& operation_attributes, const tensor_args_t& tensor_args) { const auto& bias_grad = tensor_args.bias_grad.value(); - if (tt::operations::primary::is_scalar(bias_grad)) + if (is_scalar(bias_grad)) return SingleCoreProgramFactory(); return MultiCoreProgramFactory(); } diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_multi_core_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_multi_core_program_factory.cpp index 319e6d6e7f2..d1b503b689e 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_multi_core_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_multi_core_program_factory.cpp @@ -71,7 +71,7 @@ MorehBiasAddBackwardOperation::MultiCoreProgramFactory::create( const uint32_t im1_t = 1; auto cb_data_format = datatype_to_dataformat_converter(output_grad.get_dtype()); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -88,9 +88,9 @@ MorehBiasAddBackwardOperation::MultiCoreProgramFactory::create( const ::bfloat16 bfloat_scaler_value = ::bfloat16(1.0f); const uint32_t packed_scaler_value = pack_two_bfloat16_into_uint32({bfloat_scaler_value, bfloat_scaler_value}); const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(output_grad)), packed_scaler_value}; + static_cast(is_dram(output_grad)), packed_scaler_value}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(bias_grad))}; + static_cast(is_dram(bias_grad))}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/kernels/reader_moreh_bias_backward_h.cpp"; @@ -99,9 +99,9 @@ MorehBiasAddBackwardOperation::MultiCoreProgramFactory::create( "ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/kernels/writer_moreh_bias_backward.cpp"; const auto reader_kernel_id = - tt::operations::primary::CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); + CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -118,7 +118,7 @@ MorehBiasAddBackwardOperation::MultiCoreProgramFactory::create( const auto compute_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/kernels/moreh_bias_backward_multi_core_h.cpp"; - const auto compute_kernel_1_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_1_id = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, @@ -131,7 +131,7 @@ MorehBiasAddBackwardOperation::MultiCoreProgramFactory::create( std::optional compute_kernel_2_id = std::nullopt; if (!core_group_2.ranges().empty()) { const std::vector compute_args_group_2{num_cols_per_core_group_2}; - compute_kernel_2_id = tt::operations::primary::CreateComputeKernel( + compute_kernel_2_id = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_single_core_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_single_core_program_factory.cpp index d7ec2a28105..5f47d621149 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_single_core_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/moreh_linear_backward_single_core_program_factory.cpp @@ -65,7 +65,7 @@ MorehBiasAddBackwardOperation::SingleCoreProgramFactory::create( //////////////////////////////////////////////////////////////////////////// auto cb_data_format = datatype_to_dataformat_converter(output_grad.get_dtype()); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, std::set{CoreRange(core, core)}, cb_data_format, @@ -81,9 +81,9 @@ MorehBiasAddBackwardOperation::SingleCoreProgramFactory::create( //////////////////////////////////////////////////////////////////////////// const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(output_grad))}; + static_cast(is_dram(output_grad))}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(bias_grad))}; + static_cast(is_dram(bias_grad))}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/kernels/reader_moreh_bias_backward_hw.cpp"; @@ -92,9 +92,9 @@ MorehBiasAddBackwardOperation::SingleCoreProgramFactory::create( "ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/kernels/writer_moreh_bias_backward.cpp"; const auto reader_kernel_id = - tt::operations::primary::CreateReadKernel(program, reader_kernel_file, core, reader_compile_time_args); + CreateReadKernel(program, reader_kernel_file, core, reader_compile_time_args); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, core, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, core, writer_compile_time_args); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -110,7 +110,7 @@ MorehBiasAddBackwardOperation::SingleCoreProgramFactory::create( const auto compute_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/device/kernels/moreh_bias_backward_single_core_hw.cpp"; - const auto compute_kernel_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_id = CreateComputeKernel( program, compute_kernel_file, {core, core_num, compute_kernel_args}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/moreh_linear_backward.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/moreh_linear_backward.cpp index 08450c6d0ed..3e859d075fd 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/moreh_linear_backward.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_linear_backward/moreh_linear_backward.cpp @@ -47,21 +47,21 @@ inline void moreh_linear_backward_validate( if (input_grad.has_value()) { const auto& input_grad_tensor = input_grad.value(); TT_FATAL( - tt::operations::primary::is_same_shape(input, input_grad_tensor), "both tensors should be the same shape"); + is_same_shape(input, input_grad_tensor), "both tensors should be the same shape"); } if (weight_grad.has_value()) { const auto& weight_grad_tensor = weight_grad.value(); TT_FATAL( - tt::operations::primary::is_same_shape(weight, weight_grad_tensor), + is_same_shape(weight, weight_grad_tensor), "both tensors should be the same shape"); } if (bias_grad.has_value()) { const auto& bias_grad_tensor = bias_grad.value(); TT_FATAL( - tt::operations::primary::is_scalar(bias_grad_tensor) || - tt::operations::primary::is_1d_tensor(bias_grad_tensor), + is_scalar(bias_grad_tensor) || + is_1d_tensor(bias_grad_tensor), "bias_grad tensor should be 1d or scalar"); } } diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/moreh_matmul_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/moreh_matmul_device_operation.cpp index 6bcfa265181..d39f49e66dc 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/moreh_matmul_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/moreh_matmul_device_operation.cpp @@ -23,10 +23,10 @@ void MorehMatmulOperation::validate_inputs( const auto &output = tensor_args.output; // validate tensor - tt::operations::primary::check_tensor(input, "moreh_matmul", "input", {DataType::BFLOAT16}); - tt::operations::primary::check_tensor(other, "moreh_matmul", "other", {DataType::BFLOAT16}); - tt::operations::primary::check_tensor(output, "moreh_matmul", "output", {DataType::BFLOAT16}); - tt::operations::primary::check_tensor(bias, "moreh_matmul", "bias", {DataType::BFLOAT16}); + check_tensor(input, "moreh_matmul", "input", {DataType::BFLOAT16}); + check_tensor(other, "moreh_matmul", "other", {DataType::BFLOAT16}); + check_tensor(output, "moreh_matmul", "output", {DataType::BFLOAT16}); + check_tensor(bias, "moreh_matmul", "bias", {DataType::BFLOAT16}); // check matrix dims const auto &input_shape = input.get_shape().value.without_padding(); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/moreh_matmul_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/moreh_matmul_program_factory.cpp index 6a6cf2a0552..39b1d425b0a 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/moreh_matmul_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/moreh_matmul_program_factory.cpp @@ -283,7 +283,7 @@ MorehMatmulOperation::MultiCoreProgramFactory::cached_program_t MorehMatmulOpera const uint32_t im3_t{1}; // temp for bias add const uint32_t out0_t{2}; // output - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -305,8 +305,8 @@ MorehMatmulOperation::MultiCoreProgramFactory::cached_program_t MorehMatmulOpera //////////////////////////////////////////////////////////////////////////// std::map reader_defines; std::vector reader_compile_time_args = { - static_cast(tt::operations::primary::is_dram(input)), - static_cast(tt::operations::primary::is_dram(other)), + static_cast(is_dram(input)), + static_cast(is_dram(other)), Kt, static_cast(transpose_input), static_cast(transpose_other), @@ -318,36 +318,36 @@ MorehMatmulOperation::MultiCoreProgramFactory::cached_program_t MorehMatmulOpera if (bias.has_value()) { reader_defines["FUSE_BIAS"] = "1"; - reader_compile_time_args.push_back(static_cast(tt::operations::primary::is_dram(bias))); + reader_compile_time_args.push_back(static_cast(is_dram(bias))); reader_compile_time_args.push_back(static_cast(is_scalar_bias)); log_debug( tt::LogOp, "{}:{} bias tensor. is bias dram {}", __func__, __LINE__, - tt::operations::primary::is_dram(bias)); + is_dram(bias)); } const std::vector writer_compile_time_args = { - static_cast(tt::operations::primary::is_dram(output))}; + static_cast(is_dram(output))}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/kernels/reader_moreh_matmul.cpp"; const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_matmul/device/kernels/writer_moreh_matmul.cpp"; - const auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + const auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); log_debug( tt::LogOp, "{}:{} DMVK is_dram(input): {}, is_dram(other): {}, is_dram(output): {}", __func__, __LINE__, - tt::operations::primary::is_dram(input), - tt::operations::primary::is_dram(other), - tt::operations::primary::is_dram(output)); + is_dram(input), + is_dram(other), + is_dram(output)); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -378,7 +378,7 @@ MorehMatmulOperation::MultiCoreProgramFactory::cached_program_t MorehMatmulOpera unpack_to_dest_mode[tt::CB::c_intermed0] = UnpackToDestMode::UnpackToDestFp32; } - const auto compute_kernel_1_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_1_id = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_output_tiles_per_core_group_1, compute_args_group_1}, @@ -406,7 +406,7 @@ MorehMatmulOperation::MultiCoreProgramFactory::cached_program_t MorehMatmulOpera compute_args_group_2.push_back(static_cast(is_scalar_bias)); } - compute_kernel_2_id = tt::operations::primary::CreateComputeKernel( + compute_kernel_2_id = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_output_tiles_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/moreh_matmul.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/moreh_matmul.cpp index 2a2942a8ffe..a774a1c1e7d 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/moreh_matmul.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_matmul/moreh_matmul.cpp @@ -20,8 +20,8 @@ inline bool is_dot_forward(const Tensor& input, const Tensor& other, bool transp return false; } - return tt::operations::primary::is_1d_tensor(input) && tt::operations::primary::is_1d_tensor(other) && - tt::operations::primary::is_same_shape(input, other); + return is_1d_tensor(input) && is_1d_tensor(other) && + is_same_shape(input, other); } Tensor MorehMatmul::invoke( diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_matmul_backward/moreh_matmul_backward.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_matmul_backward/moreh_matmul_backward.cpp index cc71fd66c24..2b66b4459ec 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_matmul_backward/moreh_matmul_backward.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_matmul_backward/moreh_matmul_backward.cpp @@ -19,8 +19,8 @@ inline bool is_dot_backward(const Tensor& output_grad, const Tensor& input, cons other.get_legacy_shape().rank() != 4) { return false; } - return tt::operations::primary::is_scalar(output_grad) && tt::operations::primary::is_1d_tensor(input) && - tt::operations::primary::is_1d_tensor(other) && tt::operations::primary::is_same_shape(input, other); + return is_scalar(output_grad) && is_1d_tensor(input) && + is_1d_tensor(other) && is_same_shape(input, other); } std::vector> MorehMatmulBackward::invoke( diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_device_operation.cpp index 1c919f21ec1..84843f3afc5 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_device_operation.cpp @@ -20,13 +20,13 @@ void MorehMeanOperation::validate_tensors( operation_attributes.dim); TT_FATAL(operation_attributes.divisor.has_value() == false, "divisor not supported yet."); - tt::operations::primary::check_tensor(input, "moreh_mean", "input", {DataType::BFLOAT16}); - tt::operations::primary::check_tensor(output, "moreh_mean", "output", {DataType::BFLOAT16}); + check_tensor(input, "moreh_mean", "input", {DataType::BFLOAT16}); + check_tensor(output, "moreh_mean", "output", {DataType::BFLOAT16}); - tt::operations::primary::validate_input_with_dim(input, operation_attributes.dim); + validate_input_with_dim(input, operation_attributes.dim); if (output.has_value()) { - tt::operations::primary::validate_output_with_keepdim( + validate_output_with_keepdim( input, output.value(), operation_attributes.dim, operation_attributes.keepdim); } } diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_h_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_h_program_factory.cpp index 0648b2ba59d..08ea9596f4f 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_h_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_h_program_factory.cpp @@ -19,7 +19,6 @@ MorehMeanOperation::MorehMeanHFactory::cached_program_t MorehMeanOperation::More tensor_return_value_t& output) { using namespace tt; using namespace tt::tt_metal; - using namespace tt::operations::primary; auto input = tensor_args.input; auto compute_kernel_config = @@ -52,7 +51,7 @@ MorehMeanOperation::MorehMeanHFactory::cached_program_t MorehMeanOperation::More uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, units_per_core_group_1, units_per_core_group_2] = - split_work_to_cores(core_range, units_to_divide); + split_work_to_cores_wt_core_range(core_range, units_to_divide); auto arch = input.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_nc_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_nc_program_factory.cpp index 0bd21b787a5..e79b5af1768 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_nc_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_nc_program_factory.cpp @@ -19,7 +19,6 @@ MorehMeanOperation::MorehMeanNCFactory::cached_program_t MorehMeanOperation::Mor tensor_return_value_t& output) { using namespace tt; using namespace tt::tt_metal; - using namespace tt::operations::primary; auto input = tensor_args.input; auto dim = operation_attributes.dim; @@ -67,7 +66,7 @@ MorehMeanOperation::MorehMeanNCFactory::cached_program_t MorehMeanOperation::Mor uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, units_per_core_group_1, units_per_core_group_2] = - split_work_to_cores(core_range, units_to_divide); + split_work_to_cores_wt_core_range(core_range, units_to_divide); auto arch = input.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_w_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_w_program_factory.cpp index 5aed254b1c9..baba5737d47 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_w_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/device/moreh_mean_w_program_factory.cpp @@ -19,7 +19,6 @@ MorehMeanOperation::MorehMeanWFactory::cached_program_t MorehMeanOperation::More tensor_return_value_t& output) { using namespace tt; using namespace tt::tt_metal; - using namespace tt::operations::primary; auto input = tensor_args.input; auto compute_kernel_config = @@ -51,7 +50,7 @@ MorehMeanOperation::MorehMeanWFactory::cached_program_t MorehMeanOperation::More auto units_to_divide = input.volume() / W / H * Ht; auto [num_cores, all_cores, core_group_1, core_group_2, units_per_core_group_1, units_per_core_group_2] = - split_work_to_cores(core_range, units_to_divide); + split_work_to_cores_wt_core_range(core_range, units_to_divide); auto arch = input.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/moreh_mean.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/moreh_mean.cpp index 7b94384eab6..b9cf174b85a 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_mean/moreh_mean.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_mean/moreh_mean.cpp @@ -16,7 +16,7 @@ Tensor MorehMean::invoke( const std::optional& output, const std::optional& memory_config, const std::optional& compute_kernel_config) { - ttnn::SmallVector dims = tt::operations::primary::get_dim(dim, input.get_shape().rank()); + ttnn::SmallVector dims = get_dim(dim, input.get_shape().rank()); std::sort(dims.begin(), dims.end()); auto temp_input = input; diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/moreh_mean_backward_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/moreh_mean_backward_device_operation.cpp index c989b3675de..56a1ebd2e42 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/moreh_mean_backward_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/moreh_mean_backward_device_operation.cpp @@ -17,8 +17,8 @@ void MorehMeanBackwardOperation::validate_tensors( input_grad.has_value() || operation_attributes.input_grad_shape.has_value() || operation_attributes.keepdim, "Either input_grad tensor or input_grad_shape or keepdim must be present"); - tt::operations::primary::check_tensor(output_grad, "moreh_mean_backward", "output_grad", {DataType::BFLOAT16}); - tt::operations::primary::check_tensor(input_grad, "moreh_mean_backward", "input_grad", {DataType::BFLOAT16}); + check_tensor(output_grad, "moreh_mean_backward", "output_grad", {DataType::BFLOAT16}); + check_tensor(input_grad, "moreh_mean_backward", "input_grad", {DataType::BFLOAT16}); } MorehMeanBackwardOperation::program_factory_t MorehMeanBackwardOperation::select_program_factory( @@ -45,7 +45,7 @@ MorehMeanBackwardOperation::shape_return_value_t MorehMeanBackwardOperation::com ttnn::SmallVector dimensions_pads; for (uint32_t dim = 0; dim < rank; dim++) { - if (tt::operations::primary::is_hw_dim(dim, rank)) { + if (is_hw_dim(dim, rank)) { uint32_t up32_shape = tt::round_up(input_grad_shape[dim], 32); uint32_t padding_back = up32_shape - input_grad_shape[dim]; shape.push_back(up32_shape); @@ -70,7 +70,7 @@ MorehMeanBackwardOperation::tensor_return_value_t MorehMeanBackwardOperation::cr return tensor_args.input_grad.value(); } - return tt::operations::primary::create_device_tensor( + return create_device_tensor( compute_output_shapes(operation_attributes, tensor_args), output_grad.get_dtype(), Layout::TILE, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/moreh_mean_backward_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/moreh_mean_backward_program_factory.cpp index 02fd35ab4c9..46ab6465658 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/moreh_mean_backward_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/moreh_mean_backward_program_factory.cpp @@ -119,7 +119,7 @@ MorehMeanBackwardOperation::MorehMeanBackwardFactory::create( //////////////////////////////////////////////////////////////////////////// // CircularBuffer Setup //////////////////////////////////////////////////////////////////////////// - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -135,17 +135,17 @@ MorehMeanBackwardOperation::MorehMeanBackwardFactory::create( // DataMovementKernel SetUp //////////////////////////////////////////////////////////////////////////// std::vector reader_compile_time_args = { - static_cast(tt::operations::primary::is_dram(output_grad)), input_grad_rank}; + static_cast(is_dram(output_grad)), input_grad_rank}; std::vector writer_compile_time_args = { - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/kernels/reader_moreh_mean_backward.cpp"; const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/device/kernels/writer_moreh_mean_backward.cpp"; const auto reader_kernel_id = - tt::operations::primary::CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); + CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -160,14 +160,14 @@ MorehMeanBackwardOperation::MorehMeanBackwardFactory::create( const std::vector compute_args_group_1{num_cols_per_core_group_1, need_bcast_dim[0], need_bcast_dim[1]}; const std::vector compute_args_group_2{num_cols_per_core_group_2, need_bcast_dim[0], need_bcast_dim[1]}; std::vector unpack_to_dest_mode(NUM_CIRCULAR_BUFFERS, UnpackToDestMode::Default); - auto compute_kernel_ids = tt::operations::primary::CreateComputeKernel( + auto compute_kernel_ids = CreateComputeKernel( program, compute_kernel_file, { {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, }, - tt::operations::primary::ComputeKernelConfig{ + ComputeKernelConfig{ .math_fidelity = math_fidelity, .fp32_dest_acc_en = fp32_dest_acc_en, .unpack_to_dest_mode = unpack_to_dest_mode, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/moreh_mean_backward.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/moreh_mean_backward.cpp index 88cb5946c56..6a26c25a6e3 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/moreh_mean_backward.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_mean_backward/moreh_mean_backward.cpp @@ -28,7 +28,7 @@ Tensor MorehMeanBackward::invoke( input_grad_rank += dims.size(); } } - ttnn::SmallVector dims = tt::operations::primary::get_dim(dim, input_grad_rank); + ttnn::SmallVector dims = get_dim(dim, input_grad_rank); return ttnn::prim::moreh_mean_backward( output_grad, dims, keepdim, input_grad_shape, input_grad, memory_config, compute_kernel_config); } diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step1/device/moreh_nll_loss_step1_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step1/device/moreh_nll_loss_step1_program_factory.cpp index a552aa2eb05..c95dcf8eefe 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step1/device/moreh_nll_loss_step1_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step1/device/moreh_nll_loss_step1_program_factory.cpp @@ -75,7 +75,7 @@ MorehNllLossStep1DeviceOperation::Factory::cached_program_t MorehNllLossStep1Dev const bool use_large_algorithm = cb_usage >= available_L1; if (use_large_algorithm) { - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -86,7 +86,7 @@ MorehNllLossStep1DeviceOperation::Factory::cached_program_t MorehNllLossStep1Dev {CB::c_out0, 1}, // output }); } else { - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -100,12 +100,12 @@ MorehNllLossStep1DeviceOperation::Factory::cached_program_t MorehNllLossStep1Dev // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(target)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false), + static_cast(is_dram(target)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false), static_cast(weight_has_value)}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(output))}; + static_cast(is_dram(output))}; std::map reader_defines; std::map writer_defines; @@ -126,9 +126,9 @@ MorehNllLossStep1DeviceOperation::Factory::cached_program_t MorehNllLossStep1Dev "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step1/device/kernels/" "writer_moreh_nll_loss_step1.cpp"; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, writer_defines); const auto target_addr = target.buffer()->address(); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/moreh_nll_loss_step2_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/moreh_nll_loss_step2_program_factory.cpp index 608e82b1b57..f544821fb5d 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/moreh_nll_loss_step2_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/moreh_nll_loss_step2_program_factory.cpp @@ -56,7 +56,7 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 auto fp32_dest_acc_en_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -75,14 +75,14 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input)), - static_cast(tt::operations::primary::is_dram(target)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false), - static_cast(divisor.has_value() ? tt::operations::primary::is_dram(divisor.value()) : false), + static_cast(is_dram(input)), + static_cast(is_dram(target)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false), + static_cast(divisor.has_value() ? is_dram(divisor.value()) : false), }; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(output))}; + static_cast(is_dram(output))}; std::map reader_defines; std::map writer_defines; @@ -102,14 +102,14 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 compute_defines["FP32_DEST_ACC_EN"] = 1; } - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/kernels/" "reader_moreh_nll_loss_step2_2d.cpp", all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/kernels/" "writer_moreh_nll_loss_step2_2d.cpp", @@ -117,7 +117,7 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 writer_compile_time_args, writer_defines); - const auto compute_kernel_ids = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_ids = CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/kernels/" "moreh_nll_loss_step2_kernel.cpp", @@ -235,7 +235,7 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 auto fp32_dest_acc_en_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -254,14 +254,14 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input)), - static_cast(tt::operations::primary::is_dram(target)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false), - static_cast(divisor.has_value() ? tt::operations::primary::is_dram(divisor.value()) : false), + static_cast(is_dram(input)), + static_cast(is_dram(target)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false), + static_cast(divisor.has_value() ? is_dram(divisor.value()) : false), }; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(output))}; + static_cast(is_dram(output))}; std::map reader_defines; std::map writer_defines; @@ -281,14 +281,14 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 compute_defines["FP32_DEST_ACC_EN"] = 1; } - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/kernels/" "reader_moreh_nll_loss_step2_3d.cpp", all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/kernels/" "writer_moreh_nll_loss_step2_3d.cpp", @@ -296,7 +296,7 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 writer_compile_time_args, writer_defines); - const auto compute_kernel_ids = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_ids = CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/kernels/" "moreh_nll_loss_step2_kernel.cpp", @@ -424,7 +424,7 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 auto fp32_dest_acc_en_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; uint32_t weight_num_tile = div_up(channel_size, tt::constants::TILE_WIDTH); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -443,14 +443,14 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input)), - static_cast(tt::operations::primary::is_dram(target)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false), - static_cast(divisor.has_value() ? tt::operations::primary::is_dram(divisor.value()) : false), + static_cast(is_dram(input)), + static_cast(is_dram(target)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false), + static_cast(divisor.has_value() ? is_dram(divisor.value()) : false), }; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(output))}; + static_cast(is_dram(output))}; std::map reader_defines; std::map writer_defines; @@ -470,14 +470,14 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 compute_defines["FP32_DEST_ACC_EN"] = 1; } - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/kernels/" "reader_moreh_nll_loss_step2_4d.cpp", all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/kernels/" "writer_moreh_nll_loss_step2_4d.cpp", @@ -485,7 +485,7 @@ MorehNllLossStep2DeviceOperation::Factory::cached_program_t moreh_nll_loss_step2 writer_compile_time_args, writer_defines); - const auto compute_kernel_ids = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_ids = CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss/moreh_nll_loss_step2/device/kernels/" "moreh_nll_loss_step2_kernel.cpp", diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_backward/device/moreh_nll_loss_backward_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_backward/device/moreh_nll_loss_backward_program_factory.cpp index 4042457a541..830470d5c48 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_backward/device/moreh_nll_loss_backward_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_backward/device/moreh_nll_loss_backward_program_factory.cpp @@ -54,7 +54,7 @@ MorehNllLossBackwardDeviceOperation::Factory::cached_program_t moreh_nll_loss_ba auto fp32_dest_acc_en_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; uint32_t weight_num_tile = tt::div_up(channel_size, tt::constants::TILE_WIDTH); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -71,13 +71,13 @@ MorehNllLossBackwardDeviceOperation::Factory::cached_program_t moreh_nll_loss_ba // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(target)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false), - static_cast(divisor.has_value() ? tt::operations::primary::is_dram(divisor.value()) : false), - static_cast(tt::operations::primary::is_dram(output_grad))}; + static_cast(is_dram(target)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false), + static_cast(divisor.has_value() ? is_dram(divisor.value()) : false), + static_cast(is_dram(output_grad))}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; std::map reader_defines; std::map writer_defines; @@ -107,12 +107,12 @@ MorehNllLossBackwardDeviceOperation::Factory::cached_program_t moreh_nll_loss_ba "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_backward/device/kernels/" "moreh_nll_loss_backward_kernel.cpp"; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, writer_defines); - const auto compute_kernel_ids = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_ids = CreateComputeKernel( program, compute_kernel_file, { @@ -230,7 +230,7 @@ MorehNllLossBackwardDeviceOperation::Factory::cached_program_t moreh_nll_loss_ba auto fp32_dest_acc_en_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; uint32_t weight_num_tile = tt::div_up(channel_size, tt::constants::TILE_WIDTH); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -247,13 +247,13 @@ MorehNllLossBackwardDeviceOperation::Factory::cached_program_t moreh_nll_loss_ba // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(target)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false), - static_cast(divisor.has_value() ? tt::operations::primary::is_dram(divisor.value()) : false), - static_cast(tt::operations::primary::is_dram(output_grad))}; + static_cast(is_dram(target)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false), + static_cast(divisor.has_value() ? is_dram(divisor.value()) : false), + static_cast(is_dram(output_grad))}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; std::map reader_defines; std::map writer_defines; @@ -283,12 +283,12 @@ MorehNllLossBackwardDeviceOperation::Factory::cached_program_t moreh_nll_loss_ba "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_backward/device/kernels/" "moreh_nll_loss_backward_kernel.cpp"; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, writer_defines); - const auto compute_kernel_ids = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_ids = CreateComputeKernel( program, compute_kernel_file, { @@ -404,7 +404,7 @@ MorehNllLossBackwardDeviceOperation::Factory::cached_program_t moreh_nll_loss_ba auto fp32_dest_acc_en_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; uint32_t weight_num_tile = tt::div_up(channel_size, tt::constants::TILE_WIDTH); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -421,13 +421,13 @@ MorehNllLossBackwardDeviceOperation::Factory::cached_program_t moreh_nll_loss_ba // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(target)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false), - static_cast(divisor.has_value() ? tt::operations::primary::is_dram(divisor.value()) : false), - static_cast(tt::operations::primary::is_dram(output_grad))}; + static_cast(is_dram(target)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false), + static_cast(divisor.has_value() ? is_dram(divisor.value()) : false), + static_cast(is_dram(output_grad))}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; std::map reader_defines; std::map writer_defines; @@ -457,12 +457,12 @@ MorehNllLossBackwardDeviceOperation::Factory::cached_program_t moreh_nll_loss_ba "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_backward/device/kernels/" "moreh_nll_loss_backward_kernel.cpp"; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, writer_defines); - const auto compute_kernel_ids = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_ids = CreateComputeKernel( program, compute_kernel_file, { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_unreduced_backward/device/moreh_nll_loss_unreduced_backward_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_unreduced_backward/device/moreh_nll_loss_unreduced_backward_program_factory.cpp index 4d313d324d6..07b1f972483 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_unreduced_backward/device/moreh_nll_loss_unreduced_backward_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_unreduced_backward/device/moreh_nll_loss_unreduced_backward_program_factory.cpp @@ -50,7 +50,7 @@ MorehNllLossUnreducedBackwardDeviceOperation::Factory::cached_program_t moreh_nl auto Ct = tt::div_up(channel_size, tt::constants::TILE_WIDTH); auto Nt = tt::div_up(N, tt::constants::TILE_WIDTH); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -63,12 +63,12 @@ MorehNllLossUnreducedBackwardDeviceOperation::Factory::cached_program_t moreh_nl // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(target)), - static_cast(tt::operations::primary::is_dram(output_grad)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false)}; + static_cast(is_dram(target)), + static_cast(is_dram(output_grad)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false)}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; std::map reader_defines; std::map writer_defines; @@ -88,9 +88,9 @@ MorehNllLossUnreducedBackwardDeviceOperation::Factory::cached_program_t moreh_nl "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_unreduced_backward/device/kernels/" "writer_moreh_nll_loss_unreduced_backward.cpp"; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, writer_defines); const auto target_addr = target.buffer()->address(); @@ -178,7 +178,7 @@ MorehNllLossUnreducedBackwardDeviceOperation::Factory::cached_program_t moreh_nl // create circular buffers tt::DataFormat data_format = tt::tt_metal::datatype_to_dataformat_converter(input_grad.get_dtype()); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -191,12 +191,12 @@ MorehNllLossUnreducedBackwardDeviceOperation::Factory::cached_program_t moreh_nl // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(target)), - static_cast(tt::operations::primary::is_dram(output_grad)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false)}; + static_cast(is_dram(target)), + static_cast(is_dram(output_grad)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false)}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; std::map reader_defines; std::map writer_defines; @@ -216,9 +216,9 @@ MorehNllLossUnreducedBackwardDeviceOperation::Factory::cached_program_t moreh_nl "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_unreduced_backward/device/kernels/" "writer_moreh_nll_loss_unreduced_backward.cpp"; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, writer_defines); const auto target_addr = target.buffer()->address(); @@ -305,7 +305,7 @@ MorehNllLossUnreducedBackwardDeviceOperation::Factory::cached_program_t moreh_nl // create circular buffers tt::DataFormat data_format = tt::tt_metal::datatype_to_dataformat_converter(input_grad.get_dtype()); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -318,12 +318,12 @@ MorehNllLossUnreducedBackwardDeviceOperation::Factory::cached_program_t moreh_nl // create read/wrtie kernel const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(target)), - static_cast(tt::operations::primary::is_dram(output_grad)), - static_cast(weight.has_value() ? tt::operations::primary::is_dram(weight.value()) : false)}; + static_cast(is_dram(target)), + static_cast(is_dram(output_grad)), + static_cast(weight.has_value() ? is_dram(weight.value()) : false)}; const std::vector writer_compile_time_args{ - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; std::map reader_defines; std::map writer_defines; @@ -343,9 +343,9 @@ MorehNllLossUnreducedBackwardDeviceOperation::Factory::cached_program_t moreh_nl "ttnn/cpp/ttnn/operations/moreh/moreh_nll_loss_unreduced_backward/device/kernels/" "writer_moreh_nll_loss_unreduced_backward.cpp"; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, writer_defines); const auto target_addr = target.buffer()->address(); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_device_operation.cpp index f79e228fb3e..ff20005b8d1 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_device_operation.cpp @@ -51,10 +51,10 @@ inline void validate_output_tensor_with_keepdim(const Tensor& input, const Tenso ttnn::SmallVector input_dim_wo_padding(tt::tt_metal::MAX_NUM_DIMENSIONS, 1); ttnn::SmallVector output_dim_wo_padding(tt::tt_metal::MAX_NUM_DIMENSIONS, 1); - tt::operations::primary::expand_to_max_dim(input_dim, adjusted_input_shape); - tt::operations::primary::expand_to_max_dim(output_dim, output_shape); - tt::operations::primary::expand_to_max_dim(input_dim_wo_padding, adjusted_input_shape_wo_padding); - tt::operations::primary::expand_to_max_dim(output_dim_wo_padding, output_shape_wo_padding); + expand_to_max_dim(input_dim, adjusted_input_shape); + expand_to_max_dim(output_dim, output_shape); + expand_to_max_dim(input_dim_wo_padding, adjusted_input_shape_wo_padding); + expand_to_max_dim(output_dim_wo_padding, output_shape_wo_padding); for (int i = 0; i < input_rank; ++i) { TT_FATAL(input_dim[i] == output_dim[i], "Input and output dimensions do not match at index {}.", i); @@ -97,8 +97,8 @@ void MorehNormOperation::validate_inputs( const auto& input = tensor_args.input; const auto& output = tensor_args.output; const auto dim = operation_attributes.dim; - tt::operations::primary::check_tensor(input, "moreh_norm", "input"); - tt::operations::primary::check_tensor(output, "moreh_norm", "output"); + check_tensor(input, "moreh_norm", "input"); + check_tensor(output, "moreh_norm", "output"); validate_input_tensor_with_dim(input, dim); if (output.has_value()) validate_output_tensor_with_keepdim(input, output.value(), dim, operation_attributes.keepdim); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_h.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_h.cpp index c9ca5c1fd8a..b92d76192c5 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_h.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_h.cpp @@ -80,7 +80,7 @@ MorehNormOperation::ProgramFactoryH::cached_program_t MorehNormOperation::Progra const uint32_t im5_t{1}; // Add(|x + decimal|^p) const uint32_t im6_t{1}; // Sum(|x + decimal|^p) - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -110,8 +110,8 @@ MorehNormOperation::ProgramFactoryH::cached_program_t MorehNormOperation::Progra "ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_h/kernels/" "writer_moreh_norm_h.cpp"; - const auto reader_kernels_id = tt::operations::primary::CreateReadKernel(program, reader_kernel_file, all_cores); - const auto writer_kernels_id = tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores); + const auto reader_kernels_id = CreateReadKernel(program, reader_kernel_file, all_cores); + const auto writer_kernels_id = CreateWriteKernel(program, writer_kernel_file, all_cores); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -124,7 +124,7 @@ MorehNormOperation::ProgramFactoryH::cached_program_t MorehNormOperation::Progra "ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_h/kernels/" "moreh_norm_h_kernel.cpp"; - const auto compute_kernels_id_1 = tt::operations::primary::CreateComputeKernel( + const auto compute_kernels_id_1 = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_units_per_core_group_1}, @@ -135,7 +135,7 @@ MorehNormOperation::ProgramFactoryH::cached_program_t MorehNormOperation::Progra KernelHandle compute_kernels_id_2{0}; if (!core_group_2.ranges().empty()) { - compute_kernels_id_2 = tt::operations::primary::CreateComputeKernel( + compute_kernels_id_2 = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_units_per_core_group_2}, @@ -166,7 +166,7 @@ MorehNormOperation::ProgramFactoryH::cached_program_t MorehNormOperation::Progra // reader const std::vector reader_runtime_args{ input.buffer()->address(), - static_cast(tt::operations::primary::is_dram(input)), + static_cast(is_dram(input)), *reinterpret_cast(&decimal), *reinterpret_cast(&recip_p_decimal), num_cols_per_core, @@ -179,7 +179,7 @@ MorehNormOperation::ProgramFactoryH::cached_program_t MorehNormOperation::Progra // writer const std::vector writer_runtime_args{ output.buffer()->address(), - static_cast(tt::operations::primary::is_dram(output)), + static_cast(is_dram(output)), num_cols_per_core, tile_offset}; SetRuntimeArgs(program, writer_kernels_id, core, writer_runtime_args); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_other.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_other.cpp index 47374a248a6..22cdada3016 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_other.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_other.cpp @@ -90,7 +90,7 @@ MorehNormOperation::ProgramFactoryOther::cached_program_t MorehNormOperation::Pr const uint32_t im4_t{1}; // |x|^p * exp(log(|x|) * decimal) == |x + decimal|^p const uint32_t im5_t{1}; // Add(|x + decimal|^p) - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -118,8 +118,8 @@ MorehNormOperation::ProgramFactoryOther::cached_program_t MorehNormOperation::Pr "ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_other/kernels/" "writer_moreh_norm_other.cpp"; - const auto reader_kernels_id = tt::operations::primary::CreateReadKernel(program, reader_kernel_file, all_cores); - const auto writer_kernels_id = tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores); + const auto reader_kernels_id = CreateReadKernel(program, reader_kernel_file, all_cores); + const auto writer_kernels_id = CreateWriteKernel(program, writer_kernel_file, all_cores); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -130,7 +130,7 @@ MorehNormOperation::ProgramFactoryOther::cached_program_t MorehNormOperation::Pr "ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_other/kernels/" "moreh_norm_other_kernel.cpp"; - const auto compute_kernels_id_1 = tt::operations::primary::CreateComputeKernel( + const auto compute_kernels_id_1 = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_units_per_core_group_1}, @@ -141,7 +141,7 @@ MorehNormOperation::ProgramFactoryOther::cached_program_t MorehNormOperation::Pr KernelHandle compute_kernels_id_2{0}; if (!core_group_2.ranges().empty()) { - compute_kernels_id_2 = tt::operations::primary::CreateComputeKernel( + compute_kernels_id_2 = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_units_per_core_group_2}, @@ -172,7 +172,7 @@ MorehNormOperation::ProgramFactoryOther::cached_program_t MorehNormOperation::Pr // reader const std::vector reader_runtime_args{ input.buffer()->address(), - static_cast(tt::operations::primary::is_dram(input)), + static_cast(is_dram(input)), *reinterpret_cast(&decimal), *reinterpret_cast(&recip_p_decimal), num_output_tiles_per_core, @@ -185,7 +185,7 @@ MorehNormOperation::ProgramFactoryOther::cached_program_t MorehNormOperation::Pr // writer const std::vector writer_runtime_args{ output.buffer()->address(), - static_cast(tt::operations::primary::is_dram(output)), + static_cast(is_dram(output)), num_output_tiles_per_core, tile_offset}; SetRuntimeArgs(program, writer_kernels_id, core, writer_runtime_args); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_w.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_w.cpp index bb5ca9be017..912d2c1542a 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_w.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_program_factory_w.cpp @@ -80,7 +80,7 @@ MorehNormOperation::ProgramFactoryW::cached_program_t MorehNormOperation::Progra const uint32_t im5_t{1}; // Add(|x + decimal|^p) const uint32_t im6_t{1}; // Sum(|x + decimal|^p) - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -110,8 +110,8 @@ MorehNormOperation::ProgramFactoryW::cached_program_t MorehNormOperation::Progra "ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_w/kernels/" "writer_moreh_norm_w.cpp"; - const auto reader_kernels_id = tt::operations::primary::CreateReadKernel(program, reader_kernel_file, all_cores); - const auto writer_kernels_id = tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores); + const auto reader_kernels_id = CreateReadKernel(program, reader_kernel_file, all_cores); + const auto writer_kernels_id = CreateWriteKernel(program, writer_kernel_file, all_cores); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -124,7 +124,7 @@ MorehNormOperation::ProgramFactoryW::cached_program_t MorehNormOperation::Progra "ttnn/cpp/ttnn/operations/moreh/moreh_norm/device/moreh_norm_w/kernels/" "moreh_norm_w_kernel.cpp"; - const auto compute_kernels_id_1 = tt::operations::primary::CreateComputeKernel( + const auto compute_kernels_id_1 = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_units_per_core_group_1}, @@ -135,7 +135,7 @@ MorehNormOperation::ProgramFactoryW::cached_program_t MorehNormOperation::Progra KernelHandle compute_kernels_id_2{0}; if (!core_group_2.ranges().empty()) { - compute_kernels_id_2 = tt::operations::primary::CreateComputeKernel( + compute_kernels_id_2 = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_units_per_core_group_2}, @@ -166,7 +166,7 @@ MorehNormOperation::ProgramFactoryW::cached_program_t MorehNormOperation::Progra // reader const std::vector reader_runtime_args{ input.buffer()->address(), - static_cast(tt::operations::primary::is_dram(input)), + static_cast(is_dram(input)), *reinterpret_cast(&decimal), *reinterpret_cast(&recip_p_decimal), num_units_per_core, @@ -178,7 +178,7 @@ MorehNormOperation::ProgramFactoryW::cached_program_t MorehNormOperation::Progra // writer const std::vector writer_runtime_args{ output.buffer()->address(), - static_cast(tt::operations::primary::is_dram(output)), + static_cast(is_dram(output)), num_units_per_core, Wt, tile_offset}; diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_norm_backward/device/moreh_norm_backward_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_norm_backward/device/moreh_norm_backward_device_operation.cpp index 6013b1a0b08..a289d1a8396 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_norm_backward/device/moreh_norm_backward_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_norm_backward/device/moreh_norm_backward_device_operation.cpp @@ -11,10 +11,10 @@ namespace ttnn::operations::moreh::moreh_norm_backward { void MorehNormBackwardOperation::validate_inputs( const operation_attributes_t& operation_attributes, const tensor_args_t& tensor_args) { - tt::operations::primary::check_tensor(tensor_args.input, "moreh_norm_backward", "input"); - tt::operations::primary::check_tensor(tensor_args.output, "moreh_norm_backward", "output"); - tt::operations::primary::check_tensor(tensor_args.output_grad, "moreh_norm_backward", "output_grad"); - tt::operations::primary::check_tensor(tensor_args.input_grad, "moreh_norm_backward", "input_grad"); + check_tensor(tensor_args.input, "moreh_norm_backward", "input"); + check_tensor(tensor_args.output, "moreh_norm_backward", "output"); + check_tensor(tensor_args.output_grad, "moreh_norm_backward", "output_grad"); + check_tensor(tensor_args.input_grad, "moreh_norm_backward", "input_grad"); } MorehNormBackwardOperation::program_factory_t MorehNormBackwardOperation::select_program_factory( @@ -61,7 +61,7 @@ MorehNormBackwardOperation::invoke( const std::optional& input_grad, const std::optional& memory_config, const std::optional& compute_kernel_config) { - ttnn::SmallVector dims = tt::operations::primary::get_dim(dim, input.get_legacy_shape().rank()); + ttnn::SmallVector dims = get_dim(dim, input.get_legacy_shape().rank()); std::sort(dims.begin(), dims.end()); return { operation_attributes_t{ diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_norm_backward/device/moreh_norm_backward_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_norm_backward/device/moreh_norm_backward_program_factory.cpp index 79413615e68..d12fcb9e803 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_norm_backward/device/moreh_norm_backward_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_norm_backward/device/moreh_norm_backward_program_factory.cpp @@ -135,7 +135,7 @@ MorehNormBackwardOperation::ProgramFactory::cached_program_t MorehNormBackwardOp const uint32_t im6_t{1}; const uint32_t im7_t{1}; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -170,16 +170,16 @@ MorehNormBackwardOperation::ProgramFactory::cached_program_t MorehNormBackwardOp "writer_moreh_norm_backward.cpp"; std::vector reader_compile_time_args = { - static_cast(tt::operations::primary::is_dram(input)), - static_cast(tt::operations::primary::is_dram(output)), - static_cast(tt::operations::primary::is_dram(output_grad)), + static_cast(is_dram(input)), + static_cast(is_dram(output)), + static_cast(is_dram(output_grad)), static_cast(input_grad_rank)}; std::vector writer_compile_time_args = { - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; const auto reader_kernels_id = - tt::operations::primary::CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); + CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); const auto writer_kernels_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -193,7 +193,7 @@ MorehNormBackwardOperation::ProgramFactory::cached_program_t MorehNormBackwardOp } const std::vector compute_args_group_1{num_cols_per_core_group_1, need_bcast_dim[0], need_bcast_dim[1]}; - const auto compute_kernels_id_1 = tt::operations::primary::CreateComputeKernel( + const auto compute_kernels_id_1 = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, @@ -206,7 +206,7 @@ MorehNormBackwardOperation::ProgramFactory::cached_program_t MorehNormBackwardOp if (!core_group_2.ranges().empty()) { const std::vector compute_args_group_2{ num_cols_per_core_group_2, need_bcast_dim[0], need_bcast_dim[1]}; - compute_kernels_id_2 = tt::operations::primary::CreateComputeKernel( + compute_kernels_id_2 = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/moreh_sgd_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/moreh_sgd_device_operation.cpp index a9645109a22..22d725b3734 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/moreh_sgd_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/moreh_sgd_device_operation.cpp @@ -14,21 +14,21 @@ void MorehSgdOperation::validate_inputs( auto& params_in = tensor_args.param_in; auto& grad = tensor_args.grad; - tt::operations::primary::check_tensor(params_in, "moreh_sgd", "params_in", {DataType::BFLOAT16}); - tt::operations::primary::check_tensor(grad, "moreh_sgd", "grad", {DataType::BFLOAT16}); + check_tensor(params_in, "moreh_sgd", "params_in", {DataType::BFLOAT16}); + check_tensor(grad, "moreh_sgd", "grad", {DataType::BFLOAT16}); if (tensor_args.momentum_buffer_in) { - tt::operations::primary::check_tensor( + check_tensor( *tensor_args.momentum_buffer_in, "moreh_sgd", "momentum_buffer_in", {DataType::BFLOAT16}); } if (tensor_args.param_out.has_value()) { - tt::operations::primary::check_tensor( + check_tensor( tensor_args.param_out.value(), "moreh_sgd", "param_out", {DataType::BFLOAT16}); } if (tensor_args.momentum_buffer_out.has_value()) { - tt::operations::primary::check_tensor( + check_tensor( tensor_args.momentum_buffer_out.value(), "moreh_sgd", "momentum_buffer_out", {DataType::BFLOAT16}); } } diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/moreh_sgd_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/moreh_sgd_program_factory.cpp index e8be177e76f..278366cee7f 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/moreh_sgd_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/moreh_sgd_program_factory.cpp @@ -63,7 +63,7 @@ MorehSgdOperation::ProgramFactory::cached_program_t MorehSgdOperation::ProgramFa //////////////////////////////////////////////////////////////////////////// auto data_format = tt::tt_metal::datatype_to_dataformat_converter(param_in.get_dtype()); auto intermed_cb_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -121,15 +121,15 @@ MorehSgdOperation::ProgramFactory::cached_program_t MorehSgdOperation::ProgramFa //////////////////////////////////////////////////////////////////////////// const std::vector reader_compile_time_args{ - static_cast(tt::operations::primary::is_dram(param_in)), - static_cast(tt::operations::primary::is_dram(grad)), + static_cast(is_dram(param_in)), + static_cast(is_dram(grad)), static_cast( - momentum_buffer_in.has_value() ? tt::operations::primary::is_dram(momentum_buffer_in.value()) : 0)}; + momentum_buffer_in.has_value() ? is_dram(momentum_buffer_in.value()) : 0)}; - std::vector writer_compile_time_args{static_cast(tt::operations::primary::is_dram(param_out))}; + std::vector writer_compile_time_args{static_cast(is_dram(param_out))}; if (has_momentum_buffer_out) writer_compile_time_args.push_back( - static_cast(tt::operations::primary::is_dram(momentum_buffer_out.value()))); + static_cast(is_dram(momentum_buffer_out.value()))); const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/kernels/" @@ -138,9 +138,9 @@ MorehSgdOperation::ProgramFactory::cached_program_t MorehSgdOperation::ProgramFa "ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/kernels/" "writer_moreh_sgd.cpp"; - const auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + const auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); - const auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + const auto writer_kernel_id = CreateWriteKernel( program, writer_kernel_file, all_cores, writer_compile_time_args, writer_defines); //////////////////////////////////////////////////////////////////////////// @@ -152,7 +152,7 @@ MorehSgdOperation::ProgramFactory::cached_program_t MorehSgdOperation::ProgramFa "ttnn/cpp/ttnn/operations/moreh/moreh_sgd/device/kernels/" "moreh_sgd.cpp"; - auto compute_kernel_id = tt ::operations::primary::CreateComputeKernel( + auto compute_kernel_id = CreateComputeKernel( program, compute_kernel_file, { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_c_large/softmax_c_large.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_c_large/softmax_c_large.cpp index 73a87fcf49c..5a170a6b25b 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_c_large/softmax_c_large.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_c_large/softmax_c_large.cpp @@ -34,7 +34,7 @@ MorehSoftmaxOperation::MorehSoftmaxCLargeFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_tiles); + split_work_to_cores_wt_core_range(core_range, num_tiles); auto arch = input.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -46,7 +46,7 @@ MorehSoftmaxOperation::MorehSoftmaxCLargeFactory::create( auto data_format = tt::tt_metal::datatype_to_dataformat_converter(input.get_dtype()); auto intermed_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -67,13 +67,13 @@ MorehSoftmaxOperation::MorehSoftmaxCLargeFactory::create( std::map reader_defines; std::map writer_defines; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/reader_moreh_softmax_c_large.cpp", all_cores, {src_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/writer_moreh_softmax_c_large.cpp", all_cores, @@ -101,7 +101,7 @@ MorehSoftmaxOperation::MorehSoftmaxCLargeFactory::create( } // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/moreh_softmax_c_large.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_h_large/softmax_h_large.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_h_large/softmax_h_large.cpp index 4b38c7c02e3..94d1d26bca0 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_h_large/softmax_h_large.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_h_large/softmax_h_large.cpp @@ -33,7 +33,7 @@ MorehSoftmaxOperation::MorehSoftmaxHLargeFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_cols_tiles); + split_work_to_cores_wt_core_range(core_range, num_cols_tiles); auto arch = input.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -45,7 +45,7 @@ MorehSoftmaxOperation::MorehSoftmaxHLargeFactory::create( auto data_format = tt::tt_metal::datatype_to_dataformat_converter(input.get_dtype()); auto intermed_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -68,13 +68,13 @@ MorehSoftmaxOperation::MorehSoftmaxHLargeFactory::create( std::map reader_defines; std::map writer_defines; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels//reader_moreh_softmax_h_large.cpp", all_cores, {src_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/writer_moreh_softmax_h_large.cpp", all_cores, @@ -96,7 +96,7 @@ MorehSoftmaxOperation::MorehSoftmaxHLargeFactory::create( } // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/moreh_softmax_h_large.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_h_small/softmax_h_small.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_h_small/softmax_h_small.cpp index b182ec8d63c..d2583efadd4 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_h_small/softmax_h_small.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_h_small/softmax_h_small.cpp @@ -34,7 +34,7 @@ MorehSoftmaxOperation::MorehSoftmaxHSmallFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_cols_tiles); + split_work_to_cores_wt_core_range(core_range, num_cols_tiles); auto arch = input.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -46,7 +46,7 @@ MorehSoftmaxOperation::MorehSoftmaxHSmallFactory::create( auto data_format = tt::tt_metal::datatype_to_dataformat_converter(input.get_dtype()); auto intermed_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -69,13 +69,13 @@ MorehSoftmaxOperation::MorehSoftmaxHSmallFactory::create( std::map reader_defines; std::map writer_defines; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/reader_moreh_softmax_h.cpp", all_cores, {src_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/writer_moreh_softmax_h.cpp", all_cores, @@ -97,7 +97,7 @@ MorehSoftmaxOperation::MorehSoftmaxHSmallFactory::create( } // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/moreh_softmax_h.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_w_large/softmax_w_large.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_w_large/softmax_w_large.cpp index 2622708e47f..e01d417fc05 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_w_large/softmax_w_large.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_w_large/softmax_w_large.cpp @@ -34,7 +34,7 @@ MorehSoftmaxOperation::MorehSoftmaxWLargeFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_kernel_rows); + split_work_to_cores_wt_core_range(core_range, num_kernel_rows); auto arch = input.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -46,7 +46,7 @@ MorehSoftmaxOperation::MorehSoftmaxWLargeFactory::create( auto data_format = tt::tt_metal::datatype_to_dataformat_converter(input.get_dtype()); auto intermed_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -69,13 +69,13 @@ MorehSoftmaxOperation::MorehSoftmaxWLargeFactory::create( std::map reader_defines; std::map writer_defines; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/reader_moreh_softmax_w_large.cpp", all_cores, {src_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/writer_moreh_softmax_w_large.cpp", all_cores, @@ -97,7 +97,7 @@ MorehSoftmaxOperation::MorehSoftmaxWLargeFactory::create( } // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/moreh_softmax_w_large.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_w_small/softmax_w_small.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_w_small/softmax_w_small.cpp index a43840e1949..412e1750d80 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_w_small/softmax_w_small.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/softmax_w_small/softmax_w_small.cpp @@ -34,7 +34,7 @@ MorehSoftmaxOperation::MorehSoftmaxWSmallFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_kernel_rows); + split_work_to_cores_wt_core_range(core_range, num_kernel_rows); auto arch = input.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -46,7 +46,7 @@ MorehSoftmaxOperation::MorehSoftmaxWSmallFactory::create( auto data_format = tt::tt_metal::datatype_to_dataformat_converter(input.get_dtype()); auto intermed_data_format = fp32_dest_acc_en ? tt::DataFormat::Float32 : data_format; - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -69,13 +69,13 @@ MorehSoftmaxOperation::MorehSoftmaxWSmallFactory::create( std::map reader_defines; std::map writer_defines; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/reader_moreh_softmax_w.cpp", all_cores, {src_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/writer_moreh_softmax_w.cpp", all_cores, @@ -96,7 +96,7 @@ MorehSoftmaxOperation::MorehSoftmaxWSmallFactory::create( } // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax/device/kernels/moreh_softmax_w.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_c_large/softmax_backward_c_large.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_c_large/softmax_backward_c_large.cpp index 9445e917f7b..3d65e02b5fb 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_c_large/softmax_backward_c_large.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_c_large/softmax_backward_c_large.cpp @@ -35,7 +35,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardCLargeFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_tiles); + split_work_to_cores_wt_core_range(core_range, num_tiles); auto arch = input_grad.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -46,7 +46,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardCLargeFactory::create( // create circular buffers tt::DataFormat data_format = tt::tt_metal::datatype_to_dataformat_converter(input_grad.get_dtype()); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -82,13 +82,13 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardCLargeFactory::create( compute_defines["FP32_DEST_ACC_EN"] = "1"; } - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/reader_moreh_softmax_backward_c.cpp", all_cores, {y_is_dram, dy_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/writer_moreh_softmax_backward_c.cpp", all_cores, @@ -103,7 +103,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardCLargeFactory::create( auto inner_size = outer_stride / dim_size; // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/moreh_softmax_backward_c_large.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_h_large/softmax_backward_h_large.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_h_large/softmax_backward_h_large.cpp index 7d8f06884dc..50932b704f4 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_h_large/softmax_backward_h_large.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_h_large/softmax_backward_h_large.cpp @@ -35,7 +35,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardHLargeFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_cols_tiles); + split_work_to_cores_wt_core_range(core_range, num_cols_tiles); auto arch = input_grad.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -46,7 +46,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardHLargeFactory::create( // create circular buffers tt::DataFormat data_format = tt::tt_metal::datatype_to_dataformat_converter(input_grad.get_dtype()); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -86,14 +86,14 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardHLargeFactory::create( compute_defines["FP32_DEST_ACC_EN"] = "1"; } - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/" "reader_moreh_softmax_backward_h_large.cpp", all_cores, {y_is_dram, dy_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/writer_moreh_softmax_h.cpp", all_cores, @@ -101,7 +101,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardHLargeFactory::create( writer_defines); // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/moreh_softmax_backward_h_large.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_h_small/softmax_backward_h_small.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_h_small/softmax_backward_h_small.cpp index 997d1b56259..33f3e2bc4c9 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_h_small/softmax_backward_h_small.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_h_small/softmax_backward_h_small.cpp @@ -35,7 +35,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardHSmallFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_cols_tiles); + split_work_to_cores_wt_core_range(core_range, num_cols_tiles); auto arch = input_grad.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -46,7 +46,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardHSmallFactory::create( // create circular buffers tt::DataFormat data_format = tt::tt_metal::datatype_to_dataformat_converter(input_grad.get_dtype()); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -71,13 +71,13 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardHSmallFactory::create( std::map reader_defines; std::map writer_defines; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/reader_moreh_softmax_backward_h.cpp", all_cores, {y_is_dram, dy_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/writer_moreh_softmax_h.cpp", all_cores, @@ -99,7 +99,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardHSmallFactory::create( } // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/moreh_softmax_backward_h.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_w_large/softmax_backward_w_large.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_w_large/softmax_backward_w_large.cpp index 8090c3c232f..0594fc6fb7d 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_w_large/softmax_backward_w_large.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_w_large/softmax_backward_w_large.cpp @@ -35,7 +35,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardWLargeFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_kernel_rows); + split_work_to_cores_wt_core_range(core_range, num_kernel_rows); auto arch = input_grad.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -46,7 +46,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardWLargeFactory::create( // create circular buffers tt::DataFormat data_format = tt::tt_metal::datatype_to_dataformat_converter(input_grad.get_dtype()); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -86,14 +86,14 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardWLargeFactory::create( compute_defines["FP32_DEST_ACC_EN"] = "1"; } - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/" "reader_moreh_softmax_backward_w_large.cpp", all_cores, {y_is_dram, dy_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/writer_moreh_softmax_w.cpp", all_cores, @@ -101,7 +101,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardWLargeFactory::create( writer_defines); // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/moreh_softmax_backward_w_large.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_w_small/softmax_backward_w_small.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_w_small/softmax_backward_w_small.cpp index 213741f30de..12558db92f0 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_w_small/softmax_backward_w_small.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/softmax_backward_w_small/softmax_backward_w_small.cpp @@ -35,7 +35,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardWSmallFactory::create( uint32_t core_h = core_range.end_coord.y - core_range.start_coord.y + 1; auto [num_cores, all_cores, core_group_1, core_group_2, num_tiles_per_core_group_1, num_tiles_per_core_group_2] = - tt::operations::primary::split_work_to_cores(core_range, num_kernel_rows); + split_work_to_cores_wt_core_range(core_range, num_kernel_rows); auto arch = input_grad.device()->arch(); auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -46,7 +46,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardWSmallFactory::create( // create circular buffers tt::DataFormat data_format = tt::tt_metal::datatype_to_dataformat_converter(input_grad.get_dtype()); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, data_format, @@ -70,13 +70,13 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardWSmallFactory::create( std::map reader_defines; std::map writer_defines; - auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + auto reader_kernel_id = CreateReadKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/reader_moreh_softmax_backward_w.cpp", all_cores, {y_is_dram, dy_is_dram}, reader_defines); - auto writer_kernel_id = tt::operations::primary::CreateWriteKernel( + auto writer_kernel_id = CreateWriteKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/writer_moreh_softmax_w.cpp", all_cores, @@ -97,7 +97,7 @@ MorehSoftmaxBackwardOperation::MorehSoftmaxBackwardWSmallFactory::create( compute_defines["FP32_DEST_ACC_EN"] = "1"; } // create compute kernel - tt::operations::primary::CreateComputeKernel( + CreateComputeKernel( program, "ttnn/cpp/ttnn/operations/moreh/moreh_softmax_backward/device/kernels/moreh_softmax_backward_w.cpp", { diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_h_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_h_program_factory.cpp index bb2652f4e46..5bf8d308585 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_h_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_h_program_factory.cpp @@ -26,7 +26,7 @@ MorehSumOperation::MorehSumHIntFactory::cached_program_t MorehSumOperation::More const auto cb_data_format{datatype_to_dataformat_converter(output.get_dtype())}; const auto shape{input.get_padded_shape()}; - const auto [W, H, other_dims_product] = tt::operations::primary::extract_spatial_dims(shape); + const auto [W, H, other_dims_product] = extract_spatial_dims(shape); uint32_t Wt{W / tt::constants::TILE_WIDTH}; uint32_t Ht{H / tt::constants::TILE_HEIGHT}; uint32_t HtWt{Ht * Wt}; @@ -80,7 +80,7 @@ MorehSumOperation::MorehSumHIntFactory::cached_program_t MorehSumOperation::More //////////////////////////////////////////////////////////////////////////// // CircularBuffer Setup //////////////////////////////////////////////////////////////////////////// - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -94,20 +94,20 @@ MorehSumOperation::MorehSumHIntFactory::cached_program_t MorehSumOperation::More // DataMovementKernel SetUp //////////////////////////////////////////////////////////////////////////// std::vector reader_compile_time_args = { - static_cast(tt::operations::primary::is_dram(input)), Ht, Wt}; + static_cast(is_dram(input)), Ht, Wt}; std::map reader_defines{}; if (do_mask_h) { reader_defines["DO_MASK_H"] = "1"; } - std::vector writer_compile_time_args = {static_cast(tt::operations::primary::is_dram(output))}; + std::vector writer_compile_time_args = {static_cast(is_dram(output))}; const auto reader_kernel_file{ "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_h_impl_kernels/reader_moreh_int_sum_h.cpp"}; const auto writer_kernel_file{ "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_h_impl_kernels/writer_moreh_int_sum_h.cpp"}; - const auto reader_kernel_id{tt::operations::primary::CreateReadKernel( + const auto reader_kernel_id{CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines)}; const auto writer_kernel_id{ - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args)}; + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args)}; //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -123,7 +123,7 @@ MorehSumOperation::MorehSumHIntFactory::cached_program_t MorehSumOperation::More } const auto compute_kernel_file{ "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_h_impl_kernels/moreh_int_sum_h.cpp"}; - const auto compute_kernel_1_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_1_id = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, @@ -138,7 +138,7 @@ MorehSumOperation::MorehSumHIntFactory::cached_program_t MorehSumOperation::More num_cols_per_core_group_2, // num_cols Ht, // Ht origin_H}; - compute_kernel_2_id = tt::operations::primary::CreateComputeKernel( + compute_kernel_2_id = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_nc_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_nc_program_factory.cpp index 28595685d09..d0813de1b28 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_nc_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_nc_program_factory.cpp @@ -31,7 +31,7 @@ MorehSumOperation::MorehSumNCIntFactory::cached_program_t MorehSumOperation::Mor const auto input_shape = input.get_padded_shape(); const auto input_shape_without_padding = input.get_logical_shape(); const auto [Wt, Ht, inner_tile_size, reduce_tile_size] = - tt::operations::primary::extract_and_scale_spatial_dims(input_shape, static_cast(dim)); + extract_and_scale_spatial_dims(input_shape, static_cast(dim)); const auto num_reduce_input_tile{input_shape[dim]}; const auto num_output_tiles{output.volume() / tt::constants::TILE_HW}; auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -71,7 +71,7 @@ MorehSumOperation::MorehSumNCIntFactory::cached_program_t MorehSumOperation::Mor num_cols_per_core_group_1, num_cols_per_core_group_2] = tt::tt_metal::split_work_to_cores(grid, num_output_tiles); - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -81,16 +81,16 @@ MorehSumOperation::MorehSumNCIntFactory::cached_program_t MorehSumOperation::Mor {tt::CB::c_out0, out0_t}, // output }); - std::vector reader_compile_time_args = {static_cast(tt::operations::primary::is_dram(input))}; - std::vector writer_compile_time_args = {static_cast(tt::operations::primary::is_dram(output))}; + std::vector reader_compile_time_args = {static_cast(is_dram(input))}; + std::vector writer_compile_time_args = {static_cast(is_dram(output))}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_impl_kernels/reader_moreh_sum_nc.cpp"; const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_impl_kernels/writer_moreh_sum_nc.cpp"; const auto reader_kernel_id = - tt::operations::primary::CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); + CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -102,7 +102,7 @@ MorehSumOperation::MorehSumNCIntFactory::cached_program_t MorehSumOperation::Mor } const auto compute_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_impl_kernels/moreh_int_sum_nc.cpp"; - const auto compute_kernel_1_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_1_id = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, @@ -114,7 +114,7 @@ MorehSumOperation::MorehSumNCIntFactory::cached_program_t MorehSumOperation::Mor std::optional compute_kernel_2_id = std::nullopt; if (!core_group_2.ranges().empty()) { const std::vector compute_args_group_2{num_cols_per_core_group_2, num_reduce_input_tile}; - compute_kernel_2_id = tt::operations::primary::CreateComputeKernel( + compute_kernel_2_id = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_w_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_w_program_factory.cpp index 28108a232f5..0f99c77962b 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_w_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_int_sum_w_program_factory.cpp @@ -29,7 +29,7 @@ MorehSumOperation::MorehSumWIntFactory::cached_program_t MorehSumOperation::More const auto cb_data_format{datatype_to_dataformat_converter(output.get_dtype())}; const auto shape{input.get_padded_shape()}; - const auto [W, H, other_dims_product] = tt::operations::primary::extract_spatial_dims(shape); + const auto [W, H, other_dims_product] = extract_spatial_dims(shape); uint32_t Wt{W / tt::constants::TILE_WIDTH}; uint32_t Ht{H / tt::constants::TILE_HEIGHT}; uint32_t num_tiles = input.volume() / tt::constants::TILE_HW; @@ -82,7 +82,7 @@ MorehSumOperation::MorehSumWIntFactory::cached_program_t MorehSumOperation::More //////////////////////////////////////////////////////////////////////////// // CircularBuffer Setup //////////////////////////////////////////////////////////////////////////// - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -95,20 +95,20 @@ MorehSumOperation::MorehSumWIntFactory::cached_program_t MorehSumOperation::More //////////////////////////////////////////////////////////////////////////// // DataMovementKernel SetUp //////////////////////////////////////////////////////////////////////////// - std::vector reader_compile_time_args = {static_cast(tt::operations::primary::is_dram(input))}; + std::vector reader_compile_time_args = {static_cast(is_dram(input))}; std::map reader_defines{}; if (do_mask_w) { reader_defines["DO_MASK_W"] = "1"; } - std::vector writer_compile_time_args = {static_cast(tt::operations::primary::is_dram(output))}; + std::vector writer_compile_time_args = {static_cast(is_dram(output))}; const auto reader_kernel_file{ "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_w_impl_kernels/reader_moreh_int_sum_w.cpp"}; const auto writer_kernel_file{ "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_w_impl_kernels/writer_moreh_int_sum_w.cpp"}; - const auto reader_kernel_id{tt::operations::primary::CreateReadKernel( + const auto reader_kernel_id{CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines)}; const auto writer_kernel_id{ - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args)}; + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args)}; //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -124,7 +124,7 @@ MorehSumOperation::MorehSumWIntFactory::cached_program_t MorehSumOperation::More } const auto compute_kernel_file{ "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_w_impl_kernels/moreh_int_sum_w.cpp"}; - const auto compute_kernel_1_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_1_id = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_rows_per_core_group_1, compute_args_group_1}, @@ -139,7 +139,7 @@ MorehSumOperation::MorehSumWIntFactory::cached_program_t MorehSumOperation::More num_rows_per_core_group_2, // num_rows Wt, // Wt origin_W}; - compute_kernel_2_id = tt::operations::primary::CreateComputeKernel( + compute_kernel_2_id = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_rows_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_device_operation.cpp index f66e99ca63b..64d704f13b7 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_device_operation.cpp @@ -42,13 +42,13 @@ void validate_tensors( const auto& input = tensor_args.input; auto& output = tensor_args.output; - tt::operations::primary::check_tensor(input, "moreh_sum", "input", {DataType::BFLOAT16, DataType::INT32}); - tt::operations::primary::check_tensor(output, "moreh_sum", "output", {DataType::BFLOAT16, DataType::INT32}); + check_tensor(input, "moreh_sum", "input", {DataType::BFLOAT16, DataType::INT32}); + check_tensor(output, "moreh_sum", "output", {DataType::BFLOAT16, DataType::INT32}); - tt::operations::primary::validate_input_with_dim(input, operation_attributes.dim); + validate_input_with_dim(input, operation_attributes.dim); if (output.has_value()) { - tt::operations::primary::validate_output_with_keepdim( + validate_output_with_keepdim( input, output.value(), operation_attributes.dim, operation_attributes.keepdim); } } diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_h_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_h_program_factory.cpp index 594b27f1ff0..6eadb070199 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_h_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_h_program_factory.cpp @@ -27,7 +27,7 @@ MorehSumOperation::MorehSumHFactory::cached_program_t MorehSumOperation::MorehSu float scaler = 1.0f; const auto shape = input.get_padded_shape(); - const auto [W, H, other_dims_product] = tt::operations::primary::extract_spatial_dims(shape); + const auto [W, H, other_dims_product] = extract_spatial_dims(shape); uint32_t Wt = W / tt::constants::TILE_WIDTH; uint32_t Ht = H / tt::constants::TILE_HEIGHT; @@ -75,7 +75,7 @@ MorehSumOperation::MorehSumHFactory::cached_program_t MorehSumOperation::MorehSu {0, 0}, {compute_with_storage_grid_size.x - 1, compute_with_storage_grid_size.y - 1}); auto [num_cores, all_cores, core_group_1, core_group_2, num_cols_per_core_group_1, num_cols_per_core_group_2] = - tt::operations::primary::split_work_to_cores(all_core_range, num_cols); + split_work_to_cores_wt_core_range(all_core_range, num_cols); string compute_kernel_name = "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_h_impl_kernels/moreh_sum_h.cpp"; diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_program_factory.cpp index ebd8d432981..504eaff5135 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_program_factory.cpp @@ -30,7 +30,7 @@ MorehSumOperation::MorehSumNCFactory::cached_program_t MorehSumOperation::MorehS const auto input_shape = input.get_padded_shape(); const auto input_shape_without_padding = input.get_logical_shape(); const auto [Wt, Ht, inner_tile_size, reduce_tile_size] = - tt::operations::primary::extract_and_scale_spatial_dims(input_shape, static_cast(dim)); + extract_and_scale_spatial_dims(input_shape, static_cast(dim)); const auto num_reduce_input_tile = input_shape[dim]; const auto num_output_tiles = output.volume() / tt::constants::TILE_HW; auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = @@ -69,7 +69,7 @@ MorehSumOperation::MorehSumNCFactory::cached_program_t MorehSumOperation::MorehS //////////////////////////////////////////////////////////////////////////// // CircularBuffer Setup //////////////////////////////////////////////////////////////////////////// - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -82,18 +82,18 @@ MorehSumOperation::MorehSumNCFactory::cached_program_t MorehSumOperation::MorehS //////////////////////////////////////////////////////////////////////////// // DataMovementKernel SetUp //////////////////////////////////////////////////////////////////////////// - std::vector reader_compile_time_args = {static_cast(tt::operations::primary::is_dram(input))}; + std::vector reader_compile_time_args = {static_cast(is_dram(input))}; std::map reader_defines; reader_defines["USE_FPU"] = "1"; - std::vector writer_compile_time_args = {static_cast(tt::operations::primary::is_dram(output))}; + std::vector writer_compile_time_args = {static_cast(is_dram(output))}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_impl_kernels/reader_moreh_sum_nc.cpp"; const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_impl_kernels/writer_moreh_sum_nc.cpp"; - const auto reader_kernel_id = tt::operations::primary::CreateReadKernel( + const auto reader_kernel_id = CreateReadKernel( program, reader_kernel_file, all_cores, reader_compile_time_args, reader_defines); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -111,7 +111,7 @@ MorehSumOperation::MorehSumNCFactory::cached_program_t MorehSumOperation::MorehS compute_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_nc_impl_kernels/moreh_sum_nc_gs.cpp"; } - const auto compute_kernel_1_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_1_id = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, @@ -124,7 +124,7 @@ MorehSumOperation::MorehSumNCFactory::cached_program_t MorehSumOperation::MorehS std::optional compute_kernel_2_id = std::nullopt; if (!core_group_2.ranges().empty()) { const std::vector compute_args_group_2{num_cols_per_core_group_2, num_reduce_input_tile}; - compute_kernel_2_id = tt::operations::primary::CreateComputeKernel( + compute_kernel_2_id = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_w_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_w_program_factory.cpp index 7c8eba24a72..0b5164b5a45 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_w_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_w_program_factory.cpp @@ -26,7 +26,7 @@ MorehSumOperation::MorehSumWFactory::cached_program_t MorehSumOperation::MorehSu float scaler = 1.0f; const auto shape = input.get_padded_shape(); - const auto [W, H, other_dims_product] = tt::operations::primary::extract_spatial_dims(shape); + const auto [W, H, other_dims_product] = extract_spatial_dims(shape); uint32_t HW = H * W; uint32_t Wt = W / tt::constants::TILE_WIDTH; @@ -75,7 +75,7 @@ MorehSumOperation::MorehSumWFactory::cached_program_t MorehSumOperation::MorehSu {0, 0}, {compute_with_storage_grid_size.x - 1, compute_with_storage_grid_size.y - 1}); auto [num_cores, all_cores, core_group_1, core_group_2, num_rows_per_core_group_1, num_rows_per_core_group_2] = - tt::operations::primary::split_work_to_cores(all_core_range, num_rows); + split_work_to_cores_wt_core_range(all_core_range, num_rows); string compute_kernel_name = "ttnn/cpp/ttnn/operations/moreh/moreh_sum/device/moreh_sum_w_impl_kernels/moreh_sum_w.cpp"; diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/moreh_sum.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/moreh_sum.cpp index d5ce16b7cb4..d1d38f7df9e 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum/moreh_sum.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum/moreh_sum.cpp @@ -15,7 +15,7 @@ Tensor MorehSum::invoke( const std::optional& output, const std::optional& memory_config, const std::optional& compute_kernel_config) { - ttnn::SmallVector dims = tt::operations::primary::get_dim(dim, input.get_legacy_shape().rank()); + ttnn::SmallVector dims = get_dim(dim, input.get_legacy_shape().rank()); std::sort(dims.begin(), dims.end()); auto temp_input = input; diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/moreh_sum_backward_device_operation.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/moreh_sum_backward_device_operation.cpp index a07da8910c3..8750a7596cc 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/moreh_sum_backward_device_operation.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/moreh_sum_backward_device_operation.cpp @@ -18,14 +18,14 @@ void MorehSumBackwardOperation::validate_inputs( const auto dims = operation_attributes.dims; // validate tensor - tt::operations::primary::check_tensor(output_grad, "moreh_sum_backward", "output_grad"); - tt::operations::primary::check_tensor(input_grad, "moreh_sum_backward", " input_grad"); + check_tensor(output_grad, "moreh_sum_backward", "output_grad"); + check_tensor(input_grad, "moreh_sum_backward", " input_grad"); if (!input.has_value()) { return; } - tt::operations::primary::check_tensor(input, "moreh_sum_backward", "input"); + check_tensor(input, "moreh_sum_backward", "input"); const auto& input_shape = input.value().get_legacy_shape(); auto input_shape_wo_padding = input_shape.without_padding(); auto input_rank = input_shape.rank(); diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/moreh_sum_backward_program_factory.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/moreh_sum_backward_program_factory.cpp index 670bfa0301a..0aecc1495ae 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/moreh_sum_backward_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/moreh_sum_backward_program_factory.cpp @@ -129,7 +129,7 @@ MorehSumBackwardOperation::ProgramFactory::cached_program_t MorehSumBackwardOper //////////////////////////////////////////////////////////////////////////// // CircularBuffer Setup //////////////////////////////////////////////////////////////////////////// - tt::operations::primary::CreateCircularBuffer( + CreateCircularBuffer( program, all_cores, cb_data_format, @@ -143,17 +143,17 @@ MorehSumBackwardOperation::ProgramFactory::cached_program_t MorehSumBackwardOper // DataMovementKernel SetUp //////////////////////////////////////////////////////////////////////////// std::vector reader_compile_time_args = { - static_cast(tt::operations::primary::is_dram(output_grad)), input_grad_rank}; + static_cast(is_dram(output_grad)), input_grad_rank}; std::vector writer_compile_time_args = { - static_cast(tt::operations::primary::is_dram(input_grad))}; + static_cast(is_dram(input_grad))}; const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/kernels/reader_moreh_sum_backward.cpp"; const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/kernels/writer_moreh_sum_backward.cpp"; const auto reader_kernel_id = - tt::operations::primary::CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); + CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); const auto writer_kernel_id = - tt::operations::primary::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -165,7 +165,7 @@ MorehSumBackwardOperation::ProgramFactory::cached_program_t MorehSumBackwardOper } const auto compute_kernel_file = "ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/device/kernels/moreh_sum_backward.cpp"; - const auto compute_kernel_1_id = tt::operations::primary::CreateComputeKernel( + const auto compute_kernel_1_id = CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, @@ -178,7 +178,7 @@ MorehSumBackwardOperation::ProgramFactory::cached_program_t MorehSumBackwardOper if (!core_group_2.ranges().empty()) { const std::vector compute_args_group_2{ num_cols_per_core_group_2, need_bcast_dim[0], need_bcast_dim[1]}; - compute_kernel_2_id = tt::operations::primary::CreateComputeKernel( + compute_kernel_2_id = CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, diff --git a/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/moreh_sum_backward.cpp b/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/moreh_sum_backward.cpp index 1ea51a27cc5..6f32ad0c3be 100644 --- a/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/moreh_sum_backward.cpp +++ b/ttnn/cpp/ttnn/operations/moreh/moreh_sum_backward/moreh_sum_backward.cpp @@ -18,7 +18,7 @@ Tensor MorehSumBackward::invoke( const std::optional& compute_kernel_config) { TT_FATAL((input.has_value() || input_grad.has_value()), "either input or input_grad must have a value"); uint32_t rank = input.has_value() ? input->get_shape().value.rank() : input_grad->get_shape().value.rank(); - ttnn::SmallVector dims = tt::operations::primary::get_dim(dim, rank); + ttnn::SmallVector dims = get_dim(dim, rank); std::sort(dims.begin(), dims.end()); return ttnn::prim::moreh_sum_backward( output_grad, input, dims, keepdim, input_grad, memory_config, compute_kernel_config); diff --git a/ttnn/cpp/ttnn/operations/reduction/prod/device/prod_nc_program_factory.cpp b/ttnn/cpp/ttnn/operations/reduction/prod/device/prod_nc_program_factory.cpp index 6308a96bfe0..c15490dd013 100644 --- a/ttnn/cpp/ttnn/operations/reduction/prod/device/prod_nc_program_factory.cpp +++ b/ttnn/cpp/ttnn/operations/reduction/prod/device/prod_nc_program_factory.cpp @@ -73,7 +73,7 @@ operation::ProgramWithCallbacks prod_nc_format(const Tensor &input, const Tensor //////////////////////////////////////////////////////////////////////////// // CircularBuffer Setup //////////////////////////////////////////////////////////////////////////// - CreateCircularBuffer( + ttnn::operations::CreateCircularBuffer( program, all_cores, cb_data_format, @@ -99,8 +99,8 @@ operation::ProgramWithCallbacks prod_nc_format(const Tensor &input, const Tensor const auto reader_kernel_file = "ttnn/cpp/ttnn/operations/reduction/prod/device/kernels/dataflow/reader_prod_nc.cpp"; const auto writer_kernel_file = "ttnn/cpp/ttnn/operations/eltwise/unary/device/kernels/dataflow/writer_unary_interleaved_start_id.cpp"; - const auto reader_kernel_id = CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); - const auto writer_kernel_id = CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); + const auto reader_kernel_id = ttnn::operations::CreateReadKernel(program, reader_kernel_file, all_cores, reader_compile_time_args); + const auto writer_kernel_id = ttnn::operations::CreateWriteKernel(program, writer_kernel_file, all_cores, writer_compile_time_args); //////////////////////////////////////////////////////////////////////////// // ComputeKernel SetUp @@ -109,13 +109,13 @@ operation::ProgramWithCallbacks prod_nc_format(const Tensor &input, const Tensor std::map compute_defines; const auto compute_kernel_file = "ttnn/cpp/ttnn/operations/reduction/prod/device/kernels/compute/prod_nc.cpp"; - const auto compute_kernel_1_id = CreateComputeKernel( + const auto compute_kernel_1_id = ttnn::operations::CreateComputeKernel( program, compute_kernel_file, {core_group_1, num_cols_per_core_group_1, compute_args_group_1}, compute_defines); std::optional compute_kernel_2_id = std::nullopt; if (!core_group_2.ranges().empty()) { const std::vector compute_args_group_2{num_cols_per_core_group_2}; - compute_kernel_2_id = CreateComputeKernel( + compute_kernel_2_id = ttnn::operations::CreateComputeKernel( program, compute_kernel_file, {core_group_2, num_cols_per_core_group_2, compute_args_group_2}, @@ -146,7 +146,7 @@ operation::ProgramWithCallbacks prod_nc_format(const Tensor &input, const Tensor num_tiles_per_core, input_tile_offset, tile_offset, - static_cast(is_dram(input)), + static_cast(ttnn::operations::is_dram(input)), HtWt, CHtWt, static_cast(dim) @@ -156,7 +156,7 @@ operation::ProgramWithCallbacks prod_nc_format(const Tensor &input, const Tensor program, writer_kernel_id, core, - {output.buffer()->address(), num_tiles_per_core, tile_offset, static_cast(is_dram(output))}); + {output.buffer()->address(), num_tiles_per_core, tile_offset, static_cast(ttnn::operations::is_dram(output))}); if (core_group_1.core_coord_in_core_ranges(core)) { SetRuntimeArgs(program, compute_kernel_1_id, core, {num_reduce_input_tile, num_tiles_per_core});