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.
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
🐛 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 https://github.com/ROCm/xformers/pull/1#issuecomment-2053610724 but different architecture. Related to https://github.com/facebookresearch/xformers/issues/978#issuecomment-1965460368. Hope this helps if support is planned.
To Reproduce
Steps to reproduce the behavior:
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:
gfx1030 errors:
the other gfx1030 error is too long to paste here, but it ends with:
Environment