setup.py

"""
Copyright (c) 2024 by SageAttention team.

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
"""

import os
import subprocess
import threading
from packaging.version import parse, Version
import warnings

from setuptools import setup, find_packages
import torch
from torch.utils.cpp_extension import BuildExtension, CUDAExtension, CUDA_HOME

HAS_SM80 = False
HAS_SM86 = False
HAS_SM89 = False
HAS_SM90 = False
HAS_SM120 = False

# Supported NVIDIA GPU architectures.
SUPPORTED_ARCHS = {"8.0", "8.6", "8.9", "9.0", "12.0"}

# Compiler flags.
CXX_FLAGS = ["-g", "-O3", "-fopenmp", "-lgomp", "-std=c++17", "-DENABLE_BF16"]
NVCC_FLAGS = [
    "-O3",
    "-std=c++17",
    "-U__CUDA_NO_HALF_OPERATORS__",
    "-U__CUDA_NO_HALF_CONVERSIONS__",
    "--use_fast_math",
    "--threads=8",
    "-Xptxas=-v",
    "-diag-suppress=174", # suppress the specific warning
]

ABI = 1 if torch._C._GLIBCXX_USE_CXX11_ABI else 0
CXX_FLAGS += [f"-D_GLIBCXX_USE_CXX11_ABI={ABI}"]
NVCC_FLAGS += [f"-D_GLIBCXX_USE_CXX11_ABI={ABI}"]

if CUDA_HOME is None:
    raise RuntimeError(
        "Cannot find CUDA_HOME. CUDA must be available to build the package.")

def get_nvcc_cuda_version(cuda_dir: str) -> Version:
    """Get the CUDA version from nvcc.

    Adapted from https://github.com/NVIDIA/apex/blob/8b7a1ff183741dd8f9b87e7bafd04cfde99cea28/setup.py
    """
    nvcc_output = subprocess.check_output([cuda_dir + "/bin/nvcc", "-V"],
                                          universal_newlines=True)
    output = nvcc_output.split()
    release_idx = output.index("release") + 1
    nvcc_cuda_version = parse(output[release_idx].split(",")[0])
    return nvcc_cuda_version

def filter_nvcc_flags_for_arch(nvcc_flags, arch_substrings):
    """Filter NVCC flags, only keep gencode flags for specified architectures"""
    filtered_flags = []
    skip_next = False
    for i, flag in enumerate(nvcc_flags):
        if skip_next:
            skip_next = False
            continue
        if flag == "-gencode":
            if i + 1 < len(nvcc_flags):
                arch_flag = nvcc_flags[i + 1]
                if any(sub in arch_flag for sub in arch_substrings):
                    filtered_flags.append(flag)
                    filtered_flags.append(arch_flag)
                skip_next = True
        elif flag not in ["-gencode"]:
            filtered_flags.append(flag)
    return filtered_flags

compute_capabilities = set()
cuda_architectures = os.environ.get("CUDA_ARCHITECTURES")
if cuda_architectures is not None:
    for arch in cuda_architectures.split(","):
        arch = arch.strip()
        if arch:
            compute_capabilities.add(arch)
else:
    #Iterate over all GPUs on the current machine. Also you can modify this part to specify the architecture if you want to build for specific GPU architectures.
    device_count = torch.cuda.device_count()
    for i in range(device_count):
        major, minor = torch.cuda.get_device_capability(i)
        if major < 8:
            warnings.warn(f"skipping GPU {i} with compute capability {major}.{minor}")
            continue
        compute_capabilities.add(f"{major}.{minor}")

if not compute_capabilities:
    raise RuntimeError("No GPUs found. Please specify the target GPU architectures or build on a machine with GPUs.")
else:
    
    unsupported_archs = compute_capabilities - SUPPORTED_ARCHS
    if unsupported_archs:
        warnings.warn(f"Unsupported GPU architectures detected: {unsupported_archs}. Supported architectures: {SUPPORTED_ARCHS}")
        compute_capabilities = compute_capabilities & SUPPORTED_ARCHS
        if not compute_capabilities:
            raise RuntimeError(f"No supported GPU architectures found. Detected: {compute_capabilities | unsupported_archs}, Supported: {SUPPORTED_ARCHS}")
    
    print(f"Detect GPUs with compute capabilities: {compute_capabilities}")

nvcc_cuda_version = get_nvcc_cuda_version(CUDA_HOME)
# Validate the NVCC CUDA version.
if nvcc_cuda_version < Version("12.0"):
    raise RuntimeError("CUDA 12.0 or higher is required to build the package.")
if nvcc_cuda_version < Version("12.4") and any(cc.startswith("8.9") for cc in compute_capabilities):
    raise RuntimeError(
        "CUDA 12.4 or higher is required for compute capability 8.9.")
if nvcc_cuda_version < Version("12.3") and any(cc.startswith("9.0") for cc in compute_capabilities):
    raise RuntimeError(
        "CUDA 12.3 or higher is required for compute capability 9.0.")
if nvcc_cuda_version < Version("12.8") and any(cc.startswith("12.0") for cc in compute_capabilities):
    raise RuntimeError(
        "CUDA 12.8 or higher is required for compute capability 12.0.")

# Add target compute capabilities to NVCC flags.
for capability in compute_capabilities:
    if capability.startswith("8.0"):
        HAS_SM80 = True
        num = "80"
    elif capability.startswith("8.6"):
        HAS_SM86 = True
        num = "86"
    elif capability.startswith("8.9"):
        HAS_SM89 = True
        num = "89"
    elif capability.startswith("9.0"):
        HAS_SM90 = True
        num = "90a" # need to use sm90a instead of sm90 to use wgmma ptx instruction.
    elif capability.startswith("12.0"):
        HAS_SM120 = True
        num = "120" # need to use sm120a to use mxfp8/mxfp4/nvfp4 instructions.
    NVCC_FLAGS += ["-gencode", f"arch=compute_{num},code=sm_{num}"]
    if capability.endswith("+PTX"):
        NVCC_FLAGS += ["-gencode", f"arch=compute_{num},code=compute_{num}"]

ext_modules = []

if HAS_SM80 or HAS_SM86 or HAS_SM89 or HAS_SM90 or HAS_SM120:
    sm80_sources = [
        "csrc/qattn/pybind_sm80.cpp",
        "csrc/qattn/qk_int_sv_f16_cuda_sm80.cu",
    ]

    qattn_extension_sm80 = CUDAExtension(
        name="sageattention._qattn_sm80",
        sources=sm80_sources,
        extra_compile_args={
            "cxx": CXX_FLAGS,
            "nvcc": NVCC_FLAGS,
        },
    )
    ext_modules.append(qattn_extension_sm80)

if HAS_SM89 or HAS_SM120:
    sm89_sources = [
        "csrc/qattn/pybind_sm89.cpp",
        "csrc/qattn/sm89_qk_int8_sv_f8_accum_f32_attn_inst_buf.cu",
        "csrc/qattn/sm89_qk_int8_sv_f8_accum_f16_attn_inst_buf.cu",
        "csrc/qattn/sm89_qk_int8_sv_f8_accum_f32_attn.cu",
        "csrc/qattn/sm89_qk_int8_sv_f8_accum_f32_fuse_v_scale_fuse_v_mean_attn.cu",
        "csrc/qattn/sm89_qk_int8_sv_f8_accum_f32_fuse_v_scale_attn.cu",
        "csrc/qattn/sm89_qk_int8_sv_f8_accum_f32_fuse_v_scale_attn_inst_buf.cu",
        "csrc/qattn/sm89_qk_int8_sv_f8_accum_f16_fuse_v_scale_attn_inst_buf.cu"
        #"csrc/qattn/qk_int_sv_f8_cuda_sm89.cu",
    ]

    arch_substrings = ["sm_89", "compute_89", "sm_90a", "compute_90a", "sm_120", "compute_120"]
    filtered_flags = filter_nvcc_flags_for_arch(NVCC_FLAGS, arch_substrings)

    qattn_extension_sm89 = CUDAExtension(
        name="sageattention._qattn_sm89",
        sources=sm89_sources,
        extra_compile_args={
            "cxx": CXX_FLAGS,
            "nvcc": filtered_flags if filtered_flags else NVCC_FLAGS,
        },
    )
    ext_modules.append(qattn_extension_sm89)

if HAS_SM90:
    sm90_sources = [
        "csrc/qattn/pybind_sm90.cpp",
        "csrc/qattn/qk_int_sv_f8_cuda_sm90.cu",
    ]

    arch_substrings = ["sm_90a", "compute_90a"]
    filtered_flags = filter_nvcc_flags_for_arch(NVCC_FLAGS, arch_substrings)

    qattn_extension_sm90 = CUDAExtension(
        name="sageattention._qattn_sm90",
        sources=sm90_sources,
        extra_compile_args={
            "cxx": CXX_FLAGS,
            "nvcc": filtered_flags if filtered_flags else NVCC_FLAGS,
        },
        extra_link_args=['-lcuda'],
    )
    ext_modules.append(qattn_extension_sm90)

# Fused kernels.
fused_extension = CUDAExtension(
    name="sageattention._fused",
    sources=["csrc/fused/pybind.cpp", "csrc/fused/fused.cu"],
    extra_compile_args={
        "cxx": CXX_FLAGS,
        "nvcc": NVCC_FLAGS,
    },
)
ext_modules.append(fused_extension)


parallel = None
if 'EXT_PARALLEL' in os.environ:
    try:
        parallel = int(os.getenv('EXT_PARALLEL'))
    finally:
        pass


# Prevent file conflicts when multiple extensions are compiled simultaneously
class BuildExtensionSeparateDir(BuildExtension):
    build_extension_patch_lock = threading.Lock()
    thread_ext_name_map = {}

    def finalize_options(self):
        if parallel is not None:
            self.parallel = parallel
        super().finalize_options()

    def build_extension(self, ext):
        with self.build_extension_patch_lock:
            if not getattr(self.compiler, "_compile_separate_output_dir", False):
                compile_orig = self.compiler.compile

                def compile_new(*args, **kwargs):
                    return compile_orig(*args, **{
                        **kwargs,
                        "output_dir": os.path.join(
                            kwargs["output_dir"],
                            self.thread_ext_name_map.get(threading.current_thread().ident,  f"thread_{threading.current_thread().ident}")),
                    })
                self.compiler.compile = compile_new
                self.compiler._compile_separate_output_dir = True
        self.thread_ext_name_map[threading.current_thread().ident] = ext.name

        original_build_temp = self.build_temp
        self.build_temp = os.path.join(original_build_temp, ext.name.replace(".", "_"))
        os.makedirs(self.build_temp, exist_ok=True)

        try:
            objects = super().build_extension(ext)
        finally:
            self.build_temp = original_build_temp

        return objects

setup(
    name='sageattention',
    version='2.2.0',
    author='SageAttention team',
    license='Apache 2.0 License',
    description='Accurate and efficient plug-and-play low-bit attention.',
    long_description=open('README.md', encoding='utf-8').read(),
    long_description_content_type='text/markdown',
    url='https://github.com/thu-ml/SageAttention',
    packages=find_packages(),
    python_requires='>=3.9',
    ext_modules=ext_modules,
    cmdclass={"build_ext": BuildExtensionSeparateDir} if ext_modules else {},
)