Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We鈥檒l occasionally send you account related emails.

Already on GitHub? Sign in to your account

[ROCm/xformers] error generated when compiling for gfx1100 #1026

Open
Looong01 opened this issue Apr 13, 2024 · 10 comments
Open

[ROCm/xformers] error generated when compiling for gfx1100 #1026

Looong01 opened this issue Apr 13, 2024 · 10 comments

Comments

@Looong01
Copy link

Looong01 commented Apr 13, 2024

馃悰 Bug

I put up some screenshots that I think it is important:
image
image
image
image
image
image
image

And

image
image
image
image
image
image
image

Command

1.	git clone https://github.com/ROCm/xformers
2.	cd xformers
3.	conda create -n xformers python=3.9
4.	conda activate xformers
5.	pip install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/rocm5.7
6.      pip install ninja
7.	pip install ./

Environment

PyTorch version: 2.2.0+rocm5.7
Is debug build: False
CUDA used to build PyTorch: N/A
ROCM used to build PyTorch: 5.7.31921-d1770ee1b

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.9.18 (main, Sep 11 2023, 13:41:44) [GCC 11.2.0] (64-bit runtime)
Python platform: Linux-6.5.0-21-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: Radeon RX 7900 XTX (gfx1100)
Nvidia driver version: Could not collect
cuDNN version: Could not collect
HIP runtime version: 5.7.31921
MIOpen runtime version: 2.20.0
Is XNNPACK available: True

CPU:
Architecture: x86_64
CPU op-mode(s): 32-bit, 64-bit
Address sizes: 39 bits physical, 48 bits virtual
Byte Order: Little Endian
CPU(s): 8
On-line CPU(s) list: 0-7
Vendor ID: GenuineIntel
Model name: Intel(R) Core(TM) i7-9700 CPU @ 3.00GHz
CPU family: 6
Model: 158
Thread(s) per core: 1
Core(s) per socket: 8
Socket(s): 1
Stepping: 13
CPU max MHz: 4700.0000
CPU min MHz: 800.0000
BogoMIPS: 6000.00
Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc art arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc cpuid aperfmperf pni pclmulqdq dtes64 monitor ds_cpl vmx smx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch cpuid_fault epb invpcid_single ssbd ibrs ibpb stibp ibrs_enhanced tpr_shadow flexpriority ept vpid ept_ad fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid mpx rdseed adx smap clflushopt intel_pt xsaveopt xsavec xgetbv1 xsaves dtherm ida arat pln pts hwp hwp_notify hwp_act_window hwp_epp vnmi md_clear flush_l1d arch_capabilities
Virtualization: VT-x
L1d cache: 256 KiB (8 instances)
L1i cache: 256 KiB (8 instances)
L2 cache: 2 MiB (8 instances)
L3 cache: 12 MiB (1 instance)
NUMA node(s): 1
NUMA node0 CPU(s): 0-7
Vulnerability Gather data sampling: Mitigation; Microcode
Vulnerability Itlb multihit: KVM: Mitigation: VMX disabled
Vulnerability L1tf: Not affected
Vulnerability Mds: Not affected
Vulnerability Meltdown: Not affected
Vulnerability Mmio stale data: Mitigation; Clear CPU buffers; SMT disabled
Vulnerability Retbleed: Mitigation; Enhanced IBRS
Vulnerability Spec rstack overflow: Not affected
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, RSB filling, PBRSB-eIBRS SW sequence
Vulnerability Srbds: Mitigation; Microcode
Vulnerability Tsx async abort: Mitigation; TSX disabled

Versions of relevant libraries:
[pip3] numpy==1.24.1
[pip3] pytorch-triton-rocm==2.2.0
[pip3] torch==2.2.0+rocm5.7
[pip3] torchaudio==2.2.0+rocm5.7
[pip3] torchvision==0.17.0+rocm5.7
[conda] numpy 1.24.1 pypi_0 pypi
[conda] pytorch-triton-rocm 2.2.0 pypi_0 pypi
[conda] torch 2.2.0+rocm5.7 pypi_0 pypi
[conda] torchaudio 2.2.0+rocm5.7 pypi_0 pypi
[conda] torchvision 0.17.0+rocm5.7 pypi_0 pypi

Additional context

@tenpercent @qianfengz

@gardner
Copy link

gardner commented Apr 17, 2024

Text is better than screenshots when reporting issues.

It looks like you have ROCm 6.0.1 but you are installing torch built for ROCm 5.7.

What happens when you try with the other index-url?

git clone https://github.com/ROCm/xformers
cd xformers
conda create -n xformers python=3.9
conda activate xformers
- pip install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/rocm5.7
+ pip install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/rocm6.0
pip install ninja
pip install ./

@Looong01
Copy link
Author

Text is better than screenshots when reporting issues.

It looks like you have ROCm 6.0.1 but you are installing torch built for ROCm 5.7.

What happens when you try with the other index-url?

git clone https://github.com/ROCm/xformers
cd xformers
conda create -n xformers python=3.9
conda activate xformers
- pip install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/rocm5.7
+ pip install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/rocm6.0
pip install ninja
pip install ./

No help. The same errors. Thanks anyway!

@gardner
Copy link

gardner commented Apr 17, 2024

It's a bit difficult to read which symbol it's failing on. If you paste the text from the error then it will be searchable for people in the future.

It looks like it might be a similar error to: #978 (comment)

Basically, the ROCm support is experimental. The code might have been merged a bit early / prematurely. It looks like there are some improvements to get it to work. You bug report is an important part of that. Please consider editing your first comment to remove the screenshots and replace it with the text. Thanks 馃檹

@gardner
Copy link

gardner commented Apr 17, 2024

Can you try this Dockerfile that builds the ROCm fork?

FROM rocm/pytorch:latest

WORKDIR /workspace

RUN python3 -m pip install --upgrade pip

RUN apt-get update \
  && apt-get install -y \
    wget \
    git \
    build-essential \
    ninja-build \
    git-lfs \
    libaio-dev \
  && rm -rf /var/lib/apt/lists/*

RUN pip install -U pip \
  && pip install ninja packaging pytest numpy

# See https://github.com/ROCm/xformers/pull/1/files#r1494885190
env PYTORCH_ROCM_ARCH=gfx1100
env MAX_JOBS=12
RUN pip wheel -v --no-build-isolation git+https://github.com/ROCm/xformers.git@main#egg=xformers

RUN find . -name xformers\*.whl

You can build it with:

docker build -t xformers . --progress=plain

@Looong01
Copy link
Author

mha/instances/ck_tiled_fmha_grouped_infer_fp16_no_causalmask_with_attnbias_maxk_32.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_no_causalmask_with_attnbias_maxk_64.cu -> /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_no_causalmask_with_attnbias_maxk_64.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_no_attnbias_maxk_128.cu -> /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_no_attnbias_maxk_128.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_no_attnbias_maxk_256.cu -> /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_no_attnbias_maxk_256.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_no_attnbias_maxk_32.cu -> /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_no_attnbias_maxk_32.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_no_attnbias_maxk_64.cu -> /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_no_attnbias_maxk_64.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_with_attnbias_maxk_128.cu -> /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_with_attnbias_maxk_128.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_with_attnbias_maxk_256.cu -> /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_with_attnbias_maxk_256.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_with_attnbias_maxk_32.cu -> /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_with_attnbias_maxk_32.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_with_attnbias_maxk_64.cu -> /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_grouped_infer_fp16_with_causalmask_with_attnbias_maxk_64.hip [skipped, already hipified]
/home/loong/Downloads/xformers/xformers/csrc/attention/matmul.cpp -> /home/loong/Downloads/xformers/xformers/csrc/attention/matmul.cpp [skipped, no changes]
/home/loong/Downloads/xformers/xformers/csrc/attention/sddmm.cpp -> /home/loong/Downloads/xformers/xformers/csrc/attention/sddmm.cpp [skipped, no changes]
/home/loong/Downloads/xformers/xformers/csrc/attention/sparse_softmax.cpp -> /home/loong/Downloads/xformers/xformers/csrc/attention/sparse_softmax.cpp [skipped, no changes]
/home/loong/Downloads/xformers/xformers/csrc/attention/spmm.cpp -> /home/loong/Downloads/xformers/xformers/csrc/attention/spmm.cpp [skipped, no changes]
/home/loong/Downloads/xformers/xformers/csrc/boxing_unboxing.cpp -> /home/loong/Downloads/xformers/xformers/csrc/boxing_unboxing.cpp [skipped, no changes]
/home/loong/Downloads/xformers/xformers/csrc/sequence_parallel_fused/memset_32b.cpp -> /home/loong/Downloads/xformers/xformers/csrc/sequence_parallel_fused/memset_32b.cpp [skipped, no changes]
/home/loong/Downloads/xformers/xformers/csrc/sequence_parallel_fused/synchronization.cpp -> /home/loong/Downloads/xformers/xformers/csrc/sequence_parallel_fused/synchronization.cpp [skipped, no changes]
/home/loong/Downloads/xformers/xformers/csrc/sparse24/sparse24.cpp -> /home/loong/Downloads/xformers/xformers/csrc/sparse24/sparse24.cpp [skipped, no changes]
/home/loong/Downloads/xformers/xformers/csrc/swiglu/swiglu_op.cpp -> /home/loong/Downloads/xformers/xformers/csrc/swiglu/swiglu_op.cpp [skipped, no changes]
/home/loong/Downloads/xformers/xformers/csrc/swiglu/swiglu_packedw.cpp -> /home/loong/Downloads/xformers/xformers/csrc/swiglu/swiglu_packedw.cpp [skipped, no changes]
Successfully preprocessed all matching files.
Total number of unsupported CUDA function calls: 0


Total number of replaced kernel launches: 8
running bdist_wheel
running build
running build_py
running build_ext
building 'xformers._C' extension
Emitting ninja build file /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/build.ninja...
Compiling objects...
Allowing ninja to set a default number of workers... (overridable by setting the environment variable MAX_JOBS=N)
[1/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_fmha_test.cu -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/ck_fmha_test.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc

@Looong01
Copy link
Author

Looong01 commented Apr 18, 2024

[2/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:734:13: error: static assertion failed due to requirement 'kKPack % K3 == 0'
            static_assert(kKPack % K3 == 0);
            ^             ~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:204:47: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQXKSVSCustomPolicy<true, false, false, 1, 1>::MakeVDramTileDistribution<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>' requested here
                             Policy::template MakeVDramTileDistribution<Problem>());
                                              ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:531:16: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>, ck::identity, ck::identity, ck::identity, ck::identity, ck::identity>' requested here
        return operator()(q_dram_block_window_tmp,
               ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha/ck_tiled_fmha_forward_kernel_hip.h:636:13: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>>' requested here
            FmhaPipeline{}(q_dram_window,
            ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:19:5: note: in instantiation of member function 'FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::operator()' requested here
    f(args...);
    ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:178:25: note: in instantiation of function template specialization 'kernel_wrapper<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    const auto kernel = kernel_wrapper<MaxThreadPerBlock, MinBlockPerCu, KernelImpl, Args...>;
                        ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:210:11: note: in instantiation of function template specialization 'launch_kernel<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    (void)launch_kernel<kBlockSize.x, kBlockPerCu>(
          ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:115:15: note: in instantiation of function template specialization 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, false, 256>::RunWithKernel<FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>>' requested here
              RunWithKernel<FmhaKernel>(param, stream);
              ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:232:14: note: in instantiation of member function 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, false, 256>::Run' requested here
      MaxK>::Run(param, stream);
             ^
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:26: error: constexpr if condition is not a constant expression
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                         ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:42: note: division by zero
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                                         ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: error: constexpr variable 'K1' must be initialized by a constant expression
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^    ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:56: note: division by zero
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                       ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:31: error: static assertion expression is not an integral constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                              ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:50: note: initializer of 'K1' is not a constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                                                 ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: error: non-type template argument is not a constant expression
                        Tuple<Sequence<N0, N1>, Sequence<K0, K1, K2, K3>>,
                                                             ^~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: note: initializer of 'K1' is not a constant expression
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:42: error: remainder by zero is undefined [-Werror,-Wdivision-by-zero]
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                                         ^ ~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:56: error: division by zero is undefined [-Werror,-Wdivision-by-zero]
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                       ^ ~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:833:9: error: static assertion failed due to requirement 'kKPack % K3 == 0'
        static_assert(kKPack % K3 == 0);
        ^             ~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:414:38: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQXKSVSCustomPolicy<true, false, false, 1, 1>::MakeShuffledVRegBlockDescriptor<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>' requested here
                    Policy::template MakeShuffledVRegBlockDescriptor<Problem>());
                                     ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:531:16: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>, ck::identity, ck::identity, ck::identity, ck::identity, ck::identity>' requested here
        return operator()(q_dram_block_window_tmp,
               ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha/ck_tiled_fmha_forward_kernel_hip.h:636:13: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>>' requested here
            FmhaPipeline{}(q_dram_window,
            ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:19:5: note: in instantiation of member function 'FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::operator()' requested here
    f(args...);
    ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:178:25: note: in instantiation of function template specialization 'kernel_wrapper<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    const auto kernel = kernel_wrapper<MaxThreadPerBlock, MinBlockPerCu, KernelImpl, Args...>;
                        ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:210:11: note: in instantiation of function template specialization 'launch_kernel<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    (void)launch_kernel<kBlockSize.x, kBlockPerCu>(
          ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:115:15: note: in instantiation of function template specialization 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, false, 256>::RunWithKernel<FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>>' requested here
              RunWithKernel<FmhaKernel>(param, stream);
              ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:232:14: note: in instantiation of member function 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, false, 256>::Run' requested here
      MaxK>::Run(param, stream);
             ^
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:835:22: error: constexpr if condition is not a constant expression
        if constexpr(get_warp_size() % (K2 * N0) == 0)
                     ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:835:38: note: division by zero
        if constexpr(get_warp_size() % (K2 * N0) == 0)
                                     ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:837:31: error: constexpr variable 'K1' must be initialized by a constant expression
            constexpr index_t K1 = get_warp_size() / (K2 * N0);
                              ^    ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:837:52: note: division by zero
            constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                   ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:842:85: error: non-type template argument is not a constant expression
                                               Tuple<Sequence<N0, N1>, Sequence<K0, K1, K2, K3>>,
                                                                                    ^~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:842:85: note: initializer of 'K1' is not a constant expression
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:837:31: note: declared here
            constexpr index_t K1 = get_warp_size() / (K2 * N0);
                              ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:835:38: error: remainder by zero is undefined [-Werror,-Wdivision-by-zero]
        if constexpr(get_warp_size() % (K2 * N0) == 0)
                                     ^ ~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:837:52: error: division by zero is undefined [-Werror,-Wdivision-by-zero]
            constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                   ^ ~~~~~~~~~
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:443:46: error: no matching member function for call to 'MakeShuffledVRegBlockDescriptor'
                            Policy::template MakeShuffledVRegBlockDescriptor<Problem>());
                            ~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:531:16: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>, ck::identity, ck::identity, ck::identity, ck::identity, ck::identity>' requested here
        return operator()(q_dram_block_window_tmp,
               ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha/ck_tiled_fmha_forward_kernel_hip.h:636:13: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>>' requested here
            FmhaPipeline{}(q_dram_window,
            ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:19:5: note: in instantiation of member function 'FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::operator()' requested here
    f(args...);
    ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:178:25: note: in instantiation of function template specialization 'kernel_wrapper<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    const auto kernel = kernel_wrapper<MaxThreadPerBlock, MinBlockPerCu, KernelImpl, Args...>;
                        ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:210:11: note: in instantiation of function template specialization 'launch_kernel<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    (void)launch_kernel<kBlockSize.x, kBlockPerCu>(
          ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:115:15: note: in instantiation of function template specialization 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, false, 256>::RunWithKernel<FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>>' requested here
              RunWithKernel<FmhaKernel>(param, stream);
              ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:232:14: note: in instantiation of member function 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, false, 256>::Run' requested here
      MaxK>::Run(param, stream);
             ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:818:47: note: candidate template ignored: substitution failure [with Problem = ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>]
    __host__ __device__ static constexpr auto MakeShuffledVRegBlockDescriptor()
                                              ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:734:13: error: static assertion failed due to requirement 'kKPack % K3 == 0'
            static_assert(kKPack % K3 == 0);
            ^             ~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:204:47: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQXKSVSCustomPolicy<true, false, false, 1, 1>::MakeVDramTileDistribution<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, false, true, 1>>>' requested here
                             Policy::template MakeVDramTileDistribution<Problem>());
                                              ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:531:16: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>, ck::identity, ck::identity, ck::identity, ck::identity, ck::identity>' requested here
        return operator()(q_dram_block_window_tmp,
               ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha/ck_tiled_fmha_forward_kernel_hip.h:636:13: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>>' requested here
            FmhaPipeline{}(q_dram_window,
            ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:19:5: note: in instantiation of member function 'FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::operator()' requested here
    f(args...);
    ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:178:25: note: in instantiation of function template specialization 'kernel_wrapper<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    const auto kernel = kernel_wrapper<MaxThreadPerBlock, MinBlockPerCu, KernelImpl, Args...>;
                        ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:210:11: note: in instantiation of function template specialization 'launch_kernel<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    (void)launch_kernel<kBlockSize.x, kBlockPerCu>(
          ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:115:15: note: in instantiation of function template specialization 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, false, 256>::RunWithKernel<FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>>' requested here
              RunWithKernel<FmhaKernel>(param, stream);
              ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:232:14: note: in instantiation of member function 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, false, 256>::Run' requested here
      MaxK>::Run(param, stream);
             ^
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:26: error: constexpr if condition is not a constant expression
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                         ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:42: note: division by zero
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                                         ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: error: constexpr variable 'K1' must be initialized by a constant expression
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^    ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:56: note: division by zero
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                       ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:31: error: static assertion expression is not an integral constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                              ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:50: note: initializer of 'K1' is not a constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                                                 ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: error: non-type template argument is not a constant expression
                        Tuple<Sequence<N0, N1>, Sequence<K0, K1, K2, K3>>,
                                                             ^~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: note: initializer of 'K1' is not a constant expression
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
fatal error: too many errors emitted, stopping now [-ferror-limit=]
20 errors generated when compiling for gfx1100.
[3/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:734:13: error: static assertion failed due to requirement 'kKPack % K3 == 0'
            static_assert(kKPack % K3 == 0);
            ^             ~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:204:47: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQXKSVSCustomPolicy<true, false, false, 1, 1>::MakeVDramTileDistribution<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>' requested here
                             Policy::template MakeVDramTileDistribution<Problem>());
                                              ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:531:16: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>, ck::identity, ck::identity, ck::identity, ck::identity, ck::identity>' requested here
        return operator()(q_dram_block_window_tmp,
               ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha/ck_tiled_fmha_forward_kernel_hip.h:636:13: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>>' requested here
            FmhaPipeline{}(q_dram_window,
            ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:19:5: note: in instantiation of member function 'FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::operator()' requested here
    f(args...);
    ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:178:25: note: in instantiation of function template specialization 'kernel_wrapper<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    const auto kernel = kernel_wrapper<MaxThreadPerBlock, MinBlockPerCu, KernelImpl, Args...>;
                        ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:210:11: note: in instantiation of function template specialization 'launch_kernel<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    (void)launch_kernel<kBlockSize.x, kBlockPerCu>(
          ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:115:15: note: in instantiation of function template specialization 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, true, 256>::RunWithKernel<FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>>' requested here
              RunWithKernel<FmhaKernel>(param, stream);
              ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:232:14: note: in instantiation of member function 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, true, 256>::Run' requested here
      MaxK>::Run(param, stream);
             ^
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:26: error: constexpr if condition is not a constant expression
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                         ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:42: note: division by zero
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                                         ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: error: constexpr variable 'K1' must be initialized by a constant expression
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^    ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:56: note: division by zero
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                       ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:31: error: static assertion expression is not an integral constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                              ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:50: note: initializer of 'K1' is not a constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                                                 ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: error: non-type template argument is not a constant expression
                        Tuple<Sequence<N0, N1>, Sequence<K0, K1, K2, K3>>,
                                                             ^~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: note: initializer of 'K1' is not a constant expression
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:42: error: remainder by zero is undefined [-Werror,-Wdivision-by-zero]
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                                         ^ ~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:56: error: division by zero is undefined [-Werror,-Wdivision-by-zero]
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                       ^ ~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:833:9: error: static assertion failed due to requirement 'kKPack % K3 == 0'
        static_assert(kKPack % K3 == 0);
        ^             ~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:414:38: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQXKSVSCustomPolicy<true, false, false, 1, 1>::MakeShuffledVRegBlockDescriptor<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>' requested here
                    Policy::template MakeShuffledVRegBlockDescriptor<Problem>());
                                     ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:531:16: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>, ck::identity, ck::identity, ck::identity, ck::identity, ck::identity>' requested here
        return operator()(q_dram_block_window_tmp,
               ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha/ck_tiled_fmha_forward_kernel_hip.h:636:13: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>>' requested here
            FmhaPipeline{}(q_dram_window,
            ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:19:5: note: in instantiation of member function 'FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::operator()' requested here
    f(args...);
    ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:178:25: note: in instantiation of function template specialization 'kernel_wrapper<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    const auto kernel = kernel_wrapper<MaxThreadPerBlock, MinBlockPerCu, KernelImpl, Args...>;
                        ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:210:11: note: in instantiation of function template specialization 'launch_kernel<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    (void)launch_kernel<kBlockSize.x, kBlockPerCu>(
          ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:115:15: note: in instantiation of function template specialization 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, true, 256>::RunWithKernel<FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>>' requested here
              RunWithKernel<FmhaKernel>(param, stream);
              ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:232:14: note: in instantiation of member function 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, true, 256>::Run' requested here
      MaxK>::Run(param, stream);
             ^
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:835:22: error: constexpr if condition is not a constant expression
        if constexpr(get_warp_size() % (K2 * N0) == 0)
                     ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:835:38: note: division by zero
        if constexpr(get_warp_size() % (K2 * N0) == 0)
                                     ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:837:31: error: constexpr variable 'K1' must be initialized by a constant expression
            constexpr index_t K1 = get_warp_size() / (K2 * N0);
                              ^    ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:837:52: note: division by zero
            constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                   ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:842:85: error: non-type template argument is not a constant expression
                                               Tuple<Sequence<N0, N1>, Sequence<K0, K1, K2, K3>>,
                                                                                    ^~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:842:85: note: initializer of 'K1' is not a constant expression
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:837:31: note: declared here
            constexpr index_t K1 = get_warp_size() / (K2 * N0);
                              ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:835:38: error: remainder by zero is undefined [-Werror,-Wdivision-by-zero]
        if constexpr(get_warp_size() % (K2 * N0) == 0)
                                     ^ ~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:837:52: error: division by zero is undefined [-Werror,-Wdivision-by-zero]
            constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                   ^ ~~~~~~~~~
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:443:46: error: no matching member function for call to 'MakeShuffledVRegBlockDescriptor'
                            Policy::template MakeShuffledVRegBlockDescriptor<Problem>());
                            ~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:531:16: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>, ck::identity, ck::identity, ck::identity, ck::identity, ck::identity>' requested here
        return operator()(q_dram_block_window_tmp,
               ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha/ck_tiled_fmha_forward_kernel_hip.h:636:13: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>>' requested here
            FmhaPipeline{}(q_dram_window,
            ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:19:5: note: in instantiation of member function 'FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::operator()' requested here
    f(args...);
    ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:178:25: note: in instantiation of function template specialization 'kernel_wrapper<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    const auto kernel = kernel_wrapper<MaxThreadPerBlock, MinBlockPerCu, KernelImpl, Args...>;
                        ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:210:11: note: in instantiation of function template specialization 'launch_kernel<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    (void)launch_kernel<kBlockSize.x, kBlockPerCu>(
          ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:115:15: note: in instantiation of function template specialization 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, true, 256>::RunWithKernel<FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>>' requested here
              RunWithKernel<FmhaKernel>(param, stream);
              ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:232:14: note: in instantiation of member function 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, true, 256>::Run' requested here
      MaxK>::Run(param, stream);
             ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:818:47: note: candidate template ignored: substitution failure [with Problem = ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, true, true, 1>>]
    __host__ __device__ static constexpr auto MakeShuffledVRegBlockDescriptor()
                                              ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:734:13: error: static assertion failed due to requirement 'kKPack % K3 == 0'
            static_assert(kKPack % K3 == 0);
            ^             ~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:204:47: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQXKSVSCustomPolicy<true, false, false, 1, 1>::MakeVDramTileDistribution<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, true, true, 1>>>' requested here
                             Policy::template MakeVDramTileDistribution<Problem>());
                                              ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:531:16: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, true, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>, ck::identity, ck::identity, ck::identity, ck::identity, ck::identity>' requested here
        return operator()(q_dram_block_window_tmp,
               ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha/ck_tiled_fmha_forward_kernel_hip.h:636:13: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, true, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>>' requested here
            FmhaPipeline{}(q_dram_window,
            ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:19:5: note: in instantiation of member function 'FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::operator()' requested here
    f(args...);
    ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:178:25: note: in instantiation of function template specialization 'kernel_wrapper<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    const auto kernel = kernel_wrapper<MaxThreadPerBlock, MinBlockPerCu, KernelImpl, Args...>;
                        ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:210:11: note: in instantiation of function template specialization 'launch_kernel<128, 1, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>, FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::FmhaFwdBatchModeKargs>' requested here
    (void)launch_kernel<kBlockSize.x, kBlockPerCu>(
          ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:115:15: note: in instantiation of function template specialization 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, true, 256>::RunWithKernel<FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, false, true, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>>' requested here
              RunWithKernel<FmhaKernel>(param, stream);
              ^
/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:232:14: note: in instantiation of member function 'batched_forward_causalmask_attnbias_dispatched<unsigned short, false, true, 256>::Run' requested here
      MaxK>::Run(param, stream);
             ^
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:26: error: constexpr if condition is not a constant expression
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                         ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:42: note: division by zero
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                                         ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: error: constexpr variable 'K1' must be initialized by a constant expression
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^    ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:56: note: division by zero
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                       ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:31: error: static assertion expression is not an integral constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                              ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:50: note: initializer of 'K1' is not a constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                                                 ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: error: non-type template argument is not a constant expression
                        Tuple<Sequence<N0, N1>, Sequence<K0, K1, K2, K3>>,
                                                             ^~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: note: initializer of 'K1' is not a constant expression
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
fatal error: too many errors emitted, stopping now [-ferror-limit=]
20 errors generated when compiling for gfx1100.
[4/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:734:13: error: static assertion failed due to requirement 'kKPack % K3 == 0'
            static_assert(kKPack % K3 == 0);
            ^             ~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:204:47: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQXKSVSCustomPolicy<true, false, false, 1, 1>::MakeVDramTileDistribution<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>' requested here
                             Policy::template MakeVDramTileDistribution<Problem>());
                                              ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:531:16: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>, ck::identity, ck::identity, ck::identity, ck::identity, ck::identity>' requested here
        return operator()(q_dram_block_window_tmp,
               ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha/ck_tiled_fmha_forward_kernel_hip.h:636:13: note: in instantiation of function template specialization 'ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>::operator()<ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 256>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>>, ck::Sequence<3, 4>, long, ck::Sequence<-1, -1, 32, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 32>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, const unsigned short, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int, int>, ck::Tuple<int, int>>, ck::PassThrough<int>, ck::PassThrough<int>, ck::RightPad<int, int>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>, ck::Sequence<2>, ck::Sequence<4>, ck::Sequence<3>>, ck::Tuple<ck::Sequence<1, 2>, ck::Sequence<3>, ck::Sequence<4>, ck::Sequence<5>, ck::Sequence<6>>, ck::Sequence<5, 6>, long, ck::Sequence<-1, -1, 32, -1, -1, -1, -1>, ck::Sequence<-1, -1, 1, -1, -1, -1, -1>>>, ck::Tuple<ck::integral_constant<int, 256>, ck::integral_constant<int, 32>>>, ck::tile_program::NullTileWindow<ck::Tuple<ck::integral_constant<int, 128>, ck::integral_constant<int, 128>>>, ck::tile_program::TileWindowWithStaticLengths<ck::TensorView<ck::BufferView<ck::AddressSpaceEnum::Global, float, long, true, ck::AmdBufferCoherenceEnum::DefaultCoherence>, ck::TensorDescriptor<ck::Tuple<ck::Embed<ck::Tuple<int>, ck::Tuple<int>>, ck::RightPad<int, int>>, ck::Tuple<ck::Sequence<0>, ck::Sequence<1>>, ck::Tuple<ck::Sequence<1>, ck::Sequence<2>>, ck::Sequence<2>, long, ck::Sequence<-1, 1, -1>, ck::Sequence<-1, 1, -1>>>, ck::Tuple<ck::integral_constant<int, 128>>>>' requested here
            FmhaPipeline{}(q_dram_window,
            ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/host_utility/kernel_launch_hip.hpp:19:5: note: in instantiation of member function 'FmhaFwdKernel<FmhaFwdTilePartitioner<FmhaFwdShape<256>>, ck::tile_program::block::BlockFmhaPipelineQRKSVS<ck::tile_program::block::BlockFmhaPipelineProblem<unsigned short, unsigned short, unsigned short, float, float, unsigned short, float, unsigned short, float, unsigned short, FmhaFwdShape<256>, false, ck::tile_program::block::GenericAttentionMask<true, true>, ck::tile_program::TileFmhaTraits<true, true, true, true, false, true, 1>>>, FmhaFwdEpilogue<FmhaFwdEpilogueProblem<float, unsigned short>>>::operator()' requested here
    f(args...);

@Looong01
Copy link
Author

In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:26: error: constexpr if condition is not a constant expression
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                         ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:42: note: division by zero
            if constexpr(get_warp_size() % (K2 * N0) == 0)
                                         ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: error: constexpr variable 'K1' must be initialized by a constant expression
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^    ~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:56: note: division by zero
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                                       ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:31: error: static assertion expression is not an integral constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                              ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:740:50: note: initializer of 'K1' is not a constant expression
                static_assert(kKPerBlock == K0 * K1 * K2 * K3);
                                                 ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: error: non-type template argument is not a constant expression
                        Tuple<Sequence<N0, N1>, Sequence<K0, K1, K2, K3>>,
                                                             ^~
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: note: initializer of 'K1' is not a constant expression
/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here
                constexpr index_t K1 = get_warp_size() / (K2 * N0);
                                  ^
fatal error: too many errors emitted, stopping now [-ferror-limit=]
20 errors generated when compiling for gfx1100.
[5/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:18:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_hip.hpp:10:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_attribute_mfma_hip.hpp:13:
/home/loong/Downloads/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 gfx1100.
[6/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_64.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_64.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_64.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_64.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_64.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_64.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:18:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_hip.hpp:10:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_attribute_mfma_hip.hpp:13:
/home/loong/Downloads/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 gfx1100.
[7/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_32.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_32.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_32.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_32.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_32.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_32.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:18:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_hip.hpp:10:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_attribute_mfma_hip.hpp:13:
/home/loong/Downloads/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 gfx1100.
[8/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_32.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_32.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_32.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_32.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_32.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_32.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:18:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_hip.hpp:10:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_attribute_mfma_hip.hpp:13:
/home/loong/Downloads/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 gfx1100.
[9/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_128.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_128.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_128.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_128.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_128.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_128.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:18:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_hip.hpp:10:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_attribute_mfma_hip.hpp:13:
/home/loong/Downloads/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 gfx1100.
[10/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_128.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_128.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_128.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_128.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_128.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_128.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:18:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_hip.hpp:10:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_attribute_mfma_hip.hpp:13:
/home/loong/Downloads/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 gfx1100.
[11/139] /opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_128.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_128.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
FAILED: /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_128.o
/opt/rocm-6.0.2/bin/hipcc  -I/home/loong/Downloads/xformers/xformers/csrc -I/home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/torch/csrc/api/include -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/TH -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THC -I/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/include/THH -I/opt/rocm-6.0.2/include -I/home/loong/miniconda3/envs/CV/include/python3.9 -c -c /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_128.hip -o /home/loong/Downloads/xformers/build/temp.linux-x86_64-cpython-39/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_128.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_with_attnbias_maxk_128.hip:10:
In file included from /home/loong/Downloads/xformers/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:18:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_hip.hpp:10:
In file included from /home/loong/Downloads/xformers/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_attribute_mfma_hip.hpp:13:
/home/loong/Downloads/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 gfx1100.
ninja: build stopped: subcommand failed.
Traceback (most recent call last):
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 2107, in _run_ninja_build
    subprocess.run(
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/subprocess.py", line 528, in run
    raise CalledProcessError(retcode, process.args,
subprocess.CalledProcessError: Command '['ninja', '-v']' returned non-zero exit status 1.

The above exception was the direct cause of the following exception:

Traceback (most recent call last):
  File "/home/loong/Downloads/xformers/setup.py", line 485, in <module>
    setuptools.setup(
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/__init__.py", line 103, in setup
    return distutils.core.setup(**attrs)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/core.py", line 185, in setup
    return run_commands(dist)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/core.py", line 201, in run_commands
    dist.run_commands()
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/dist.py", line 969, in run_commands
    self.run_command(cmd)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/dist.py", line 989, in run_command
    super().run_command(command)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/dist.py", line 988, in run_command
    cmd_obj.run()
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/wheel/bdist_wheel.py", line 364, in run
    self.run_command("build")
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/cmd.py", line 318, in run_command
    self.distribution.run_command(command)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/dist.py", line 989, in run_command
    super().run_command(command)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/dist.py", line 988, in run_command
    cmd_obj.run()
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/command/build.py", line 131, in run
    self.run_command(cmd_name)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/cmd.py", line 318, in run_command
    self.distribution.run_command(command)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/dist.py", line 989, in run_command
    super().run_command(command)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/dist.py", line 988, in run_command
    cmd_obj.run()
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/command/build_ext.py", line 88, in run
    _build_ext.run(self)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/command/build_ext.py", line 345, in run
    self.build_extensions()
  File "/home/loong/Downloads/xformers/setup.py", line 442, in build_extensions
    super().build_extensions()
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 870, in build_extensions
    build_ext.build_extensions(self)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/command/build_ext.py", line 467, in build_extensions
    self._build_extensions_serial()
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/command/build_ext.py", line 493, in _build_extensions_serial
    self.build_extension(ext)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/command/build_ext.py", line 249, in build_extension
    _build_ext.build_extension(self, ext)
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/setuptools/_distutils/command/build_ext.py", line 548, in build_extension
    objects = self.compiler.compile(
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 683, in unix_wrap_ninja_compile
    _write_ninja_file_and_compile_objects(
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 1783, in _write_ninja_file_and_compile_objects
    _run_ninja_build(
  File "/home/loong/miniconda3/envs/CV/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 2123, in _run_ninja_build
    raise RuntimeError(message) from e
RuntimeError: Error compiling objects for extension

@lhl
Copy link

lhl commented Apr 20, 2024

FYI, I have a W7900 (gfx1100) that I am trying to build with and getting errors with both upstream and using the ROCm fork.

PyTorch/HIP version:

$ pip install --pre torch torchvision torchaudio --index-url https://download.pytorch.org/whl/nightly/rocm6.0
$ python -c "import torch; print(torch.__version__); print(torch.version.hip)"
2.4.0.dev20240420+rocm6.0
6.0.32830-d62f6a171

Errors:

$ pip wheel -v --no-build-isolation git+https://github.com/ROCm/xformers.git@main#egg=xformers

  [18/156] /opt/rocm-6.0.0/bin/hipcc  -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha -I/tmp/pip-wheel-12djh
_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include -I/home/lhl
/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/torch/csrc/api/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/to
rch/include/TH -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THC -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THH -I/opt/rocm-6.0.0/include -I/home/lhl/miniforge3/envs/exll
amav2/include/python3.11 -c -c /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip -o /tmp/pip-wheel-12djh_kv/xf
ormers_5dafbdb166d3479187353a2b7fd8f12b/build/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLA
S_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werr
or -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gp
u-rdc                                                                                                                                                                                                                                              
  FAILED: /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/build/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.o                 
  /opt/rocm-6.0.0/bin/hipcc  -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha -I/tmp/pip-wheel-12djh_kv/xform
ers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include -I/home/lhl/miniforg
e3/envs/exllamav2/lib/python3.11/site-packages/torch/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/torch/csrc/api/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/inclu
de/TH -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THC -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THH -I/opt/rocm-6.0.0/include -I/home/lhl/miniforge3/envs/exllamav2/inc
lude/python3.11 -c -c /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip -o /tmp/pip-wheel-12djh_kv/xformers_5d
afbdb166d3479187353a2b7fd8f12b/build/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCU
DA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Wover
loaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc    
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip:10:                                 
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:                                                                                 
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:                                 
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:                   
  /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:734:13: error: static assertion failed due t
o requirement 'kKPack % K3 == 0'                                                                                                                                                                                                                   
              static_assert(kKPack % K3 == 0);                                                                                                                                                                                                     
              ^             ~~~~~~~~~~~~~~~~

...

  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip:10:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
  /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:736:26: error: constexpr if condition is not
 a constant expression
              if constexpr(get_warp_size() % (K2 * N0) == 0) 
                           ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

...

  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip:10:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
  /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:835:22: error: constexpr if condition is not
 a constant expression
          if constexpr(get_warp_size() % (K2 * N0) == 0)
                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

...

  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_256.hip:10:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
  /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:443:46: error: no matching member function for call to 'Ma
keShuffledVRegBlockDescriptor'
                              Policy::template MakeShuffledVRegBlockDescriptor<Problem>());
                              ~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

...
(well, a lot)

 /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:744:62: note: initializer of 'K1' is not a c
onstant expression                                                                                                                                                                                                                                 
  /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:738:35: note: declared here                 
                  constexpr index_t K1 = get_warp_size() / (K2 * N0);                                                                                                                                                                              
                                    ^                                                                                                                                                                                                              
  fatal error: too many errors emitted, stopping now [-ferror-limit=]         
  20 errors generated when compiling for gfx1100. 

...

  [19/156] /opt/rocm-6.0.0/bin/hipcc  -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha -I/tmp/pip-wheel-12djh
_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include -I/home/lhl
/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/torch/csrc/api/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/to
rch/include/TH -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THC -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THH -I/opt/rocm-6.0.0/include -I/home/lhl/miniforge3/envs/exll
amav2/include/python3.11 -c -c /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.hip -o /tmp/pip-wheel-12djh_kv/
xformers_5dafbdb166d3479187353a2b7fd8f12b/build/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHI
PBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -
Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fn
o-gpu-rdc    
  FAILED: /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/build/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.o               
  /opt/rocm-6.0.0/bin/hipcc  -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha -I/tmp/pip-wheel-12djh_kv/xform
ers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include -I/home/lhl/miniforg
e3/envs/exllamav2/lib/python3.11/site-packages/torch/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/torch/csrc/api/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/inclu
de/TH -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THC -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THH -I/opt/rocm-6.0.0/include -I/home/lhl/miniforge3/envs/exllamav2/inc
lude/python3.11 -c -c /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.hip -o /tmp/pip-wheel-12djh_kv/xformers_
5dafbdb166d3479187353a2b7fd8f12b/build/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 
-DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -W
overloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_with_causalmask_no_attnbias_maxk_256.hip:10:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:                                                
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:19:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_default_policy_hip.hpp:7:
  /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qx_ks_vs_custom_policy_hip.hpp:734:13: error: static assertion failed due t
o requirement 'kKPack % K3 == 0'                                                                                                                                                                                                                   
              static_assert(kKPack % K3 == 0);
              ^             ~~~~~~~~~~~~~~~~

...

  [21/156] /opt/rocm-6.0.0/bin/hipcc  -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha -I/tmp/pip-wheel-12djh
_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include -I/home/lhl
/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/torch/csrc/api/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/to
rch/include/TH -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THC -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THH -I/opt/rocm-6.0.0/include -I/home/lhl/miniforge3/envs/exll
amav2/include/python3.11 -c -c /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/attention_forward_generic_ck_tiled.hip -o /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/build
/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/attention_forward_generic_ck_tiled.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -
std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11
_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc                                                                                     
  [22/156] /opt/rocm-6.0.0/bin/hipcc  -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha -I/tmp/pip-wheel-12djh
_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include -I/home/lhl
/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/torch/csrc/api/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/to
rch/include/TH -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THC -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THH -I/opt/rocm-6.0.0/include -I/home/lhl/miniforge3/envs/exll
amav2/include/python3.11 -c -c /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.hip -o /tmp/pip-wheel-12djh_kv/xfo
rmers_5dafbdb166d3479187353a2b7fd8f12b/build/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_
V2 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror
 -Woverloaded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-
rdc
  FAILED: /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/build/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.o
  /opt/rocm-6.0.0/bin/hipcc  -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha -I/tmp/pip-wheel-12djh_kv/xform
ers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/example/91_tile_program/xformers_fmha -I/tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include -I/home/lhl/miniforg
e3/envs/exllamav2/lib/python3.11/site-packages/torch/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/torch/csrc/api/include -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/inclu
de/TH -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THC -I/home/lhl/miniforge3/envs/exllamav2/lib/python3.11/site-packages/torch/include/THH -I/opt/rocm-6.0.0/include -I/home/lhl/miniforge3/envs/exllamav2/inc
lude/python3.11 -c -c /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.hip -o /tmp/pip-wheel-12djh_kv/xformers_5da
fbdb166d3479187353a2b7fd8f12b/build/temp.linux-x86_64-cpython-311/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.o -fPIC -D__HIP_PLATFORM_AMD__=1 -DUSE_ROCM=1 -DHIPBLAS_V2 -DCUDA
_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 -O3 -std=c++17 --offload-arch=native -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -DCK_FMHA_FWD_FAST_EXP2=1 -fgpu-flush-denormals-to-zero -Werror -Woverlo
aded-virtual -DBUILD_PYTHON_PACKAGE -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_C -D_GLIBCXX_USE_CXX11_ABI=0 -fno-gpu-rdc
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/instances/ck_tiled_fmha_batched_forward_bp16_no_causalmask_no_attnbias_maxk_64.hip:10:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/xformers/csrc/attention/hip_fmha/ck_tiled_fmha_batched_forward_hip.h:23:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/block_tile_pipeline/block_fmha_pipeline_qr_ks_vs_hip.hpp:18:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_hip.hpp:10:
  In file included from /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/third_party/composable_kernel_tiled/include/ck/tile_program/warp_tile/warp_gemm_attribute_mfma_hip.hpp:13:
  /tmp/pip-wheel-12djh_kv/xformers_5dafbdb166d3479187353a2b7fd8f12b/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 tar
get 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 gfx1100.

A lot of similar errors so I've elided but definitely busted. Maybe @tenpercent has some feedback/thoughts on getting the ROCm fork running?

I am on a clean mamba environment with Ubuntu 22.04 LTS HWE and ROCm installed via: https://rocm.docs.amd.com/projects/install-on-linux/en/latest/how-to/native-install/ubuntu.html

@lukedupin
Copy link

My setup on arch is failing, but the docker example worked for me.

@tenpercent
Copy link
Contributor

tenpercent commented May 29, 2024

It's not supposed to work on gfx10xx or gfx11xx yet (as we haven't had developers' bandwidth to support it). The root cause is different matrix core instructions being used between different gfx families, as the compiler error is trying to hint.
The installation should work if you could get access to any MI2xx or MI3xx card (the datacenter GPU series)
I'll also leave a helpful link for gfx version deciphering

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

5 participants