Open Looong01 opened 4 months ago
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 ./
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!
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: https://github.com/facebookresearch/xformers/pull/978#issuecomment-1965489838
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 🙏
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
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
[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...);
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
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
My setup on arch is failing, but the docker example worked for me.
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
🐛 Bug
I put up some screenshots that I think it is important:
And
Command
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