Hi, I tried installing flash attention on a 7900 XTX using the recently implemented RDNA3 ck support, but building fails at the end. Below is the bottom of the log, shortened since it is over 10000 lines long:
[2764/2764] /opt/rocm/bin/hipcc -I/mnt/NM620/flash-attention/csrc/composable_kernel/include -I/mnt/NM620/flash-attention/csrc/composable_kernel/library/include -I/mnt/NM620/flash-attention/csrc/composable_kernel/example/ck_tile/01_fmha -I/home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include -I/home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/THH -I/opt/rocm/include -I/home/alex/.pyenv/versions/3.13.13/include/python3.13 -c -c /mnt/NM620/flash-attention/csrc/flash_attn_ck/mha_varlen_fwd.hip -o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_varlen_fwd.o -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -fPIC -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -DHIP_ENABLE_WARP_SYNC_BUILTINS=1 --offload-arch=gfx1100 -O3 -std=c++20 -Wno-unknown-warning-option -fbracket-depth=1024 -DCK_TILE_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -DCK_ENABLE_BF16 -DCK_ENABLE_BF8 -DCK_ENABLE_FP16 -DCK_ENABLE_FP32 -DCK_ENABLE_FP64 -DCK_ENABLE_FP8 -DCK_ENABLE_INT8 -DCK_USE_XDL -DUSE_PROF_API=1 -D__HIP_PLATFORM_HCC__=1 -DCK_TILE_FLOAT_TO_BFLOAT16_DEFAULT=0 -mllvm --lsr-drop-solution=1 -fno-offload-uniform-block -mllvm -enable-post-misched=0 -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false -DHIPIFY_V2 -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=flash_attn_2_cuda -fno-gpu-rdc
In file included from /mnt/NM620/flash-attention/csrc/flash_attn_ck/mha_varlen_fwd.hip:7:
In file included from /mnt/NM620/flash-attention/csrc/flash_attn_ck/mha_fwd_head_grouping_utils_hip.hpp:8:
In file included from /mnt/NM620/flash-attention/csrc/flash_attn_ck/flash_common_hip.hpp:9:
In file included from /home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/torch/csrc/api/include/torch/python.h:4:
In file included from /home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/torch/csrc/api/include/torch/detail/static.h:4:
In file included from /home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/torch/csrc/utils/variadic.h:4:
In file included from /home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/ATen/core/Tensor.h:4:
In file included from /home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/ATen/core/TensorBody.h:17:
In file included from /home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/c10/core/Scalar.h:10:
In file included from /home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/c10/core/ScalarType.h:14:
In file included from /home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/c10/util/complex.h:8:
In file included from /home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/include/torch/headeronly/util/complex.h:9:
In file included from /opt/rocm/include/thrust/complex.h:25:
In file included from /opt/rocm/include/thrust/detail/config.h:22:
In file included from /opt/rocm/include/thrust/detail/config/config.h:23:
In file included from /opt/rocm/include/thrust/detail/config/libcxx.h:44:
In file included from /opt/rocm/include/cuda/std/version:31:
In file included from /opt/rocm/include/cuda/std/detail/__config:67:
/opt/rocm/include/cuda/std/detail/libcxx/include/__config:371:12: warning: Assuming 100 MHz realtime clock rate (TSC) for gfx1100/gfx1101 (according to the RDNA3 ISA). Timing-related APIs (e.g., chrono) or sleep instructions may behave incorrectly! [-W#warnings]
371 | # warning Assuming 100 MHz realtime clock rate (TSC) for gfx1100/gfx1101 (according to the RDNA3 ISA). Timing-related APIs (e.g., chrono) or sleep instructions may behave incorrectly!
| ^
1 warning generated when compiling for gfx1100.
In file included from /mnt/NM620/flash-attention/csrc/flash_attn_ck/mha_varlen_fwd.hip:7:
In file included from /mnt/NM620/flash-attention/csrc/flash_attn_ck/mha_fwd_head_grouping_utils_hip.hpp:8:
/mnt/NM620/flash-attention/csrc/flash_attn_ck/flash_common_hip.hpp:31:1: warning: ignored 'inline' attribute on kernel function 'ParsePhiloxCudaState' [-Wcuda-compat]
31 | inline __global__ void ParsePhiloxCudaState(at::PhiloxCudaState arg, uint64_t* rng_state)
| ^
1 warning generated when compiling for host.
g++ -fno-strict-overflow -Wsign-compare -DNDEBUG -g -O3 -Wall -shared -L/home/alex/.pyenv/versions/3.13.13/lib -Wl,-rpath,/home/alex/.pyenv/versions/3.13.13/lib -L/home/alex/.pyenv/versions/3.13.13/lib -Wl,-rpath,/home/alex/.pyenv/versions/3.13.13/lib /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_api.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x0_batch_o2_npad_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x0_batch_o2_pd_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x0_batch_o2_ps_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x0_batch_o2_psd_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x0_group_o2_ps_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x0_group_o2_psd_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x64_batch_o2_npad_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x64_batch_o2_pd_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x64_batch_o2_ps_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x64_batch_o2_psd_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x64_group_o2_ps_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_bf16_b64x64_group_o2_psd_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x0_batch_o2_npad_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x0_batch_o2_pd_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x0_batch_o2_ps_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x0_batch_o2_psd_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x0_group_o2_ps_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x0_group_o2_psd_ndeterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x64_batch_o2_npad_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x64_batch_o2_pd_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x64_batch_o2_ps_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x64_batch_o2_psd_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x64_group_o2_ps_deterministic_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_convert_dq_d128_fp16_b64x64_group_o2_psd_deterministic_gfx11.o /mnt/NM620/
... (continues like this for thousands of lines) ...
.linux-x86_64-cpython-313/build/fmha_fwd_splitkv_d64_fp16_group_b64x64x32x64x32x64_r4x1x1_r4x1x1_w16x16x16_w16x16x16_qr_vr_psskddv_nlogits_nbias_mask_lse_nsquant_npagedkv_nsink_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_fwd_splitkv_d64_fp16_group_b64x64x32x64x32x64_r4x1x1_r4x1x1_w16x16x16_w16x16x16_qr_vr_psskddv_nlogits_nbias_mask_lse_nsquant_pagedkv_nsink_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_fwd_splitkv_d64_fp16_group_b64x64x32x64x32x64_r4x1x1_r4x1x1_w16x16x16_w16x16x16_qr_vr_psskddv_nlogits_nbias_nmask_lse_nsquant_npagedkv_nsink_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_fwd_splitkv_d64_fp16_group_b64x64x32x64x32x64_r4x1x1_r4x1x1_w16x16x16_w16x16x16_qr_vr_psskddv_nlogits_nbias_nmask_lse_nsquant_pagedkv_nsink_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/flash_api.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/flash_common.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_bwd.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_fwd.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_fwd_kvcache.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_varlen_bwd.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_varlen_fwd.o -L/home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/lib -L/opt/rocm/lib -L/opt/rocm/hip/lib -L/home/alex/.pyenv/versions/3.13.13/lib -lc10 -ltorch -ltorch_cpu -ltorch_python -lamdhip64 -lc10_hip -ltorch_hip -o build/lib.linux-x86_64-cpython-313/flash_attn_2_cuda.cpython-313-x86_64-linux-gnu.so
/usr/bin/ld: cannot find /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_d256_bf16_batch_b16x64x256x16x256x16x32x256x256_r1x4x1_r4x1x1_r1x4x1_w16x16x16_w16x16x16_o-1_maxq0_pd1dv1_alibi_ndbias_mask_dropout_wg16_ndeterministic_ntrload_gfx11.o: No such file or directory
/usr/bin/ld: cannot find /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_d256_bf16_group_b16x64x256x16x256x16x32x256x256_r1x4x1_r4x1x1_r1x4x1_w16x16x16_w16x16x16_o-1_maxq0_pd1dv1_alibi_ndbias_mask_dropout_wg16_deterministic_ntrload_gfx11.o: No such file or directory
/usr/bin/ld: cannot find /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_d256_fp16_batch_b16x64x256x16x256x16x32x256x256_r1x4x1_r4x1x1_r1x4x1_w16x16x16_w16x16x16_o-1_maxq0_pd1dv1_alibi_ndbias_mask_ndropout_ndeterministic_ntrload_gfx11.o: No such file or directory
collect2: error: ld returned 1 exit status
error: command '/usr/bin/g++' failed with exit code 1
[end of output]
note: This error originates from a subprocess, and is likely not a problem with pip.
ERROR: Failed building wheel for flash_attn
Failed to build flash_attn
error: failed-wheel-build-for-install
× Failed to build installable wheels for some pyproject.toml based projects
╰─> flash_attn
Hi, I tried installing flash attention on a 7900 XTX using the recently implemented RDNA3 ck support, but building fails at the end. Below is the bottom of the log, shortened since it is over 10000 lines long:
... (continues like this for thousands of lines) ...
.linux-x86_64-cpython-313/build/fmha_fwd_splitkv_d64_fp16_group_b64x64x32x64x32x64_r4x1x1_r4x1x1_w16x16x16_w16x16x16_qr_vr_psskddv_nlogits_nbias_mask_lse_nsquant_npagedkv_nsink_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_fwd_splitkv_d64_fp16_group_b64x64x32x64x32x64_r4x1x1_r4x1x1_w16x16x16_w16x16x16_qr_vr_psskddv_nlogits_nbias_mask_lse_nsquant_pagedkv_nsink_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_fwd_splitkv_d64_fp16_group_b64x64x32x64x32x64_r4x1x1_r4x1x1_w16x16x16_w16x16x16_qr_vr_psskddv_nlogits_nbias_nmask_lse_nsquant_npagedkv_nsink_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_fwd_splitkv_d64_fp16_group_b64x64x32x64x32x64_r4x1x1_r4x1x1_w16x16x16_w16x16x16_qr_vr_psskddv_nlogits_nbias_nmask_lse_nsquant_pagedkv_nsink_gfx11.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/flash_api.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/flash_common.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_bwd.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_fwd.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_fwd_kvcache.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_varlen_bwd.o /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/csrc/flash_attn_ck/mha_varlen_fwd.o -L/home/alex/.pyenv/versions/3.13.13/lib/python3.13/site-packages/torch/lib -L/opt/rocm/lib -L/opt/rocm/hip/lib -L/home/alex/.pyenv/versions/3.13.13/lib -lc10 -ltorch -ltorch_cpu -ltorch_python -lamdhip64 -lc10_hip -ltorch_hip -o build/lib.linux-x86_64-cpython-313/flash_attn_2_cuda.cpython-313-x86_64-linux-gnu.so
/usr/bin/ld: cannot find /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_d256_bf16_batch_b16x64x256x16x256x16x32x256x256_r1x4x1_r4x1x1_r1x4x1_w16x16x16_w16x16x16_o-1_maxq0_pd1dv1_alibi_ndbias_mask_dropout_wg16_ndeterministic_ntrload_gfx11.o: No such file or directory
/usr/bin/ld: cannot find /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_d256_bf16_group_b16x64x256x16x256x16x32x256x256_r1x4x1_r4x1x1_r1x4x1_w16x16x16_w16x16x16_o-1_maxq0_pd1dv1_alibi_ndbias_mask_dropout_wg16_deterministic_ntrload_gfx11.o: No such file or directory
/usr/bin/ld: cannot find /mnt/NM620/flash-attention/build/temp.linux-x86_64-cpython-313/build/fmha_bwd_d256_fp16_batch_b16x64x256x16x256x16x32x256x256_r1x4x1_r4x1x1_r1x4x1_w16x16x16_w16x16x16_o-1_maxq0_pd1dv1_alibi_ndbias_mask_ndropout_ndeterministic_ntrload_gfx11.o: No such file or directory
collect2: error: ld returned 1 exit status
error: command '/usr/bin/g++' failed with exit code 1
[end of output]
note: This error originates from a subprocess, and is likely not a problem with pip.
ERROR: Failed building wheel for flash_attn
Failed to build flash_attn
error: failed-wheel-build-for-install
× Failed to build installable wheels for some pyproject.toml based projects
╰─> flash_attn