Build with CUDA 12.6 and VS 2022 17.11.1 broken on Windows

[levicki]

🐛 Bug

I am trying to build xformers from the main branch and no matter whether I try 0.0.27.post2 tag or the latest commit I am getting build errors.

Command

git clone https://github.com/facebookresearch/xformers.git cd xformers git submodule update --init --recursive python -m venv venv venv\scripts\activate pip install --no-color torch==2.4.0+cu124 torchvision==0.19.0+cu124 --extra-index-url https://download.pytorch.org/whl/cu124 pip install --no-color -r requirements.txt pip install --no-color wheel set TORCH_CUDA_ARCH_LIST=8.9 set DISTUTILS_USE_SDK=1 python setup.py build python setup.py bdist_wheel

To Reproduce

Steps to reproduce the behavior:

1. Install Visual Studio 2022 17.11.1 with C++ Desktop Development workload
2. Install Python 3.11.9
3. Install CUDA toolkit 12.6
4. Use the commands above to build xformers from VS2022 x64 command prompt

ninja: build stopped: subcommand failed.
Traceback (most recent call last):
  File "C:\BUILD\xformers\venv\Lib\site-packages\torch\utils\cpp_extension.py", line 2105, in _run_ninja_build
    subprocess.run(
  File "C:\PROGRAMS\Python\Lib\subprocess.py", line 571, in run
    raise CalledProcessError(retcode, process.args,
subprocess.CalledProcessError: Command '['ninja', '-v', '-j', '12']' returned non-zero exit status 1.

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

Traceback (most recent call last):
  File "C:\BUILD\xformers\setup.py", line 558, in <module>
    setuptools.setup(
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\__init__.py", line 87, in setup
    return distutils.core.setup(**attrs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\core.py", line 185, in setup
    return run_commands(dist)
           ^^^^^^^^^^^^^^^^^^
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\core.py", line 201, in run_commands
    dist.run_commands()
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\dist.py", line 968, in run_commands
    self.run_command(cmd)
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\dist.py", line 1217, in run_command
    super().run_command(command)
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\dist.py", line 987, in run_command
    cmd_obj.run()
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\command\build.py", line 132, in run
    self.run_command(cmd_name)
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\cmd.py", line 319, in run_command
    self.distribution.run_command(command)
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\dist.py", line 1217, in run_command
    super().run_command(command)
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\dist.py", line 987, in run_command
    cmd_obj.run()
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\command\build_ext.py", line 84, in run
    _build_ext.run(self)
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\command\build_ext.py", line 346, in run
    self.build_extensions()
  File "C:\BUILD\xformers\setup.py", line 515, in build_extensions
    super().build_extensions()
  File "C:\BUILD\xformers\venv\Lib\site-packages\torch\utils\cpp_extension.py", line 866, in build_extensions
    build_ext.build_extensions(self)
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\command\build_ext.py", line 466, in build_extensions
    self._build_extensions_serial()
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\command\build_ext.py", line 492, in _build_extensions_serial
    self.build_extension(ext)
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\command\build_ext.py", line 246, in build_extension
    _build_ext.build_extension(self, ext)
  File "C:\BUILD\xformers\venv\Lib\site-packages\setuptools\_distutils\command\build_ext.py", line 547, in build_extension
    objects = self.compiler.compile(
              ^^^^^^^^^^^^^^^^^^^^^^
  File "C:\BUILD\xformers\venv\Lib\site-packages\torch\utils\cpp_extension.py", line 838, in win_wrap_ninja_compile
    _write_ninja_file_and_compile_objects(
  File "C:\BUILD\xformers\venv\Lib\site-packages\torch\utils\cpp_extension.py", line 1785, in _write_ninja_file_and_compile_objects
    _run_ninja_build(
  File "C:\BUILD\xformers\venv\Lib\site-packages\torch\utils\cpp_extension.py", line 2121, in _run_ninja_build
    raise RuntimeError(message) from e
RuntimeError: Error compiling objects for extension

The stack trace doesn't seem to capture what's really going on, here are the relevent errors from build output before that traceback:

C:/BUILD/xformers/third_party/cutlass/include\cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp(136): warning C4346: 'SharedStorage': dependent name is not a type
C:/BUILD/xformers/third_party/cutlass/include\cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp(136): note: prefix the qualified-id with 'typename' to indicate a type
C:/BUILD/xformers/third_party/cutlass/include\cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp(136): note: the template instantiation context (the oldest one first) is
C:/BUILD/xformers/third_party/cutlass/include\cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp(60): note: while compiling class template partial specialization 'cutlass::gemm::kernel::GemmUniversal<ProblemShape_,CollectiveMainloop_,CollectiveEpilogue_,TileScheduler_,enable_if<std::is_base_of_v<cutlass::gemm::KernelTmaWarpSpecializedPingpong,CollectiveMainloop_::DispatchPolicy::Schedule>,void>::type>'
C:/BUILD/xformers/third_party/cutlass/include\cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp(124): note: while compiling class 'cutlass::gemm::kernel::GemmUniversal<ProblemShape_,CollectiveMainloop_,CollectiveEpilogue_,TileScheduler_,enable_if<std::is_base_of_v<cutlass::gemm::KernelTmaWarpSpecializedPingpong,CollectiveMainloop_::DispatchPolicy::Schedule>,void>::type>::SharedStorage'
C:/BUILD/xformers/third_party/cutlass/include\cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp(133): note: while compiling class 'cutlass::gemm::kernel::GemmUniversal<ProblemShape_,CollectiveMainloop_,CollectiveEpilogue_,TileScheduler_,enable_if<std::is_base_of_v<cutlass::gemm::KernelTmaWarpSpecializedPingpong,CollectiveMainloop_::DispatchPolicy::Schedule>,void>::type>::SharedStorage::PipelineStorage'
C:/BUILD/xformers/third_party/cutlass/include\cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp(136): error C2061: syntax error: identifier 'SharedStorage'
C:/BUILD/xformers/third_party/cutlass/include\cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp(140): error C3646: 'math_wg_order': unknown override specifier
C:/BUILD/xformers/third_party/cutlass/include\cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp(140): error C4430: missing type specifier - int assumed. Note: C++ does not support default-int

Expected behavior

Build should finish successfully

Environment

• PyTorch Version (e.g., 1.0): 2.4.0
• OS (e.g., Linux): Windows 11 (Version 10.0.22631.4037)
• How you installed PyTorch (conda, pip, source): see command
• Build command you used (if compiling from source): see command
• Python version: 3.11.9
• CUDA/cuDNN version: 12.6
• GPU models and configuration: RTX 4090
• Any other relevant information: This building method worked with all xformers<=v0.0.27

Additional context

N/A