[rocm-libraries] ROCm/rocm-libraries#7256 (commit 1fc20eb)
Skip numeric drop-out when PComputeWindow is a null_tile_window in Bl… (#7256) The BlockDropout implementation already provides very complete logic for generating random numbers and executing dropout for the P tensor after first attention Gemm with capability to support both Warp-Gemm 32x32 and 16x16 as well as to run on both wave32 and wave64 arch. But in some situation, we only need the block-layer process to generate random numbers, rather than simultaneously execute dropout in real-time on the vgpr tile. For example, xformers' `test_mem_eff_attention.py::test_dropout_ck` requires the host reference implementation of `attention forward with dropout` to use the same random numbers to compare & verify the device side implementation of `attention forward with dropout`, so a standalone kernel to generate random numbers only is required. This PR will enable xformers's random_val generating kernel (in file `ck_tiled_rand_uniform_kernel.h`) to depend on BlockDropout's `Run()` operator completely to generate random numbers for a `[MPerBlock, NPerBlock]` tile during the tile iteration, no need to replicate the logic of BlockDropout in the xformers kernel
Q
Qianfeng committed
0c63d6e776c227b63a45ab8c63daf68dbe2cb319
Parent: 5003f7e
Committed by GitHub <noreply@github.com>
on 5/13/2026, 9:41:25 AM