Fails to build for Radeon RX 6950 XT (gfx1030)

[d-ber]

🐛 Bug

Fails to build on a 6950xt with rocm 6.0 on ubuntu 22.04.
Hi, I know only MI series GPU are supported at the moment, I just wanted to report the errors in case they could be useful. Similar to #1 (comment) but different architecture. Related to facebookresearch#978 (comment). Hope this helps if support is planned.

To Reproduce

Steps to reproduce the behavior:

git clone https://github.com/ROCm/xformers
cd xformers
python -m venv venv
source venv/bin/activate
pip3 install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/rocm6.0
git submodule update --init --recursive
pip install .

Full pip install output in two pastebins because of size limit: https://pastebin.com/3Nez5AjE and https://pastebin.com/ETr9496x
It produces errors for both gfx1030 and gfx1036.
gfx1036 error:

In file included from /home/xformers/third_party/composable_kernel_tiled/include/ck/utility/buffer_view_declare.hpp:9:
/home/xformers/third_party/composable_kernel_tiled/include/ck/utility/amd_buffer_addressing.hpp:33:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
  wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
                                             ^
/home/xformers/third_party/composable_kernel_tiled/include/ck/utility/amd_buffer_addressing.hpp:48:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
  wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
                                             ^
2 errors generated when compiling for gfx1036.

gfx1030 errors:

/home/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_attribute_mfma_impl_hip.hpp:116:17: error: '__builtin_amdgcn_mfma_f32_32x32x8bf16_1k' needs target feature mai-insts
      c_vec = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(a_vec, b_vec, c_vec, 0, 0, 0);
              ^
1 error generated when compiling for gfx1030.

the other gfx1030 error is too long to paste here, but it ends with:

In file included from /home/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_bwd_dot_do_o_default_policy_hip.hpp:7:
/home/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_bwd_pipeline_default_policy_hip.hpp:1133:27: error: constexpr variable 'M1' must be initialized by a constant expression
      constexpr index_t M1 = get_warp_size() / (M2 * N0);
                        ^    ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_bwd_pipeline_default_policy_hip.hpp:1133:48: note: division by zero
      constexpr index_t M1 = get_warp_size() / (M2 * N0);
                                             ^
/home/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_bwd_pipeline_default_policy_hip.hpp:1138:63: error: non-type template argument is not a constant expression
                                         Tuple<Sequence<M0, M1, M2, M3>, Sequence<N0, N1>>,
                                                            ^~
/home/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_bwd_pipeline_default_policy_hip.hpp:1138:63: note: initializer of 'M1' is not a constant expression
/home/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_bwd_pipeline_default_policy_hip.hpp:1133:27: note: declared here
      constexpr index_t M1 = get_warp_size() / (M2 * N0);
                        ^
/home/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_bwd_pipeline_default_policy_hip.hpp:1133:48: warning: division by zero is undefined [-Wdivision-by-zero]
      constexpr index_t M1 = get_warp_size() / (M2 * N0);
                                             ^ ~~~~~~~~~
fatal error: too many errors emitted, stopping now [-ferror-limit=]
5 warnings and 20 errors generated when compiling for gfx1030.

Environment

PyTorch version: 2.3.0+rocm6.0
Is debug build: False
CUDA used to build PyTorch: N/A
ROCM used to build PyTorch: 6.0.32830-d62f6a171

OS: Ubuntu 22.04.4 LTS (x86_64)
GCC version: (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0
Clang version: Could not collect
CMake version: version 3.22.1
Libc version: glibc-2.35

Python version: 3.10.12 (main, Nov 20 2023, 15:14:05) [GCC 11.4.0] (64-bit runtime)
Python platform: Linux-6.5.0-35-generic-x86_64-with-glibc2.35
Is CUDA available: True
CUDA runtime version: Could not collect
CUDA_MODULE_LOADING set to: LAZY
GPU models and configuration: AMD Radeon RX 6950 XT (gfx1030)
Nvidia driver version: Could not collect
cuDNN version: Could not collect
HIP runtime version: 6.0.32830
MIOpen runtime version: 3.0.0
Is XNNPACK available: True

CPU:
Architecture:                       x86_64
CPU op-mode(s):                     32-bit, 64-bit
Address sizes:                      48 bits physical, 48 bits virtual
Byte Order:                         Little Endian
CPU(s):                             24
On-line CPU(s) list:                0-23
Vendor ID:                          AuthenticAMD
Model name:                         AMD Ryzen 9 7900X 12-Core Processor
CPU family:                         25
Model:                              97
Thread(s) per core:                 2
Core(s) per socket:                 12
Socket(s):                          1
Stepping:                           2
CPU max MHz:                        5733,0000
CPU min MHz:                        400,0000
BogoMIPS:                           9399.41
Flags:                              fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 ht syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm constant_tsc rep_good amd_lbr_v2 nopl nonstop_tsc cpuid extd_apicid aperfmperf rapl pni pclmulqdq monitor ssse3 fma cx16 sse4_1 sse4_2 x2apic movbe popcnt aes xsave avx f16c rdrand lahf_lm cmp_legacy svm extapic cr8_legacy abm sse4a misalignsse 3dnowprefetch osvw ibs skinit wdt tce topoext perfctr_core perfctr_nb bpext perfctr_llc mwaitx cpb cat_l3 cdp_l3 hw_pstate ssbd mba perfmon_v2 ibrs ibpb stibp ibrs_enhanced vmmcall fsgsbase bmi1 avx2 smep bmi2 erms invpcid cqm rdt_a avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves cqm_llc cqm_occup_llc cqm_mbm_total cqm_mbm_local avx512_bf16 clzero irperf xsaveerptr rdpru wbnoinvd cppc arat npt lbrv svm_lock nrip_save tsc_scale vmcb_clean flushbyasid decodeassists pausefilter pfthreshold avic v_vmsave_vmload vgif x2avic v_spec_ctrl vnmi avx512vbmi umip pku ospke avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq rdpid overflow_recov succor smca fsrm flush_l1d
Virtualization:                     AMD-V
L1d cache:                          384 KiB (12 instances)
L1i cache:                          384 KiB (12 instances)
L2 cache:                           12 MiB (12 instances)
L3 cache:                           64 MiB (2 instances)
NUMA node(s):                       1
NUMA node0 CPU(s):                  0-23
Vulnerability Gather data sampling: Not affected
Vulnerability Itlb multihit:        Not affected
Vulnerability L1tf:                 Not affected
Vulnerability Mds:                  Not affected
Vulnerability Meltdown:             Not affected
Vulnerability Mmio stale data:      Not affected
Vulnerability Retbleed:             Not affected
Vulnerability Spec rstack overflow: Vulnerable: Safe RET, no microcode
Vulnerability Spec store bypass:    Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:           Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:           Mitigation; Enhanced / Automatic IBRS; IBPB conditional; STIBP always-on; RSB filling; PBRSB-eIBRS Not affected; BHI Not affected
Vulnerability Srbds:                Not affected
Vulnerability Tsx async abort:      Not affected

Versions of relevant libraries:
[pip3] numpy==1.26.3
[pip3] pytorch-triton-rocm==2.3.0
[pip3] torch==2.3.0+rocm6.0
[pip3] torchaudio==2.3.0+rocm6.0
[pip3] torchvision==0.18.0+rocm6.0
[conda] Could not collect

[1219jjh]

same error here. Was there any progress?

[1823616178]

Same mistake. Mine is 6600XT

[derfasthirnlosenick]

Still the case. Can provide the errors, or is there no hope of xformers getting to run on gfx1030?