[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