Merge feature/qutlass-nvfp4-gemm into QLORA-2

Combines NVMe weight streaming pipeline (QLORA-2) with NVFP4
quantization and CUTLASS/hand-written GEMM dispatch (qutlass).

Resolved conflicts in:
- _ops.py: Appended NVFP4 op definitions after kbit ops
- backends/cuda/ops.py: Appended NVFP4 kernel registrations after kbit kernels
- functional.py: Appended NVFP4 quantization functions after kbit functions
- nn/modules.py: Preserved LinearKbit + prepare_model_for_kbit_training,
  then LinearNVFP4 class
- .gitignore: Merged both additions

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

Author: TimDettmers

.github/workflows/test-runner.yml

@@ -26,7 +26,7 @@ on:
       gpu_type:
         type: string
         default: ""
-        description: "GPU type for CUDA testing: T4, L40S"
+        description: "GPU type for CUDA testing: T4, A10,L40S"
       # cpu_type currently only affects linux x64 CPU testing to select specific CPU architectures
       cpu_type:
         type: string
@@ -65,11 +65,14 @@ jobs:
                   T4)
                     TEST_RUNNER="bandb-aws-g4dn-4xlarge-plus-use1-public-80"
                     ;;
+                  A10)
+                    TEST_RUNNER="bandb-aws-g5-4xlarge-plus-use1-public-80"
+                    ;;
                   L40S)
                     TEST_RUNNER="bandb-aws-g6e-4xlarge-plus-use1-public-80"
                     ;;
                   *)
-                    echo "::error::Must specify gpu_type (T4 or L40S) for linux-x64 cuda backend"
+                    echo "::error::Must specify gpu_type (T4, A10, L40S) for linux-x64 cuda backend"
                     exit 1
                     ;;
                 esac
@@ -164,7 +167,7 @@ jobs:
         run: bash .github/scripts/build-cuda.sh
         env:
           cuda_version: ${{ inputs.cuda_version }}
-          cuda_targets: "75;89"
+          cuda_targets: "75;80;89"
 
       - name: Upload build artifact
         uses: actions/upload-artifact@v4

.github/workflows/tests-nightly.yml

@@ -20,7 +20,7 @@ jobs:
         platform: [linux-x64, linux-aarch64, macos, windows]
         # default runners don't have AVX-512 support, but icelake does
         cpu_type: ["", icelake]
-        torch_version: ["2.3.1", "2.8.0", "2.9.1"]
+        torch_version: ["2.3.1", "2.9.1", "2.10.0"]
 
         exclude:
           # aarch64 minimum torch version is 2.5.1
@@ -56,7 +56,7 @@ jobs:
       matrix:
         # Linux x64 cross-product
         platform: [linux-x64]
-        gpu_type: [T4, L40S]
+        gpu_type: [T4, A10, L40S]
         cuda_version: ["11.8.0", "12.6.3", "12.8.1", "13.0.2"]
 
         include:
@@ -65,13 +65,13 @@ jobs:
             torch_version: "2.3.1"
             pypi_index: "https://download.pytorch.org/whl/cu118"
           - cuda_version: "12.6.3"
-            torch_version: "2.7.1"
+            torch_version: "2.8.0"
             pypi_index: "https://download.pytorch.org/whl/cu126"
           - cuda_version: "12.8.1"
-            torch_version: "2.8.0"
+            torch_version: "2.9.1"
             pypi_index: "https://download.pytorch.org/whl/cu128"
           - cuda_version: "13.0.2"
-            torch_version: "2.9.1"
+            torch_version: "2.10.0"
             pypi_index: "https://download.pytorch.org/whl/cu130"
 
           # Windows CUDA Tests - T4 GPU (CUDA 11.8 only, multiple torch versions)

.github/workflows/tests-pr.yml

@@ -31,7 +31,7 @@ jobs:
         platform: [linux-x64, linux-aarch64, macos]
         # default runners don't have AVX-512 support, but icelake does
         cpu_type: ["", icelake]
-        torch_version: ["2.3.1", "2.9.1"]
+        torch_version: ["2.3.1", "2.10.0"]
 
         exclude:
           # aarch64 minimum torch version is 2.5.1
@@ -64,7 +64,7 @@ jobs:
       fail-fast: false
       matrix:
         platform: [linux-x64]
-        gpu_type: [T4, L40S]
+        gpu_type: [T4, A10, L40S]
         cuda_version: ["11.8.0", "12.8.1", "13.0.2"]
 
         include:
@@ -73,10 +73,10 @@ jobs:
             torch_version: "2.3.1"
             pypi_index: "https://download.pytorch.org/whl/cu118"
           - cuda_version: "12.8.1"
-            torch_version: "2.8.0"
+            torch_version: "2.9.1"
             pypi_index: "https://download.pytorch.org/whl/cu128"
           - cuda_version: "13.0.2"
-            torch_version: "2.9.1"
+            torch_version: "2.10.0"
             pypi_index: "https://download.pytorch.org/whl/cu130"
 
           # Windows CUDA test - single configuration

.gitignore

@@ -160,3 +160,4 @@ cuda-spec.md
 cuda-spec-additions.md
 spec.md
 spec_details.md
+agents/*_issues.json

.gitmodules

@@ -0,0 +1,3 @@
+[submodule "third_party/cutlass"]
+	path = third_party/cutlass
+	url = https://github.com/NVIDIA/cutlass.git

CLAUDE.md

@@ -1,25 +1,64 @@
-# Coordinating agent work on GitHub issues
+# MANDATORY: Use git worktrees for all branch work
 
-To analyze open issues, generate prompts, and launch parallel worker agents, follow `agents/coordinator_guide.md`. This uses the GitHub issue tools in `~/git/lab_tools/github/` — see `agents/github_tools_guide.md` for the bitsandbytes-specific reference.
+NEVER work on a fix or feature branch inside the main `~/git/bitsandbytes` checkout. Always create a worktree first:
 
-# Parallel sessions
+```bash
+cd ~/git/bitsandbytes
+git worktree add ~/git/bnb-fix-<NUMBER> -b fix/issue-<NUMBER>
+cd ~/git/bnb-fix-<NUMBER>
+```
+
+This keeps the main checkout clean and allows parallel sessions. If you are already inside a worktree directory, you do not need to create another one.
+
+**Before creating a worktree**, check the worktree registry for existing ones — see the Git Worktrees section in `~/.claude/CLAUDE.md`. Bitsandbytes-specific naming conventions: `agents/worktree_guide.md`. General worktree guide: `~/git/lab_tools/worktree_guide.md`.
+
+# MANDATORY: Check for existing PRs before starting work
 
-To work on multiple branches at once, use git worktrees:
+Before working on any issue, check whether a PR already exists:
 
 ```bash
-git worktree add ../bitsandbytes-<branch-name> -b <branch-name>
-cd ../bitsandbytes-<branch-name>
-claude
+gh pr list --search "issue-number OR keyword" --state open
 ```
 
-Full guide: `agents/worktree_guide.md`
+If a PR exists, review and build on it instead of starting from scratch. Do not create duplicate work.
 
-# Testing
+# MANDATORY: Run linting before every pull request
 
-Run the test suite with 4 parallel workers (optimal for any machine):
+Before pushing a PR branch, you MUST run the full pre-commit suite. CI will reject PRs that fail any check:
 
 ```bash
-pytest tests/ -v --tb=short -n 4
+pre-commit run --all-files
 ```
 
-Best practices, benchmark data, and known architecture-specific issues: `agents/testing_guide.md`
+This runs ruff, ruff format, typos, trailing-whitespace, clang-format, and all other CI lint hooks. Review and commit any changes it makes. Do NOT run only `ruff check` and `ruff format` — those are just 2 of 10 hooks. Full details: `agents/linting_guide.md`
+
+# Testing: only run relevant tests
+
+Do NOT run the full test suite — it takes 10+ minutes. Instead, run only the tests that cover the code you changed:
+
+```bash
+pytest tests/test_relevant_file.py -v --tb=short -k "relevant_test_name"
+```
+
+The full suite will be run separately. Best practices and known issues: `agents/testing_guide.md`
+
+# Agent Dispatch (the "Dispatcher" role)
+
+To triage open GitHub issues, generate prompt files, and launch parallel worker agents, read `agents/dispatch_guide.md`. If told "you're the Dispatcher" or "please read the Dispatch Guide," that's what this refers to. The dispatch workflow uses the GitHub issue tools in `agents/` — see `agents/github_tools_guide.md` for the bitsandbytes-specific reference.
+
+# Issue maintenance and triage
+
+To identify and close stale, duplicate, or resolved issues: `agents/issue_maintenance_guide.md`. Common closeable patterns (old CUDA setup, Windows pre-support, third-party app issues, etc.) are cataloged in `agents/issue_patterns.md`.
+
+# Pull request review
+
+When tasked with reviewing a pull request, you MUST read these guides before starting the review:
+
+1. `agents/pr_review_guide.md` — The complete review workflow (classification, checklists, verdict format, and posting instructions). This is the primary guide; follow its steps sequentially.
+2. `agents/architecture_guide.md` — Codebase architecture and patterns
+3. `agents/code_standards.md` — Code quality expectations
+4. `agents/api_surface.md` — Public API catalog (for detecting breaking changes)
+5. `agents/downstream_integrations.md` — How Transformers, PEFT, Accelerate, TGI, and vLLM depend on bitsandbytes (for assessing downstream impact)
+6. `agents/security_guide.md` — Trust model and security checklist (especially for external contributor PRs)
+
+For CUDA kernel changes, also read `agents/kbit_gemm_context.md`. The PR review guide references all of these at the appropriate steps.

CMakeLists.txt

@@ -10,8 +10,38 @@
 #                        Separate by semicolons, i.e. `-DCOMPUTE_CAPABILITY=89;90;100;120`
 #                        Check your compute capability here: https://developer.nvidia.com/cuda-gpus
 #  - PTXAS_VERBOSE: Pass the `-v` option to the PTX Assembler
+#  - ROCM_VERSION: Override the ROCm version shortcode used in the output library name.
+#                  Useful when PyTorch was built against a different ROCm version than the
+#                  system install. For example, `-DROCM_VERSION=70` produces
+#                  libbitsandbytes_rocm70.so even if the system has ROCm 7.2.
 cmake_minimum_required(VERSION 3.22.1)
 
+# On Windows with HIP backend, auto-detect compilers from ROCM_PATH before project()
+if(WIN32 AND COMPUTE_BACKEND STREQUAL "hip")
+    if(DEFINED ENV{ROCM_PATH})
+        set(ROCM_PATH $ENV{ROCM_PATH})
+    endif()
+    if(ROCM_PATH AND NOT DEFINED CMAKE_CXX_COMPILER)
+        set(CMAKE_CXX_COMPILER "${ROCM_PATH}/lib/llvm/bin/clang++.exe")
+    endif()
+    if(ROCM_PATH AND NOT DEFINED CMAKE_HIP_COMPILER)
+        set(CMAKE_HIP_COMPILER "${ROCM_PATH}/lib/llvm/bin/clang++.exe")
+    endif()
+    # On Windows, the HIP compiler needs explicit paths to find device libraries.
+    if(ROCM_PATH)
+        find_path(ROCM_DEVICE_LIB_PATH
+            NAMES oclc_abi_version_400.bc ocml.bc
+            PATHS "${ROCM_PATH}/amdgcn/bitcode"
+                  "${ROCM_PATH}/lib/llvm/amdgcn/bitcode"
+            NO_DEFAULT_PATH
+        )
+        set(CMAKE_HIP_FLAGS "--rocm-path=${ROCM_PATH}")
+        if(ROCM_DEVICE_LIB_PATH)
+            set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} --rocm-device-lib-path=${ROCM_DEVICE_LIB_PATH}")
+        endif()
+    endif()
+endif()
+
 project(bitsandbytes LANGUAGES CXX)
 
 # If run without specifying a build type, default to using the Release configuration:
@@ -197,20 +227,80 @@ if(BUILD_CUDA)
 
     list(APPEND SRC_FILES ${CUDA_FILES})
 
+    # SM_120a NVFP4 GEMM kernel: requires compute_120a for block-scaled MMA
+    # Only include if 120 or 121 is in the target architectures
+    # Check both COMPUTE_CAPABILITY (may have been popped) and _LATEST_CAPABILITY
+    set(_HAS_SM120 FALSE)
+    foreach(_cap IN LISTS COMPUTE_CAPABILITY)
+        if(_cap MATCHES "^12[01]$")
+            set(_HAS_SM120 TRUE)
+        endif()
+    endforeach()
+    if(_LATEST_CAPABILITY MATCHES "^12[01]$")
+        set(_HAS_SM120 TRUE)
+    endif()
+    if(_HAS_SM120)
+        # Build as separate OBJECT library with its own CUDA_ARCHITECTURES
+        # to avoid conflict with the global architecture settings
+        set(_NVFP4_SM120_SOURCES csrc/kernels_nvfp4_sm120.cu)
+
+        # CUTLASS-based NVFP4 GEMM (requires CUDA 12.8+)
+        if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL "12.8" AND EXISTS "${CMAKE_SOURCE_DIR}/third_party/cutlass/include")
+            list(APPEND _NVFP4_SM120_SOURCES
+                csrc/qutlass/gemm_nvfp4_sm120.cu
+                csrc/qutlass/scale_reorder.cu
+                csrc/qutlass/fused_quantize_nv.cu
+            )
+            set(_HAS_CUTLASS_NVFP4 TRUE)
+            message(STATUS "CUTLASS NVFP4 SM_120a GEMM + fused quantize enabled")
+        else()
+            set(_HAS_CUTLASS_NVFP4 FALSE)
+            message(STATUS "CUTLASS NVFP4 GEMM disabled (needs CUDA >= 12.8 and third_party/cutlass)")
+        endif()
+
+        add_library(nvfp4_sm120a OBJECT ${_NVFP4_SM120_SOURCES})
+        set_target_properties(nvfp4_sm120a PROPERTIES
+            CUDA_ARCHITECTURES "120a"
+            POSITION_INDEPENDENT_CODE ON
+            CUDA_SEPARABLE_COMPILATION OFF
+        )
+        target_compile_options(nvfp4_sm120a PRIVATE
+            $<$<COMPILE_LANGUAGE:CUDA>:--use_fast_math>
+        )
+
+        if(_HAS_CUTLASS_NVFP4)
+            target_include_directories(nvfp4_sm120a PRIVATE
+                "${CMAKE_SOURCE_DIR}/third_party/cutlass/include"
+                "${CMAKE_SOURCE_DIR}/third_party/cutlass/tools/util/include"
+                "${CMAKE_SOURCE_DIR}/csrc/qutlass/include"
+            )
+            target_compile_options(nvfp4_sm120a PRIVATE
+                $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>
+                $<$<COMPILE_LANGUAGE:CUDA>:-std=c++17>
+                $<$<COMPILE_LANGUAGE:CUDA>:-O3>
+                $<$<COMPILE_LANGUAGE:CUDA>:-DNDEBUG>
+                $<$<COMPILE_LANGUAGE:CUDA>:-DQUTLASS_DISABLE_PYBIND>
+            )
+        endif()
+
+        message(STATUS "NVFP4 SM_120a GEMM kernel enabled")
+    endif()
+
     string(APPEND BNB_OUTPUT_NAME "_cuda${CUDA_VERSION_SHORT}")
     add_compile_definitions(BUILD_CUDA)
 elseif(BUILD_HIP)
-    enable_language(HIP)
-    message(STATUS "HIP Compiler: ${CMAKE_HIP_COMPILER}")
+    # Set target architectures before enable_language(HIP), which would otherwise
+    # auto-detect a single GPU and override the defaults.
     if(DEFINED BNB_ROCM_ARCH)
       set(CMAKE_HIP_ARCHITECTURES ${BNB_ROCM_ARCH})
-    else()
-      if (NOT AMDGPU_TARGETS AND NOT CMAKE_HIP_ARCHITECTURES)
-        set(CMAKE_HIP_ARCHITECTURES "gfx90a;gfx942;gfx1100;gfx1101;gfx1150;gfx1151;gfx1200;gfx1201")
-      elseif (AMDGPU_TARGETS AND NOT CMAKE_HIP_ARCHITECTURES)
-        set(CMAKE_HIP_ARCHITECTURES ${AMDGPU_TARGETS})
-      endif()
+    elseif(AMDGPU_TARGETS AND NOT CMAKE_HIP_ARCHITECTURES)
+      set(CMAKE_HIP_ARCHITECTURES ${AMDGPU_TARGETS})
+    elseif(NOT CMAKE_HIP_ARCHITECTURES)
+      set(CMAKE_HIP_ARCHITECTURES "gfx90a;gfx942;gfx1100;gfx1101;gfx1150;gfx1151;gfx1200;gfx1201")
     endif()
+
+    enable_language(HIP)
+    message(STATUS "HIP Compiler: ${CMAKE_HIP_COMPILER}")
     message(STATUS "HIP Targets: ${CMAKE_HIP_ARCHITECTURES}")
 
     list(APPEND SRC_FILES ${HIP_FILES})
@@ -222,7 +312,15 @@ elseif(BUILD_HIP)
     string(REGEX MATCH "[0-9]+\\.[0-9]+" HIP_VERSION "${HIP_CONFIG_VERSION}")
     string(REPLACE "." "" HIP_VERSION_SHORT "${HIP_VERSION}")
 
-    string(APPEND BNB_OUTPUT_NAME "${HIP_VERSION_SHORT}")
+    # Expose a cache variable that the user can set to override the ROCm version in the library name
+    set(ROCM_VERSION "${HIP_VERSION_SHORT}" CACHE STRING "Expected ROCm Version Shortcode")
+
+    message(STATUS "ROCm Version: ${HIP_VERSION_SHORT} (from hipconfig)")
+    if(NOT ROCM_VERSION STREQUAL "${HIP_VERSION_SHORT}")
+        message(WARNING "Overriding ROCm version in library name: ${HIP_VERSION_SHORT} -> ${ROCM_VERSION}")
+    endif()
+
+    string(APPEND BNB_OUTPUT_NAME "${ROCM_VERSION}")
     add_compile_definitions(__HIP_PLATFORM_AMD__)
     add_compile_definitions(__HIP_PLATFORM_HCC__)
     add_compile_definitions(BUILD_HIP)
@@ -263,6 +361,8 @@ endif()
 if(WIN32)
     # Export all symbols
     set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON)
+    # Prevent Windows SDK min/max macros from conflicting with std::min/std::max
+    add_compile_definitions(NOMINMAX)
 endif()
 
 if(MSVC)
@@ -274,6 +374,11 @@ add_library(bitsandbytes SHARED ${SRC_FILES})
 target_compile_features(bitsandbytes PUBLIC cxx_std_17)
 target_include_directories(bitsandbytes PUBLIC csrc)
 
+# Link NVFP4 SM_120a object library if available
+if(TARGET nvfp4_sm120a)
+    target_sources(bitsandbytes PRIVATE $<TARGET_OBJECTS:nvfp4_sm120a>)
+endif()
+
 if (BUILD_CPU)
     if (OpenMP_CXX_FOUND)
         target_link_libraries(bitsandbytes PRIVATE OpenMP::OpenMP_CXX)
@@ -316,10 +421,11 @@ if(BUILD_CUDA)
     )
 endif()
 if(BUILD_HIP)
-    if(NOT DEFINED ENV{ROCM_PATH})
-      set(ROCM_PATH /opt/rocm)
-    else()
+    # Determine ROCM_PATH from environment variable, fallback to /opt/rocm on Linux
+    if(DEFINED ENV{ROCM_PATH})
       set(ROCM_PATH $ENV{ROCM_PATH})
+    else()
+      set(ROCM_PATH /opt/rocm)
     endif()
     list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH})
     macro(find_package_and_print_version PACKAGE_NAME)
@@ -331,14 +437,23 @@ if(BUILD_HIP)
     find_package_and_print_version(hipsparse REQUIRED)
 
     ## hacky way of excluding hip::amdhip64 (with it linked many tests unexpectedly fail e.g. adam8bit because of inaccuracies)
-    set_target_properties(hip::host PROPERTIES INTERFACE_LINK_LIBRARIES "")
-    set_target_properties(hip-lang::host PROPERTIES INTERFACE_LINK_LIBRARIES "")
-    set(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES "")
+    ## On Windows, we need to link amdhip64 explicitly
+    if(NOT WIN32)
+        set_target_properties(hip::host PROPERTIES INTERFACE_LINK_LIBRARIES "")
+        set_target_properties(hip-lang::host PROPERTIES INTERFACE_LINK_LIBRARIES "")
+        set(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES "")
+    endif()
 
     target_include_directories(bitsandbytes PRIVATE ${CMAKE_SOURCE_DIR} ${CMAKE_SOURCE_DIR}/include ${ROCM_PATH}/include /include)
     target_link_directories(bitsandbytes PRIVATE ${ROCM_PATH}/lib /lib)
     target_link_libraries(bitsandbytes PUBLIC roc::hipblas hip::hiprand roc::hipsparse)
 
+    # On Windows, rocblas is not pulled in transitively by roc::hipblas
+    # and is needed because ops_hip.cuh uses rocblas_handle directly.
+    if(WIN32)
+        target_link_libraries(bitsandbytes PUBLIC rocblas)
+    endif()
+
     target_compile_definitions(bitsandbytes PUBLIC BNB_USE_HIP)
     set_source_files_properties(${HIP_FILES} PROPERTIES LANGUAGE HIP)
     set_target_properties(bitsandbytes PROPERTIES LINKER_LANGUAGE CXX)

NOTICE.md

@@ -1 +1,5 @@
 The majority of bitsandbytes is licensed under MIT, however portions of the project are available under separate license terms: PyTorch is licensed under the BSD license.
+
+The NVFP4 GEMM kernel in `csrc/qutlass/` is derived from [QuTLASS](https://github.com/IST-DASLab/qutlass) by Roberto L. Castro (IST Austria), licensed under the Apache License 2.0.
+
+[CUTLASS](https://github.com/NVIDIA/cutlass) by NVIDIA is included as a submodule in `third_party/cutlass/`, licensed under the BSD 3-Clause License.