[llvm] [PromoteMemToReg] Insert store undef when removing lifetime markers (PR #191909)
Marcin Radomski via llvm-commits
llvm-commits at lists.llvm.org
Mon Sep 7 04:00:23 PDT 2026
dextero wrote:
This appears to cause >10x build time regressions in some heavily templated CUDA kernels. Bisect pointed to this particular PR as the culprit.
Tested with CUDA 12.8.
LLVM commit | Wall time
-|-
fdb39b1112a6db0343de169b79820768bd16c7d9 (without thisPR) | 26.72s
cbc5a226cbf8d1d37d1ba8e55ce6973e8ef739cd (including this PR) | 275.24s
<details>
<summary>Minimized repro code</summary>
Requires https://github.com/Dao-AILab/flash-attention
```cpp
#include <flash_bwd_launch_template.h>
template __global__ void flash::flash_bwd_dq_dk_dv_loop_seqk_parallel_kernel<
Flash_bwd_kernel_traits<
/*kBlockM=*/256,
/*kBlockN=*/64,
/*kHeadDim=*/64,
/*kNWarps=*/8,
/*AtomLayoutMSdP=*/4,
/*AtomLayoutNdKV=*/2,
/*AtomLayoutMdQ=*/2,
/*Is_Q_in_regs=*/false,
/*Share_Q_K_smem=*/true,
cutlass::bfloat16_t,
Flash_kernel_traits<256, 64, 64, 8, cutlass::bfloat16_t>
>,
/*Is_dropout=*/false,
/*Is_causal=*/true,
/*Is_local=*/false,
/*Has_alibi=*/false,
/*Is_even_MN=*/false,
/*Is_even_K=*/false,
/*Return_softmax=*/true
>(flash::Flash_bwd_params);
```
</details>
<details>
<summary>Repro build command</summary>
```
# fill those in as appropriate
CLANG=clang++
CUDA_PATH=/path/to/cuda/dir
FLASH_ATTENTION_DIR=~/tools/flash-attention
time $CLANG -x cuda \
--cuda-path=$CUDA_PATH \
--cuda-gpu-arch=sm_100a \
--cuda-device-only \
-O3 \
-std=c++20 \
-DFLASH_KERNEL_IMPL \
-D_CCCL_NO_SYSTEM_HEADER \
-D_CCCL_DISABLE_FLOAT128 \
-DCUTLASS_ENABLE_TENSOR_CORE_MMA=1 \
-DCUTE_USE_PACKED_TUPLE=1 \
-D__CUDA_NO_HALF_OPERATORS__ \
-D__CUDA_NO_HALF_CONVERSIONS__ \
-D__CUDA_NO_BFLOAT16_CONVERSIONS__ \
-D__CUDA_NO_HALF2_OPERATORS__ \
-I $FLASH_ATTENTION_DIR/stubs \
-I $FLASH_ATTENTION_DIR/csrc/flash_attn/src \
-I $FLASH_ATTENTION_DIR/csrc/cutlass/include \
-c test.cu \
-o test.o
```
</details>
https://github.com/llvm/llvm-project/pull/191909
More information about the llvm-commits
mailing list