Not able to install 2.0

[akanyaani]

Tried pip install and setup.py install both

/home/yellow/flash-attention/csrc/cutlass/include/cute/stride.hpp(112): warning: calling a __host__ function("__builtin_unreachable") from a __host__ __device__ function("cute::crd2idx< ::cute::tuple< ::cute::Underscore, ::cute::Underscore, int > , ::cute::tuple< ::cute::tuple< ::cute::constant<int, (int)2> , ::cute::constant<int, (int)2> , ::cute::constant<int, (int)2> > , ::cute::constant<int, (int)2> , ::cute::tuple< ::cute::constant<int, (int)2> , ::cute::constant<int, (int)2> > > , ::cute::tuple< ::cute::tuple< ::cute::constant<int, (int)1> , ::cute::constant<int, (int)2> , ::cute::constant<int, (int)4> > , ::cute::constant<int, (int)8> , ::cute::tuple< ::cute::constant<int, (int)16> , ::cute::constant<int, (int)32> > > > ") is not allowed

Killed

txas info : Used 255 registers, 576 bytes cmem[0]
ninja: build stopped: subcommand failed.
Traceback (most recent call last):
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 1893, in _run_ninja_build
 subprocess.run(
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/subprocess.py", line 524, 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/yellow/flash-attention/setup.py", line 201, in <module>
 setup(
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/__init__.py", line 107, in setup
 return distutils.core.setup(**attrs)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/core.py", line 185, in setup
 return run_commands(dist)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/core.py", line 201, in run_commands
 dist.run_commands()
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/dist.py", line 969, in run_commands
 self.run_command(cmd)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/dist.py", line 1234, in run_command
 super().run_command(command)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/dist.py", line 988, in run_command
 cmd_obj.run()
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/command/install.py", line 80, in run
 self.do_egg_install()
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/command/install.py", line 129, in do_egg_install
 self.run_command('bdist_egg')
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/cmd.py", line 318, in run_command
 self.distribution.run_command(command)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/dist.py", line 1234, in run_command
 super().run_command(command)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/dist.py", line 988, in run_command
 cmd_obj.run()
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/command/bdist_egg.py", line 164, in run
 cmd = self.call_command('install_lib', warn_dir=0)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/command/bdist_egg.py", line 150, in call_command
 self.run_command(cmdname)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/cmd.py", line 318, in run_command
 self.distribution.run_command(command)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/dist.py", line 1234, in run_command
 super().run_command(command)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/dist.py", line 988, in run_command
 cmd_obj.run()
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/command/install_lib.py", line 11, in run
 self.build()
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/command/install_lib.py", line 111, in build
 self.run_command('build_ext')
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/cmd.py", line 318, in run_command
 self.distribution.run_command(command)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/dist.py", line 1234, in run_command
 super().run_command(command)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/dist.py", line 988, in run_command
 cmd_obj.run()
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/command/build_ext.py", line 84, in run
 _build_ext.run(self)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/command/build_ext.py", line 345, in run
 self.build_extensions()
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 843, in build_extensions
 build_ext.build_extensions(self)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/command/build_ext.py", line 467, in build_extensions
 self._build_extensions_serial()
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/command/build_ext.py", line 493, in _build_extensions_serial
 self.build_extension(ext)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/command/build_ext.py", line 246, in build_extension
 _build_ext.build_extension(self, ext)
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/setuptools/_distutils/command/build_ext.py", line 548, in build_extension
 objects = self.compiler.compile(
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 658, in unix_wrap_ninja_compile
 _write_ninja_file_and_compile_objects(
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 1574, in _write_ninja_file_and_compile_objects
 _run_ninja_build(
 File "/home/yellow/anaconda3/envs/fla2/lib/python3.9/site-packages/torch/utils/cpp_extension.py", line 1909, in _run_ninja_build
 raise RuntimeError(message) from e
RuntimeError: Error compiling objects for extension

[lms21a]

Same issue--not a one off

[tmm1]

I'm seeing a similar issue:

 ptxas info : Function properties for _Z25flash_bwd_dot_do_o_kernelILb1E23Flash_bwd_kernel_traitsILi128ELi64ELi128E
Li8ELi2ELi4ELi2ELb0ELb0EN7cutlass6half_tE19Flash_kernel_traitsILi128ELi64ELi128ELi8ES2_EEEv16Flash_bwd_params 
 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
 ptxas info : Used 49 registers, 688 bytes cmem[0] 
 ninja: build stopped: subcommand failed. 
 Traceback (most recent call last): 
 File "/home/tmm1/.local/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 1893, in _run_ninja_build 
 subprocess.run( 
 File "/usr/lib/python3.10/subprocess.py", line 524, in run 
 raise CalledProcessError(retcode, process.args, 
 subprocess.CalledProcessError: Command '['ninja', '-v']' returned non-zero exit status 1. 

[timmoon10]

I also reproduce the error. Setting MAX_JOBS=1 in the environment fixes it for me, so it seems that compilation has become resource-intensive enough for a parallel Ninja build to overwhelm many systems. It's a long-term question, but I wonder if the current approach of statically-compiled CUDA kernels is sustainable. Perhaps there is value to considering JIT compilation, e.g. with Triton or NVRTC?

[tridao]

Yeah I personally don't like the fact that we're templating so heavily (for dropout / no dropout, causal / not causal, different head dimensions, whether seqlen is divisible by 128 or not, different GPU types). The goal has been to get maximum performance, perhaps at the expense of compilation time.

• Agree with you that JIT compilation is interesting. I don't have any experience there however.
• Another way is to have pre-built wheels that folks can just download. I'll get to that once I'm done fixing some of the edge cases with the backward pass.

[timmoon10]

Yea, it's definitely a hard problem that we've also been hitting in Transformer Engine.

It'll take some engineering effort, but I've found NVRTC to be a nice way to avoid the combinatorial explosion from templating. I've found that with the right wrapper classes (see rtc.h and rtc.cpp), it can be straightforward to write and launch JIT kernels (see how transpose.cu calls the kernel in rtc/transpose.cu). It does impose compilation time the first time each kernel is called though, and it is quite annoying to include external headers (including the CUDA Toolkit and C++ Standard Library).

[lms21a]

I also reproduce the error. Setting MAX_JOBS=1 in the environment fixes it for me, so it seems that compilation has become resource-intensive enough for a parallel Ninja build to overwhelm many systems. It's a long-term question, but I wonder if the current approach of statically-compiled CUDA kernels is sustainable. Perhaps there is value to considering JIT compilation, e.g. with Triton or NVRTC?

Unfortunately for me, I still get the same error even if MAX_JOBS=1 is set. Also tried building from source with the same error. Any temporary solution available?

[tmm1]

I tried with MAX_JOBS=1 and see:

 /tmp/pip-install-mvmpcn8m/flash-attn_39e3b6d4eaad444b8beae75cad8aecd2/csrc/flash_attn/src/flash_bwd_launch_template.h
(179): error: expression must have a constant value
 Note #2767-D: the value of *this cannot be used as a constant
 detected during instantiation of "void run_mha_bwd_hdim32<T>(Flash_bwd_params &, cudaStream_t, __nv_bool) [
with T=cutlass::bfloat16_t]"
 /tmp/pip-install-mvmpcn8m/flash-attn_39e3b6d4eaad444b8beae75cad8aecd2/csrc/flash_attn/src/flash_bwd_hdim32_bf16_sm80.
cu(15): here
 
 /tmp/pip-install-mvmpcn8m/flash-attn_39e3b6d4eaad444b8beae75cad8aecd2/csrc/flash_attn/src/flash_bwd_launch_template.h
(179): error: expression must have a constant value
 Note #2767-D: the value of *this cannot be used as a constant
 detected during instantiation of "void run_mha_bwd_hdim32<T>(Flash_bwd_params &, cudaStream_t, __nv_bool) [
with T=cutlass::bfloat16_t]"
 /tmp/pip-install-mvmpcn8m/flash-attn_39e3b6d4eaad444b8beae75cad8aecd2/csrc/flash_attn/src/flash_bwd_hdim32_bf16_sm80.
cu(15): here
 
 2 errors detected in the compilation of "/tmp/pip-install-mvmpcn8m/flash-attn_39e3b6d4eaad444b8beae75cad8aecd2/csrc/f
lash_attn/src/flash_bwd_hdim32_bf16_sm80.cu".

EDIT: That looks like #343

[MiladInk]

seeing the same problem here:

4, in run
 _build_ext.run(self)
 File "/home/mila/a/aghajohm/scratch/.conda/envs/competenlp/lib/python3.10/site-packages/setuptools/_distutils/command/build_ext.py", line 345, in run
 self.build_extensions()
 File "/home/mila/a/aghajohm/scratch/.conda/envs/competenlp/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 843, in build_extensions
 build_ext.build_extensions(self)
 File "/home/mila/a/aghajohm/scratch/.conda/envs/competenlp/lib/python3.10/site-packages/setuptools/_distutils/command/build_ext.py", line 467, in build_extensions
 self._build_extensions_serial()
 File "/home/mila/a/aghajohm/scratch/.conda/envs/competenlp/lib/python3.10/site-packages/setuptools/_distutils/command/build_ext.py", line 493, in _build_extensions_serial
 self.build_extension(ext)
 File "/home/mila/a/aghajohm/scratch/.conda/envs/competenlp/lib/python3.10/site-packages/setuptools/command/build_ext.py", line 246, in build_extension
 _build_ext.build_extension(self, ext)
 File "/home/mila/a/aghajohm/scratch/.conda/envs/competenlp/lib/python3.10/site-packages/setuptools/_distutils/command/build_ext.py", line 548, in build_extension
 objects = self.compiler.compile(
 File "/home/mila/a/aghajohm/scratch/.conda/envs/competenlp/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 658, in unix_wrap_ninja_compile
 _write_ninja_file_and_compile_objects(
 File "/home/mila/a/aghajohm/scratch/.conda/envs/competenlp/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 1574, in _write_ninja_file_and_compile_objects
 _run_ninja_build(
 File "/home/mila/a/aghajohm/scratch/.conda/envs/competenlp/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 1909, in _run_ninja_build
 raise RuntimeError(message) from e
RuntimeError: Error compiling objects for extension

I am setting MAX_JOBS=1 to see if it improves performance.

update: setting MAX_JOBS=1 solved it for me. It got installed albeit slowly.

[Desein-Yang]

I got the same error and I modify the setup.py files, ninja -v --> ninja --version. Then I met /cognitive_comp/yangqi/images/flash-attention-main/build/temp.linux-x86_64-cpython-38/csrc/flash_attn/src/flash_fwd_hdim96_fp16_sm80.o: No such file or directory

[mtisz]

Same issue here..

[ilikenwf]

I got the same error and I modify the setup.py files, ninja -v --> ninja --version. Then I met /cognitive_comp/yangqi/images/flash-attention-main/build/temp.linux-x86_64-cpython-38/csrc/flash_attn/src/flash_fwd_hdim96_fp16_sm80.o: No such file or directory

Same here.

Despite even attempting a git submodule update --init --recursive --force I still run into this issue...so if anyone finds anything, thank you very much.

[mtisz]

Pretty sure this is an issue with CUDA 12 -- but I don't know whether there's a timeline to support CUDA 12

[BoxiangW]

Same issue here

[BoxiangW]

Same issue here

I solved it by updating CUDA version, hope it helps