Merge branch 'main' into cpu_fused_kernel

Author: jiqing-feng

.github/scripts/build-rocm.sh

@@ -7,18 +7,18 @@ set -xeuo pipefail
 bnb_rocm_arch="gfx90a;gfx942;gfx1100;gfx1101"
 
 # ROCm 6.4+ - Add gfx1200/gfx1201. Note we assume >=6.4.1.
-[[ "${rocm_version}" == 6.4.* || "${rocm_version}" == 7.*.* ]] && bnb_rocm_arch="${bnb_rocm_arch};gfx1200;gfx1201"
+[[ "${rocm_version}" == 6.4.* || "${rocm_version}" == 7.* ]] && bnb_rocm_arch="${bnb_rocm_arch};gfx1200;gfx1201"
 
 # ROCm 7.0+ - Add gfx950
-[[ "${rocm_version}" == 7.*.* ]] && bnb_rocm_arch="${bnb_rocm_arch};gfx950"
+[[ "${rocm_version}" == 7.* ]] && bnb_rocm_arch="${bnb_rocm_arch};gfx950"
 
 if [ "${build_os:0:6}" == ubuntu ]; then
     image=rocm/dev-ubuntu-22.04:${rocm_version}-complete
     echo "Using image $image"
     docker run --rm --platform "linux/$build_arch" -i \
         -w /src -v "$PWD:/src" "$image" sh -c \
         "apt-get update \
-      && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends cmake \
+      && pip install cmake==3.31.6 \
       && cmake -DCOMPUTE_BACKEND=hip -DBNB_ROCM_ARCH=\"${bnb_rocm_arch}\" . \
       && cmake --build ."
 fi

.github/workflows/python-package.yml

@@ -137,7 +137,7 @@ jobs:
       matrix:
         os: [ubuntu-22.04]
         arch: [x86_64]
-        rocm_version: ["6.2.4", "6.3.4", "6.4.4", "7.0.2"]
+        rocm_version: ["6.2.4", "6.3.4", "6.4.4", "7.0.2", "7.1"]
     runs-on: ${{ matrix.os }}
     steps:
       - uses: actions/checkout@v4

.github/workflows/tests.yml

@@ -54,8 +54,7 @@ jobs:
   build-cuda:
     strategy:
       matrix:
-        # TODO: Add 13.0.1 when we have runners with new enough drivers.
-        cuda_version: ["11.8.0", "12.6.3", "12.8.1", "12.9.1"]
+        cuda_version: ["11.8.0", "12.6.3", "12.8.1", "13.0.1"]
         os: [ubuntu-22.04, ubuntu-22.04-arm]
         include:
           - os: ubuntu-22.04
@@ -349,26 +348,20 @@ jobs:
         os: [ubuntu-22.04, windows-2025]
         arch: [x86_64]
         gpu: [T4, L40S]
-        cuda_version: ["11.8.0", "12.6.3", "12.8.1", "12.9.1"] #, "13.0.1"]
+        cuda_version: ["11.8.0", "12.6.3", "12.8.1", "13.0.1"]
         include:
           - cuda_version: "11.8.0"
             torch_version: "2.3.1"
             pypi_index: "https://download.pytorch.org/whl/cu118"
           - cuda_version: "12.6.3"
-            torch_version: "2.6.0"
+            torch_version: "2.7.1"
             pypi_index: "https://download.pytorch.org/whl/cu126"
-          - cuda_version: "12.9.1"
-            torch_version: "2.8.0"
-            pypi_index: "https://download.pytorch.org/whl/cu129"
           - cuda_version: "12.8.1"
-            torch_version: "2.9.0"
-            pypi_index: "https://download.pytorch.org/whl/test/cu128"
-
-          # Note: Currently our runners do not have new enough drivers for CUDA 13.
-          # Add this when supported.
-          # - cuda_version: "13.0.1"
-          #   torch_version: "2.9.0"
-          #   pypi_index: "https://download.pytorch.org/whl/test/cu130"
+            torch_version: "2.8.0"
+            pypi_index: "https://download.pytorch.org/whl/cu128"
+          - cuda_version: "13.0.1"
+            torch_version: "2.9.1"
+            pypi_index: "https://download.pytorch.org/whl/cu130"
 
 
           # Linux L40S runners

csrc/common.cuh

@@ -2,9 +2,6 @@
 
 // TODO: Let's make some of these constexpr and put in a namespace.
 
-#define BNB_CC_MAXWELL 500
-#define BNB_CC_MAXWELL2 520
-#define BNB_CC_MAXWELL2_X1 530
 #define BNB_CC_PASCAL 600
 #define BNB_CC_PASCAL_X2 620
 #define BNB_CC_VOLTA 700
@@ -17,7 +14,6 @@
 #define BNB_CC_HOPPER 900
 #define BNB_CC_BLACKWELL 1000
 
-#define BNB_FP16_AVAILABLE (__CUDA_ARCH__ >= BNB_CC_MAXWELL2_X1)
 #define BNB_FP16_MMA_AVAILABLE (__CUDA_ARCH__ >= BNB_CC_VOLTA)
 #define BNB_INT8_MMA_AVAILABLE (__CUDA_ARCH__ >= BNB_CC_VOLTA_XAVIER)
 #define BNB_BF16_AVAILABLE (__CUDA_ARCH__ >= BNB_CC_AMPERE)

csrc/kernels.cu

@@ -1767,15 +1767,7 @@ template <typename T, int THREADS, int SPARSE_DECOMP>
 __launch_bounds__(1024, BNB_MAX_THREADS_PER_SM / 1024) __global__
     void kInt8VectorQuant(T* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols) {
 
-    // For sm50/sm52 and CUDA < 12.2 we need to do the reduction in fp32.
-    // Otherwise `T` is `fp16`. This can be removed when Maxwell is dropped.
-#if (__CUDACC_VER_MAJOR__ >= 12 && __CUDACC_VER_MINOR >= 2) || BNB_FP16_AVAILABLE
-    using TReduction = T;
-#else
-    using TReduction = float;
-#endif
-
-    using BlockReduceT = cub::BlockReduce<TReduction, THREADS>;
+    using BlockReduceT = cub::BlockReduce<T, THREADS>;
 
     // One block per row.
     // Threads load column values in a striped arrangement.
@@ -1785,27 +1777,27 @@ __launch_bounds__(1024, BNB_MAX_THREADS_PER_SM / 1024) __global__
     // We then do a blockwise reduction to determine the row's absmax.
 
     __shared__ typename BlockReduceT::TempStorage temp_storage;
-    __shared__ TReduction smem_row_absmax;
+    __shared__ T smem_row_absmax;
 
     const int row_id = blockIdx.x;
     const T* row_data = A + (row_id * cols);
 
     // Threads will read the row values in a striped access pattern and find a local absmax.
-    TReduction row_local_absmax = -FLT_MIN;
+    T row_local_absmax = -FLT_MIN;
     for (int i = threadIdx.x; i < cols; i += THREADS) {
-        const TReduction absval = fabsf(__ldcs(&(row_data[i])));
+        const T absval = fabsf(__ldcs(&(row_data[i])));
 
         // For sparse decomposition, values outside of the threshold are not to be
         // included when calculating the row's absmax.
         if constexpr (SPARSE_DECOMP) {
-            row_local_absmax = fmaxf(row_local_absmax, absval < TReduction(threshold) ? absval : row_local_absmax);
+            row_local_absmax = fmaxf(row_local_absmax, absval < T(threshold) ? absval : row_local_absmax);
         } else {
             row_local_absmax = fmaxf(row_local_absmax, absval);
         }
     }
 
     // Reduce thread-local absmax across the block.
-    const TReduction row_absmax = BlockReduceT(temp_storage).Reduce(row_local_absmax, CUB_REDUCTIONOP_MAX, cols);
+    const T row_absmax = BlockReduceT(temp_storage).Reduce(row_local_absmax, CUB_REDUCTIONOP_MAX, cols);
     if (threadIdx.x == 0) {
         // Save our block's absmax to shared memory for the quantization step.
         rowStats[row_id] = smem_row_absmax = row_absmax;

csrc/kernels.hip

@@ -2613,9 +2613,9 @@ template <typename T, int THREADS, int BITS> __global__ void kgemm_4bit_inferenc
 {
 
   // per threadblock:
-  // load step-by-step in chunks of [BNB_WARP_SIZE,warps]: 1xBNB_WARP_SIZE * [BNB_WARP_SIZE,warps] -> [1,warps]
+  // load step-by-step in chunks of [warp_size,warps]: 1xwarp_size * [warp_size,warps] -> [1,warps]
   // 4 warps -> 4 loads per iter
-  // 1 x BNB_WARP_SIZE * BNB_WARP_SIZE x 4 -> 1x4 outputs per thread block
+  // 1xwarp_size * warp_sizex4 -> 1x4 outputs per thread block
   typedef hipcub::WarpReduce<float, BNB_WARP_SIZE> WarpReduce;
   __shared__ typename WarpReduce::TempStorage temp_storage[THREADS/BNB_WARP_SIZE];
 

docs/source/installation.mdx

@@ -40,12 +40,6 @@ The library can be built using CUDA Toolkit versions as old as **11.8**.
 | NF4/FP4 quantization            | 6.0+            | Pascal (GTX 10X0 series, P100) or newer GPUs|
 
 
-> [!WARNING]
-> Support for Maxwell GPUs is deprecated and will be removed in a future release.
-> Maxwell support is not included in PyPI distributions from `v0.48.0` on and must be built from source.
-> For the best results, a Turing generation device or newer is recommended.
-
-
 ### Installation via PyPI[[cuda-pip]]
 
 This is the most straightforward and recommended installation option.
@@ -81,7 +75,7 @@ For Linux and Windows systems, compiling from source allows you to customize the
 <hfoptions id="source">
 <hfoption id="Linux">
 
-To compile from source, you need CMake >= **3.22.1** and Python >= **3.9** installed. Make sure you have a compiler installed to compile C++ (`gcc`, `make`, headers, etc.). It is recommended to use GCC 9 or newer.
+To compile from source, you need CMake >= **3.22.1** and Python >= **3.10** installed. Make sure you have a compiler installed to compile C++ (`gcc`, `make`, headers, etc.). It is recommended to use GCC 11 or newer.
 
 For example, to install a compiler and CMake on Ubuntu:
 
@@ -133,7 +127,7 @@ The currently distributed `bitsandbytes` packages are built with the following c
 | **OS**             | **oneAPI Toolkit** | **Kernel Implementation** |
 |--------------------|------------------|----------------------|
 | **Linux x86-64**   | 2025.1.3         | SYCL + Triton        |
-| **Windows x86-64** | N/A              | SYCL |
+| **Windows x86-64** | 2025.1.3         | SYCL + Triton        |
 
 The Linux build has a minimum glibc version of 2.34.
 
@@ -197,12 +191,12 @@ pip install -e .
 The currently distributed preview `bitsandbytes` are built with the following configurations:
 
 | **OS**             | **ROCm** | **Targets**
-|--------------------|----------|---------------------------|
-| **Linux x86-64**   | 6.1.2    | gfx90a / gfx942 / gfx1100
-| **Linux x86-64**   | 6.2.4    | gfx90a / gfx942 / gfx1100
-| **Linux x86-64**   | 6.3.4    | gfx90a / gfx942 / gfx1100
-| **Linux x86-64**   | 6.4.4    | gfx90a / gfx942 / gfx1100
-| **Linux x86-64**   | 7.0.0    | gfx90a / gfx942 / gfx1100
+|--------------------|----------|---------------------------------------------------------------------|
+| **Linux x86-64**   | 6.2.4    | CDNA: gfx90a, gfx942 / RDNA: gfx1100, gfx1101
+| **Linux x86-64**   | 6.3.4    | CDNA: gfx90a, gfx942 / RDNA: gfx1100, gfx1101
+| **Linux x86-64**   | 6.4.4    | CDNA: gfx90a, gfx942 / RDNA: gfx1100, gfx1101, gfx1200, gfx1201
+| **Linux x86-64**   | 7.0.2    | CDNA: gfx90a, gfx942, gfx950 / RDNA: gfx1100 / gfx1101 / gfx1200 / gfx1201
+| **Linux x86-64**   | 7.1.0    | CDNA: gfx90a, gfx942, gfx950 / RDNA: gfx1100 / gfx1101 / gfx1200 / gfx1201
 
 **Windows is not currently supported.**
 

tests/test_functional.py

@@ -491,7 +491,7 @@ def test_dim3_igemm(self, seq_dim, hidden_dim, batch_dim):
     @pytest.mark.parametrize("hidden_dim", [32, 1024 * 4], ids=id_formatter("hidden_dim"))
     @pytest.mark.parametrize("batch_dim", [2, 16], ids=id_formatter("batch_dim"))
     @pytest.mark.parametrize("transpose", TRUE_FALSE, ids=id_formatter("transpose"))
-    @pytest.mark.skipif(HIP_ENVIRONMENT, reason="this test is not supported on ROCm yet")
+    @pytest.mark.skipif(ROCM_WARP_SIZE_64, reason="this test is not supported on ROCm yet")
     def test_minmax_igemm(self, seq_dim, hidden_dim, batch_dim, transpose):
         def min_max(x):
             maxA = torch.amax(x, dim=2, keepdim=True)
@@ -1205,7 +1205,7 @@ def test_4bit_compressed_stats(self, device, quant_type, blocksize, dtype):
     @pytest.mark.skipif(not get_available_devices(no_cpu=True), reason="No accelerator device")
     @pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.bfloat16], ids=describe_dtype)
     @pytest.mark.parametrize("quant_type", ["fp4", "nf4"])
-    @pytest.mark.parametrize("blocksize", [64, 128] if not HIP_ENVIRONMENT else [128], ids=id_formatter("blocksize"))
+    @pytest.mark.parametrize("blocksize", [64, 128] if not ROCM_WARP_SIZE_64 else [128], ids=id_formatter("blocksize"))
     def test_4bit_quant_large(self, device, dtype, quant_type, blocksize):
         """
         Test that we can successfully quantize a large tensor. Note that the following limitations apply:
@@ -1428,7 +1428,7 @@ def test_gemv_4bit(self, device, dim, dtype, storage_type, quant_storage, double
     @pytest.mark.parametrize("device", get_available_devices())
     @pytest.mark.parametrize("storage_type", ["nf4", "fp4"], ids=["nf4", "fp4"])
     @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16, torch.float32], ids=describe_dtype)
-    @pytest.mark.skipif(HIP_ENVIRONMENT, reason="this test is not supported on ROCm yet")
+    @pytest.mark.skipif(ROCM_WARP_SIZE_64, reason="this test is not supported on ROCm yet")
     def test_gemv_eye_4bit(self, device, storage_type, dtype):
         if device == "hpu" and not is_supported_on_hpu(storage_type, dtype):
             pytest.skip("This configuration is not supported on HPU.")

tests/test_linear8bitlt.py

@@ -9,7 +9,7 @@
 import torch
 
 import bitsandbytes as bnb
-from bitsandbytes.cextension import HIP_ENVIRONMENT
+from bitsandbytes.cextension import ROCM_WARP_SIZE_64
 from bitsandbytes.nn.modules import Linear8bitLt
 from tests.helpers import (
     TRUE_FALSE,
@@ -234,7 +234,7 @@ def test_linear8bit_serialization(linear8bit):
 @pytest.mark.parametrize("fullgraph", TRUE_FALSE, ids=id_formatter("fullgraph"))
 @pytest.mark.parametrize("mode", ["default", "reduce-overhead"], ids=id_formatter("mode"))
 @pytest.mark.skipif(torch.__version__ < (2, 4), reason="Not supported in torch < 2.4")
-@pytest.mark.skipif(HIP_ENVIRONMENT, reason="this test is not supported on ROCm yet")
+@pytest.mark.skipif(ROCM_WARP_SIZE_64, reason="this test is not supported on ROCm yet")
 def test_linear8bitlt_torch_compile(device, threshold, bias, fullgraph, mode):
     if device == "cuda" and platform.system() == "Windows":
         pytest.skip("Triton is not officially supported on Windows")

tests/test_parametrize.py

@@ -3,7 +3,7 @@
 import torch.nn as nn
 
 from bitsandbytes import functional as F
-from bitsandbytes.cextension import HIP_ENVIRONMENT
+from bitsandbytes.cextension import ROCM_WARP_SIZE_64
 from bitsandbytes.nn.parametrize import (
     Bnb4bitParametrization,
     replace_parameter_4bit,
@@ -39,7 +39,7 @@ def __init__(self, device="cpu", dtype=torch.float32):
 @pytest.mark.parametrize("compress_statistics", TRUE_FALSE, ids=id_formatter("compress_statistics"))
 @pytest.mark.parametrize(
     "blocksize",
-    [64, 128, 256] if not HIP_ENVIRONMENT else [128, 256],
+    [64, 128, 256] if not ROCM_WARP_SIZE_64 else [128, 256],
 )
 def test_replace_parameter_4bit(device, dtype, quant_type, compress_statistics, blocksize):
     """Test basic parameter replacement with 4-bit quantization on different dtypes."""
@@ -267,7 +267,7 @@ def test_quant_state_preservation(device, dtype):
 
     module = ParametrizeTestModule(device=device, dtype=dtype)
 
-    blocksize = 128 if HIP_ENVIRONMENT else 64
+    blocksize = 128 if ROCM_WARP_SIZE_64 else 64
 
     # Apply parametrization with specific settings
     replace_parameter_4bit(module, "weight_2d", quant_type="nf4", compress_statistics=True, blocksize=blocksize)
@@ -326,7 +326,7 @@ def test_multiple_parameters(device, dtype):
 @pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.bfloat16], ids=describe_dtype)
 @pytest.mark.parametrize(
     "blocksize",
-    [64, 128, 256] if not HIP_ENVIRONMENT else [128, 256],
+    [64, 128, 256] if not ROCM_WARP_SIZE_64 else [128, 256],
 )
 def test_different_blocksizes(device, dtype, blocksize):
     """Test parametrization with different block sizes to verify flexibility."""