From 725e138553a5552024ac6facf3ea16181e45afd2 Mon Sep 17 00:00:00 2001
From: Qianfeng Zhang <Qianfeng.Zhang@amd.com>
Date: Sat, 9 May 2026 16:06:33 +0000
Subject: [PATCH 1/2] Skip numeric drop-out when PComputeWindow is a
 null_tile_window in BlockDropout Run()

---
 .../ck_tile/ops/fmha/block/block_dropout.hpp  | 38 ++++++++++---------
 1 file changed, 21 insertions(+), 17 deletions(-)

diff --git a/projects/composablekernel/include/ck_tile/ops/fmha/block/block_dropout.hpp b/projects/composablekernel/include/ck_tile/ops/fmha/block/block_dropout.hpp
index 37c1fe480568..fd2af816241f 100644
--- a/projects/composablekernel/include/ck_tile/ops/fmha/block/block_dropout.hpp
+++ b/projects/composablekernel/include/ck_tile/ops/fmha/block/block_dropout.hpp
@@ -381,24 +381,28 @@ struct BlockDropout
                     store_tile(randval_dram_window, randval_store);
                 }
                 move_tile_window(randval_dram_window, {0, kNPerStep});
-                // Drop values of P based on the generated probabilities
-                constexpr auto randval_spans = decltype(randval)::get_distributed_spans();
-                sweep_tile_span(randval_spans[number<0>{}], [&](auto idx0) {
-                    sweep_tile_span(randval_spans[number<1>{}], [&](auto idx1) {
-                        constexpr auto p_idx0 =
-                            tile_distributed_index<i_m0 * MIterPerWarp +
-                                                   idx0.impl_.template at<0>()>{};
-                        constexpr auto p_idx1 =
-                            tile_distributed_index<i_n0,
-                                                   idx1.impl_.template at<1>(),
-                                                   idx1.impl_.template at<2>()>{};
-                        constexpr auto p_idx = ck_tile::make_tuple(p_idx0, p_idx1);
-                        constexpr auto r_idx = ck_tile::make_tuple(idx0, idx1);
-                        p_compute(p_idx)     = randval[r_idx] <= p_undrop_in_uint8_t
-                                                   ? p_compute[p_idx] * rp_undrop
-                                                   : PComputeDataType(0);
+
+                if constexpr(!is_null_tile_window_v<PComputeWindow>)
+                {
+                    // Drop values of P based on the generated probabilities
+                    constexpr auto randval_spans = decltype(randval)::get_distributed_spans();
+                    sweep_tile_span(randval_spans[number<0>{}], [&](auto idx0) {
+                        sweep_tile_span(randval_spans[number<1>{}], [&](auto idx1) {
+                            constexpr auto p_idx0 =
+                                tile_distributed_index<i_m0 * MIterPerWarp +
+                                                       idx0.impl_.template at<0>()>{};
+                            constexpr auto p_idx1 =
+                                tile_distributed_index<i_n0,
+                                                       idx1.impl_.template at<1>(),
+                                                       idx1.impl_.template at<2>()>{};
+                            constexpr auto p_idx = ck_tile::make_tuple(p_idx0, p_idx1);
+                            constexpr auto r_idx = ck_tile::make_tuple(idx0, idx1);
+                            p_compute(p_idx)     = randval[r_idx] <= p_undrop_in_uint8_t
+                                                       ? p_compute[p_idx] * rp_undrop
+                                                       : PComputeDataType(0);
+                        });
                     });
-                });
+                };
             });
             move_tile_window(randval_dram_window, {kMPerStep, -kNPerBlock});
         });

From c2eed01ce3cf2b821bed3170a2c454bfaa78099b Mon Sep 17 00:00:00 2001
From: Qianfeng Zhang <Qianfeng.Zhang@amd.com>
Date: Mon, 11 May 2026 10:24:22 +0000
Subject: [PATCH 2/2] Remove the semicolon following the constexpr() statement
 to align the coding style

---
 .../include/ck_tile/ops/fmha/block/block_dropout.hpp            | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/projects/composablekernel/include/ck_tile/ops/fmha/block/block_dropout.hpp b/projects/composablekernel/include/ck_tile/ops/fmha/block/block_dropout.hpp
index fd2af816241f..78d68a482eb9 100644
--- a/projects/composablekernel/include/ck_tile/ops/fmha/block/block_dropout.hpp
+++ b/projects/composablekernel/include/ck_tile/ops/fmha/block/block_dropout.hpp
@@ -402,7 +402,7 @@ struct BlockDropout
                                                        : PComputeDataType(0);
                         });
                     });
-                };
+                }
             });
             move_tile_window(randval_dram_window, {kMPerStep, -kNPerBlock});
         });