Fast tilize in tiled out case for pool ops (tenstorrent#28382)

Use `fast_tilize_*` instead od `tilize_*` API for tilizing pool's output
(better perf). Currently it does not support bfp4_b output so it is
handled with classic tilize, opened issue tenstorrent#28380.

### Checklist
- [x] [All post
commit](https://github.com/tenstorrent/tt-metal/actions/runs/17725852995)
- [x] [Nightly tt-metal L2
tests](https://github.com/tenstorrent/tt-metal/actions/runs/17725837243)
(vae sdxl already regressed, not related to this change)

Author: skrsticTT

ttnn/cpp/ttnn/operations/pool/generic/device/kernels/compute/compute_pool_2d.cpp

@@ -55,7 +55,7 @@ void MAIN {
     constexpr uint32_t pre_tilize_cb_id = get_compile_time_arg_val(19);
     constexpr bool is_output_tiled = get_compile_time_arg_val(20);  // 1 = TILED, 0 = ROW_MAJOR
     constexpr bool is_output_block_format = (bool)get_compile_time_arg_val(21);
-
+    constexpr bool is_output_bfp4_b = (bool)get_compile_time_arg_val(22);
 
     constexpr uint32_t topk_output_tiles = 1;
     constexpr uint32_t topk_cb_tile_idx = 0;
@@ -228,11 +228,17 @@ void MAIN {
                         unary_op_init_common(pre_tilize_cb_id, out_cb_id);
                         tensix_sync();
 
-                        tilize_init(pre_tilize_cb_id, in_ntiles_c, out_cb_id);
-                        tilize_block(pre_tilize_cb_id, in_ntiles_c, out_cb_id);
-
+                        // Skip fast_tilize path for bfp4_b output until #28380 is closed
+                        if constexpr (is_output_bfp4_b) {
+                            tilize_init(pre_tilize_cb_id, in_ntiles_c, out_cb_id);
+                            tilize_block(pre_tilize_cb_id, in_ntiles_c, out_cb_id);
+                            tilize_uninit(pre_tilize_cb_id, out_cb_id);
+                        } else {
+                            fast_tilize_init(pre_tilize_cb_id, in_ntiles_c, out_cb_id);
+                            fast_tilize_block(pre_tilize_cb_id, in_ntiles_c, out_cb_id);
+                            fast_tilize_uninit(pre_tilize_cb_id, out_cb_id);
+                        }
                         cb_push_back(out_cb_id, in_ntiles_c);
-                        tilize_uninit(pre_tilize_cb_id, out_cb_id);
 
                         if constexpr (is_output_block_format) {
                             tensix_sync();

ttnn/cpp/ttnn/operations/pool/generic/device/pool_multi_core_program_factory.cpp

@@ -418,6 +418,7 @@ Pool2D::MultiCore::cached_program_t pool2d_multi_core_sharded_with_halo_v2_impl_
 
     const bool is_output_tiled = output_layout == Layout::TILE;
     const bool is_output_block_format = is_block_float(outputs[0].dtype());
+    const bool is_output_bfp4_b = outputs[0].dtype() == DataType::BFLOAT4_B;
 
     // Conditionally allocate temporary CB - only needed for TILED output
     uint32_t pre_tilize_cb_id = 32;  // default invalid CB ID
@@ -618,7 +619,8 @@ Pool2D::MultiCore::cached_program_t pool2d_multi_core_sharded_with_halo_v2_impl_
         (uint32_t)return_indices,       // 18
         pre_tilize_cb_id,               // 19
         is_output_tiled,                // 20
-        is_output_block_format};        // 21
+        is_output_block_format,         // 21
+        is_output_bfp4_b};              // 22
 
     auto compute_config = tt::tt_metal::ComputeConfig{
         .math_fidelity = MathFidelity::HiFi4,

ttnn/cpp/ttnn/operations/pool/grid_sample/device/grid_sample_program_factory.cpp

@@ -155,6 +155,7 @@ tt::tt_metal::operation::ProgramWithCallbacks grid_sample_program_factory(
 
     const bool is_output_tiled = false;
     const bool is_output_block_format = false;
+    const bool is_output_bfp4_b = false;
     const uint32_t pre_tilize_cb_id =
         32;  // Unused CB for pool compute kernel for grid sample, we don't have tiled output in gridsample
 
@@ -182,7 +183,8 @@ tt::tt_metal::operation::ProgramWithCallbacks grid_sample_program_factory(
             false,                                               // 18: Return Indices (unused)
             pre_tilize_cb_id,                                    // 19: Pre-tilize CB (unused)
             is_output_tiled,                                     // 20: is_output_tiled (unused)
-            is_output_block_format                               // 21: is_output_block_format (unused)
+            is_output_block_format,                              // 21: is_output_block_format (unused)
+            is_output_bfp4_b                                     // 22: is_output_bfp4_b (unused)
         };
 
         compute_kernel_group_1 = tt::tt_metal::CreateKernel(
@@ -221,7 +223,8 @@ tt::tt_metal::operation::ProgramWithCallbacks grid_sample_program_factory(
             false,                                               // 18: Return Indices (unused)
             pre_tilize_cb_id,                                    // 19: Pre-tilize CB (unused)
             is_output_tiled,                                     // 20: is_output_tiled (unused)
-            is_output_block_format                               // 21: is_output_block_format (unused)
+            is_output_block_format,                              // 21: is_output_block_format (unused)
+            is_output_bfp4_b                                     // 22: is_output_bfp4_b (unused)
         };
 
         compute_kernel_group_2 = tt::tt_metal::CreateKernel(