Add BF16 support for int8_vectorwise_quant / LLM.int8 activation quant

bitsandbytes/autograd/_functions.py

@@ -124,20 +124,16 @@ def forward(
 
         input_shape = A.shape
 
-        # Cast A to fp16
-        if A.dtype != torch.float16 and not _is_compiling():
-            logger.warning("MatMul8bitLt: inputs will be cast from %s to float16 during quantization", A.dtype)
-
         if len(A.shape) == 3:
             A = A.reshape(-1, A.shape[-1])
 
         # 1. Quantize A. Note that as a side-effect, outliers are suppressed in CA/CAt.
         if ctx.needs_input_grad[1]:
             # Slower path
-            CA, CAt, SCA, SCAt, outlier_cols = F.int8_double_quant(A.to(torch.float16), threshold=state.threshold)
+            CA, CAt, SCA, SCAt, outlier_cols = F.int8_double_quant(A, threshold=state.threshold)
         else:
             # Fast path
-            CA, SCA, outlier_cols = F.int8_vectorwise_quant(A.to(torch.float16), threshold=state.threshold)
+            CA, SCA, outlier_cols = F.int8_vectorwise_quant(A, threshold=state.threshold)
             CAt = SCAt = None
 
         has_grad = False
@@ -152,7 +148,7 @@ def forward(
                 state.reset_grads()
 
                 # 2. Quantize B
-                state.CB, state.SCB, _ = F.int8_vectorwise_quant(B.to(torch.float16))
+                state.CB, state.SCB, _ = F.int8_vectorwise_quant(B)
 
         # Handle sparse decomposition
         if state.threshold > 0.0:
@@ -258,7 +254,7 @@ def forward(ctx, A, B, out=None, bias=None, state=MatmulLtState):
 
             if (state.is_training and not has_grad) or state.CB is None or state.SCB is None:
                 state.reset_grads()
-                state.CB, state.SCB, _ = F.int8_vectorwise_quant(B.to(torch.float16))
+                state.CB, state.SCB, _ = F.int8_vectorwise_quant(B)
                 B = state.CB
 
         CB = state.CB.data.to(A.dtype).mul_(state.SCB.unsqueeze(1).mul(1.0 / 127.0))

bitsandbytes/backends/cuda/ops.py

@@ -56,10 +56,9 @@ def _setup_ctypes(names, argtypes, restype=None):
 
 # int8 vectorwise quant: (A, out, row_stats, threshold, rows, cols, stream)
 _setup_ctypes(
-    ["cint8_vector_quant"],
+    ["cint8_vector_quant", "cint8_vector_quant_bf16"],
     [ct.c_void_p] * 3 + [ct.c_float, ct.c_int32, ct.c_int32, ct.c_void_p],
 )
-
 # 4-bit/8-bit blockwise quantize: (code, A, absmax, out, blocksize, n)
 _setup_ctypes(
     [f"cquantize_blockwise_{d}_{q}" for d in ("fp32", "bf16", "fp16") for q in ("nf4", "fp4")]
@@ -214,8 +213,8 @@ def _(
 
 @register_kernel("bitsandbytes::int8_vectorwise_quant", "cuda")
 def _(A: torch.Tensor, threshold=0.0):
-    if A.dtype != torch.float16:
-        raise ValueError(f"A must be float16, got {A.dtype}")
+    if A.dtype not in (torch.float16, torch.bfloat16):
+        raise ValueError(f"A must be float16 or bfloat16, got {A.dtype}")
     if threshold < 0.0:
         raise ValueError("threshold must be non-negative")
 
@@ -237,8 +236,9 @@ def _(A: torch.Tensor, threshold=0.0):
             # Needed for torch.compile support.
             outlier_cols = torch.empty(0, device=A.device, dtype=torch.int64)
 
+    fn = lib.cint8_vector_quant_bf16 if A.dtype == torch.bfloat16 else lib.cint8_vector_quant
     with _cuda_device_of(A):
-        lib.cint8_vector_quant(
+        fn(
             A.data_ptr(),
             out_row.data_ptr(),
             row_stats.data_ptr(),

csrc/kernels.cu

@@ -1323,53 +1323,47 @@ __launch_bounds__(256, 3) __global__ void kOptimizerStatic8bit1StateBlockwise(
 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) {
-
-    using BlockReduceT = bnb_cub::BlockReduce<T, THREADS>;
-
     // One block per row.
     // Threads load column values in a striped arrangement.
     // e.g. t0 reads row[0], row[0+nthreads], ..
     // and  t1 reads row[1], row[1+nthreads], ..
     // Each thread will determine its local absmax.
     // We then do a blockwise reduction to determine the row's absmax.
-
+    using BlockReduceT = bnb_cub::BlockReduce<float, THREADS>;
     __shared__ typename BlockReduceT::TempStorage temp_storage;
-    __shared__ T smem_row_absmax;
+    __shared__ float 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.
-    T row_local_absmax = -FLT_MIN;
+    float row_local_absmax = -FLT_MIN;
     for (int i = threadIdx.x; i < cols; i += THREADS) {
-        const T absval = fabsf(__ldcs(&(row_data[i])));
-
+        const float absval = fabsf((float)__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 < T(threshold) ? absval : row_local_absmax);
+            row_local_absmax = fmaxf(row_local_absmax, absval < threshold ? absval : row_local_absmax);
         } else {
             row_local_absmax = fmaxf(row_local_absmax, absval);
         }
     }
 
     // Reduce thread-local absmax across the block.
-    const T row_absmax = BlockReduceT(temp_storage).Reduce(row_local_absmax, BNB_MAX_OP, cols);
+    const float row_absmax = BlockReduceT(temp_storage).Reduce(row_local_absmax, BNB_MAX_OP, 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;
     }
     __syncthreads();
 
     // Quantize row-wise.
     const float scale = __fdividef(127.0f, smem_row_absmax);
     for (int i = threadIdx.x; i < cols; i += THREADS) {
-        float val = row_data[i];
-
+        float val = (float)row_data[i];
         if constexpr (SPARSE_DECOMP) {
             // For sparse decomposition, we do not want to quantize the outliers.
             // Instead they're zeroed out.
-            out[row_id * cols + i] = fabs(val) < threshold ? __float2int_rn(val * scale) : 0;
+            out[row_id * cols + i] = fabsf(val) < threshold ? __float2int_rn(val * scale) : 0;
         } else {
             out[row_id * cols + i] = __float2int_rn(val * scale);
         }
@@ -1382,7 +1376,12 @@ template __global__ void kInt8VectorQuant<half, 1024, 0>(
 template __global__ void kInt8VectorQuant<half, 1024, 1>(
     half* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols
 );
-
+template __global__ void kInt8VectorQuant<bnb_bfloat16, 1024, 0>(
+    bnb_bfloat16* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols
+);
+template __global__ void kInt8VectorQuant<bnb_bfloat16, 1024, 1>(
+    bnb_bfloat16* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols
+);
 #define MM_DEQUANT_CONST 6.200012e-05f // 1.0f/(127.0f*127.0f)
 
 template <int ITEMS_PER_THREAD, int THREADS>

csrc/ops.cu

@@ -421,13 +421,14 @@ void dequant_mm_int32_fp16(
     BNB_CHECK_RETURN(BNB_PEEK_LAST_ERROR());
 }
 
+template <typename T>
 void int8VectorQuant(
-    half* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols, bnb_stream_t stream
+    T* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols, bnb_stream_t stream
 ) {
     if (threshold == 0.0) {
-        kInt8VectorQuant<half, 1024, 0><<<rows, 1024, 0, stream>>>(A, out, rowStats, threshold, rows, cols);
+        kInt8VectorQuant<T, 1024, 0><<<rows, 1024, 0, stream>>>(A, out, rowStats, threshold, rows, cols);
     } else {
-        kInt8VectorQuant<half, 1024, 1><<<rows, 1024, 0, stream>>>(A, out, rowStats, threshold, rows, cols);
+        kInt8VectorQuant<T, 1024, 1><<<rows, 1024, 0, stream>>>(A, out, rowStats, threshold, rows, cols);
     }
     BNB_CHECK_RETURN(BNB_PEEK_LAST_ERROR());
 }
@@ -481,6 +482,13 @@ template void gemm_4bit_inference_naive<float, 32>(
     int ldc, int blocksize, bnb_stream_t stream
 );
 
+template void int8VectorQuant<half>(
+    half* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols, bnb_stream_t stream
+);
+template void int8VectorQuant<bnb_bfloat16>(
+    bnb_bfloat16* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols, bnb_stream_t stream
+);
+
 template int igemmlt<32, 0>(
     bnb_blasLt_handle_t ltHandle, int m, int n, int k, const int8_t* A, const int8_t* B, void* C, float* row_scale,
     int lda, int ldb, int ldc, bnb_stream_t stream

csrc/ops.cuh

@@ -136,8 +136,9 @@ void cutlass_igemm(
 void dequant_mm_int32_fp16(
     int* A, float* rowStats, float* colStats, half* out, half* bias, int numRows, int numCols, bnb_stream_t stream
 );
+template <typename T>
 void int8VectorQuant(
-    half* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols, bnb_stream_t stream
+    T* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols, bnb_stream_t stream
 );
 
 template <typename T, int BITS>

csrc/pythonInterface.cpp

@@ -557,7 +557,14 @@ void cdequant_mm_int32_fp16(
 void cint8_vector_quant(
     half* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols, cudaStream_t stream
 ) {
-    int8VectorQuant(A, out, rowStats, threshold, rows, cols, stream);
+    int8VectorQuant<half>(A, out, rowStats, threshold, rows, cols, stream);
+}
+
+void cint8_vector_quant_bf16(
+    __nv_bfloat16* __restrict__ A, int8_t* out, float* rowStats, float threshold, int rows, int cols,
+    cudaStream_t stream
+) {
+    int8VectorQuant<__nv_bfloat16>(A, out, rowStats, threshold, rows, cols, stream);
 }
 
 void* cget_managed_ptr(size_t bytes) {

tests/test_autograd.py

@@ -76,7 +76,7 @@ def test_matmullt(
                 if not transpose[0] and not transpose[1]:
                     B2 = B2.t().contiguous()
 
-                state.CB, state.SCB, _ = bnb.functional.int8_vectorwise_quant(B2.to(torch.float16))
+                state.CB, state.SCB, _ = bnb.functional.int8_vectorwise_quant(B2)
                 B2 = state.CB
 
             if not transpose[0] and transpose[1]:

tests/test_functional.py

@@ -397,22 +397,20 @@ def test_int8_linear_matmul(self, device, dim1, dim2, dim3, dim4, dims, ldb):
     @pytest.mark.parametrize("dim3", [32], ids=id_formatter("dim3"))
     @pytest.mark.parametrize("dim4", [32], ids=id_formatter("dim4"))
     @pytest.mark.parametrize("dims", (2,), ids=id_formatter("dims"))
-    def test_int8_linear_matmul_half(self, device, dim1, dim2, dim3, dim4, dims):
+    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16], ids=describe_dtype)
+    def test_int8_linear_matmul_half(self, device, dim1, dim2, dim3, dim4, dims, dtype):
         for i in range(k):
             if dims == 2:
-                A = torch.normal(0, 0.5, size=(dim1, dim3), device=device).half()
+                A = torch.normal(0, 0.5, size=(dim1, dim3), device=device).to(dtype)
             elif dims == 3:
-                A = torch.normal(0, 0.5, size=(dim1, dim2, dim3), device=device).half()
-            B = torch.randn((dim4, dim3), device=device).half()
+                A = torch.normal(0, 0.5, size=(dim1, dim2, dim3), device=device).to(dtype)
+            B = torch.randn((dim4, dim3), device=device).to(dtype)
             torch.nn.init.xavier_uniform_(B)
             C1 = torch.matmul(A, B.t())
-
             A = A.view(-1, A.shape[-1])
-
             CA, statsA, _ = F.int8_vectorwise_quant(A)
             CB, statsB, _ = F.int8_vectorwise_quant(B)
             output = F.int8_mm_dequant(F.int8_linear_matmul(CA, CB), statsA, statsB)
-
             torch.testing.assert_close(C1.view(-1, C1.shape[-1]), output, atol=0.025, rtol=0.05)
 
     @pytest.mark.parametrize("device", get_available_devices())

tests/test_linear8bitlt.py

@@ -13,6 +13,7 @@
 from bitsandbytes.nn.modules import Linear8bitLt
 from tests.helpers import (
     TRUE_FALSE,
+    describe_dtype,
     get_available_devices,
     id_formatter,
     torch_load_from_buffer,
@@ -364,3 +365,32 @@ def test_linear8bitlt_device_movement(device):
 
     # Accelerator outputs should match both times.
     torch.testing.assert_close(out_accelerator_2, out_accelerator, rtol=1e-8, atol=1e-8)
+
+
+@pytest.mark.parametrize("device", get_available_devices(no_cpu=True))
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16], ids=describe_dtype)
+@pytest.mark.parametrize("threshold", [0.0, 6.0], ids=id_formatter("threshold"))
+def test_linear8bitlt_forward_dtypes(device, dtype, threshold):
+    # Exercises the activation-quant path end-to-end through Linear8bitLt for both fp16 and bf16.
+    linear = torch.nn.Linear(32, 96)
+    linear_custom = Linear8bitLt(
+        linear.in_features,
+        linear.out_features,
+        linear.bias is not None,
+        has_fp16_weights=False,
+        threshold=threshold,
+    )
+    linear_custom.weight = bnb.nn.Int8Params(
+        linear.weight.data.clone(),
+        requires_grad=False,
+        has_fp16_weights=False,
+    )
+    linear_custom.bias = linear.bias
+    linear_custom = linear_custom.to(device)
+
+    x = torch.randn(4, 32, dtype=dtype, device=device)
+    out = linear_custom(x)
+
+    assert out.dtype == dtype
+    assert out.shape == (4, 96)
+    assert out.isfinite().all()

tests/test_modules.py

@@ -7,7 +7,7 @@
 from torch import nn
 
 import bitsandbytes as bnb
-from tests.helpers import get_available_devices, id_formatter, is_supported_on_hpu
+from tests.helpers import describe_dtype, get_available_devices, id_formatter, is_supported_on_hpu
 
 
 @contextlib.contextmanager
@@ -61,15 +61,15 @@ def assert_all_approx_close(a, b, atol=1e-8, rtol=1e-5, count=10):
 
 
 @pytest.mark.parametrize("device", get_available_devices())
+@pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16], ids=describe_dtype)
 @pytest.mark.parametrize("threshold", [0.0, 3.0], ids=id_formatter("threshold"))
-def test_linear8bitlt_inference(device, threshold):
-    l1 = bnb.nn.Linear8bitLt(32, 64, threshold=threshold, has_fp16_weights=False).to(device).half()
+def test_linear8bitlt_inference(device, dtype, threshold):
+    l1 = bnb.nn.Linear8bitLt(32, 64, threshold=threshold, has_fp16_weights=False).to(device).to(dtype)
     assert l1.weight.device.type == device
     assert l1.weight.dtype == torch.int8
-
     l1.eval()
     for i in range(100):
-        b1 = torch.randn(16, 8, 32, device=device).half()
+        b1 = torch.randn(16, 8, 32, device=device, dtype=dtype)
         o1 = l1(b1)
         if i == 1:
             assert l1.state.CB is not None

tests/test_ops.py

@@ -36,29 +36,26 @@ def test_int8_linear_matmul_out(self, device):
 
         opcheck(torch.ops.bitsandbytes.int8_linear_matmul.out, (A, B, out))
 
+    @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16], ids=describe_dtype)
     @pytest.mark.parametrize("threshold", [0.0, 6.0])
     @pytest.mark.parametrize("device", get_available_devices())
-    def test_int8_vectorwise_quant(self, threshold, device):
-        A = torch.randn(10, 20, dtype=torch.float16, device=device)
+    def test_int8_vectorwise_quant(self, dtype, threshold, device):
+        A = torch.randn(10, 20, dtype=dtype, device=device)
         A[1][0] = 1000.0
-
         out_row, row_stats, outlier_cols = torch.ops.bitsandbytes.int8_vectorwise_quant(A, threshold=threshold)
-
         assert out_row.shape == (10, 20)
         assert out_row.dtype == torch.int8
         assert out_row.device == A.device
         assert row_stats.shape == (10,)
         assert row_stats.dtype == torch.float32
         assert row_stats.device == A.device
-
         if threshold > 0.0:
             assert outlier_cols is not None
             assert outlier_cols.dim() == 1
             assert outlier_cols.shape[0] <= A.shape[1]
             assert outlier_cols.device == A.device
         else:
             assert outlier_cols is None
-
         opcheck(torch.ops.bitsandbytes.int8_vectorwise_quant, (A,))
         opcheck(torch.ops.bitsandbytes.int8_vectorwise_quant, (A, threshold))