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

@@ -48,18 +48,18 @@ def _setup_ctypes(names, argtypes, restype=None):
     restype=ct.c_int32,
 )
 
+# int8 mm dequant: (A, row_stats, col_stats, out, bias, numRows, numCols, stream)
 # int8 mm dequant: (A, row_stats, col_stats, out, bias, numRows, numCols, stream)
 _setup_ctypes(
-    ["cdequant_mm_int32_fp16"],
+    ["cdequant_mm_int32_fp16", "cdequant_mm_int32_bf16"],
     [ct.c_void_p] * 5 + [ct.c_int32, ct.c_int32, ct.c_void_p],
 )
 
 # 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")]
@@ -187,14 +187,18 @@ def _(
 
     # Note: cuda kernel only currently supports fp16 output.
     # We'll later cast to desired dtype if needed.
-    out = torch.empty_like(A, dtype=torch.float16)
+    out_dtype = dtype or torch.float16
+    if out_dtype not in (torch.float16, torch.bfloat16):
+        raise ValueError(f"dtype must be float16 or bfloat16, got {out_dtype}")
+
+    out = torch.empty_like(A, dtype=out_dtype)
 
-    # Note: fused bias in the kernel is only supported for fp16
-    # TODO(matthewdouglas): Consider supporting bf16 fused bias
-    bias_ptr = bias.data_ptr() if bias is not None and bias.dtype == torch.float16 else None
+    # Fuse bias in the kernel when it already matches the output dtype; otherwise add it afterward.
+    bias_ptr = bias.data_ptr() if bias is not None and bias.dtype == out_dtype else None
+    fn = lib.cdequant_mm_int32_bf16 if out_dtype == torch.bfloat16 else lib.cdequant_mm_int32_fp16
 
     with _cuda_device_of(A):
-        lib.cdequant_mm_int32_fp16(
+        fn(
             A.data_ptr(),
             row_stats.data_ptr(),
             col_stats.data_ptr(),
@@ -206,16 +210,16 @@ def _(
         )
 
     # Add bias separately if not fused in kernel
-    if bias is not None and bias.dtype != torch.float16:
+    if bias is not None and bias.dtype != out_dtype:
         out.add_(bias)
 
-    return out.to(dtype or torch.float16)
+    return out
 
 
 @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 +241,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(),

bitsandbytes/functional.py

@@ -1561,6 +1561,7 @@ def int8_mm_dequant(
     col_stats: torch.Tensor,
     out: Optional[torch.Tensor] = None,
     bias: Optional[torch.Tensor] = None,
+    dtype: Optional[torch.dtype] = None,
 ):
     """Performs dequantization on the result of a quantized int8 matrix multiplication.
 
@@ -1570,11 +1571,15 @@ def int8_mm_dequant(
         col_stats (`torch.Tensor`): The column-wise quantization statistics for the rhs operand of the matrix multiplication.
         out (`torch.Tensor`, *optional*): A pre-allocated tensor to store the output of the operation.
         bias (`torch.Tensor`, *optional*): An optional bias vector to add to the result.
+        dtype (`torch.dtype`, *optional*): The desired output dtype. Supports `torch.float16` and
+            `torch.bfloat16`. Defaults to `torch.float16`.
 
     Returns:
-        `torch.Tensor`: The dequantized result with an optional bias, with dtype `torch.float16`.
+        `torch.Tensor`: The dequantized result with an optional bias, with the requested dtype.
     """
-    result = torch.ops.bitsandbytes.int8_mm_dequant.default(A, row_stats, col_stats, dtype=torch.float16, bias=bias)
+    result = torch.ops.bitsandbytes.int8_mm_dequant.default(
+        A, row_stats, col_stats, dtype=dtype or torch.float16, bias=bias
+    )
 
     # TODO(matthewdouglas): Deprecate out kwarg
     if out is not None:

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,21 +1376,26 @@ 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>
+template <typename T, int ITEMS_PER_THREAD, int THREADS>
 __global__ void kdequant_mm_int32_fp16(
-    int* __restrict__ const A, float* __restrict__ const rowStats, float* __restrict__ const colStats, half* out,
-    half* __restrict__ const bias, const int numRows, const int numCols, const int n
+    int* __restrict__ const A, float* __restrict__ const rowStats, float* __restrict__ const colStats, T* out,
+    T* __restrict__ const bias, const int numRows, const int numCols, const int n
 ) {
     const int n_out = numRows * numCols;
 
     int block_offset = blockIdx.x * THREADS * ITEMS_PER_THREAD;
     int thread_offset = threadIdx.x * ITEMS_PER_THREAD;
 
     int local_values[ITEMS_PER_THREAD];
-    half local_output[ITEMS_PER_THREAD];
+    T local_output[ITEMS_PER_THREAD];
 
     float local_rowStats[ITEMS_PER_THREAD];
     float local_colStats[ITEMS_PER_THREAD];
@@ -1415,7 +1414,7 @@ __global__ void kdequant_mm_int32_fp16(
 
         local_colStats[j] = col_idx >= numCols ? 0.0f : __ldg(&colStats[col_idx]);
         local_rowStats[j] = row_idx >= numRows ? 0.0f : __ldg(&rowStats[row_idx]);
-        local_biasValue[j] = ((bias == nullptr) || col_idx >= numCols) ? 0.0f : __half2float(bias[col_idx]);
+        local_biasValue[j] = ((bias == nullptr) || col_idx >= numCols) ? 0.0f : (float)(bias[col_idx]);
     }
 
     // Each block loads THREADS * ITEMS_PER_THREAD values from A
@@ -1425,9 +1424,8 @@ __global__ void kdequant_mm_int32_fp16(
 
 #pragma unroll ITEMS_PER_THREAD
     for (int j = 0; j < ITEMS_PER_THREAD; ++j) {
-        local_output[j] = __float2half(
-            fmaf(local_values[j] * local_rowStats[j] * local_colStats[j], MM_DEQUANT_CONST, local_biasValue[j])
-        );
+        local_output[j] =
+            (T)(fmaf(local_values[j] * local_rowStats[j] * local_colStats[j], MM_DEQUANT_CONST, local_biasValue[j]));
     }
 
 #pragma unroll ITEMS_PER_THREAD
@@ -1596,10 +1594,14 @@ template __global__ void kgemm_4bit_inference_naive<float, 128, 32>(
     float* out, int lda, int ldb, int ldc, int blocksize
 );
 
-template __global__ void kdequant_mm_int32_fp16<4, 512>(
+template __global__ void kdequant_mm_int32_fp16<half, 4, 512>(
     int* __restrict__ const A, float* __restrict__ const rowStats, float* __restrict__ const colStats, half* out,
     half* __restrict__ const bias, const int numRows, const int numCols, const int n
 );
+template __global__ void kdequant_mm_int32_fp16<bnb_bfloat16, 4, 512>(
+    int* __restrict__ const A, float* __restrict__ const rowStats, float* __restrict__ const colStats,
+    bnb_bfloat16* out, bnb_bfloat16* __restrict__ const bias, const int numRows, const int numCols, const int n
+);
 
 template __device__ unsigned char dQuantize<0>(float* smem_code, const float rand, float x);
 template __device__ unsigned char dQuantize<1>(float* smem_code, const float rand, float x);

csrc/kernels.cuh

@@ -65,10 +65,10 @@ __global__ void kOptimizerStatic8bit1StateBlockwise(
     const float gnorm_scale, const bool skip_zeros, const int n
 );
 
-template <int ITEMS_PER_THREAD, int THREADS>
+template <typename T, int ITEMS_PER_THREAD, int THREADS>
 __global__ void kdequant_mm_int32_fp16(
-    int* __restrict__ const A, float* __restrict__ const rowStats, float* __restrict__ const colStats, half* out,
-    half* __restrict__ const bias, const int numRows, const int numCols, const int n
+    int* __restrict__ const A, float* __restrict__ const rowStats, float* __restrict__ const colStats, T* out,
+    T* __restrict__ const bias, const int numRows, const int numCols, const int n
 );
 
 template <typename T, int THREADS, int SPARSE_DECOMP>

csrc/ops.cu

@@ -407,27 +407,29 @@ int fill_up_to_nearest_multiple(int value, int multiple) {
     return value + (value % multiple == 0 ? 0 : (multiple - (value % multiple)));
 }
 
+template <typename T>
 void dequant_mm_int32_fp16(
-    int* A, float* rowStats, float* colStats, half* out, half* bias, int numRows, int numCols, bnb_stream_t stream
+    int* A, float* rowStats, float* colStats, T* out, T* bias, int numRows, int numCols, bnb_stream_t stream
 ) {
     const int threads = 512;
     const int num_per_thread = 4;
     const int num_per_block = threads * num_per_thread;
     const int n = numRows * numCols;
     const int num_blocks = (n + num_per_block - 1) / num_per_block;
 
-    kdequant_mm_int32_fp16<num_per_thread, threads>
+    kdequant_mm_int32_fp16<T, num_per_thread, threads>
         <<<num_blocks, threads, 0, stream>>>(A, rowStats, colStats, out, bias, numRows, numCols, n);
     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 +483,21 @@ 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 void dequant_mm_int32_fp16<half>(
+    int* A, float* rowStats, float* colStats, half* out, half* bias, int numRows, int numCols, bnb_stream_t stream
+);
+template void dequant_mm_int32_fp16<bnb_bfloat16>(
+    int* A, float* rowStats, float* colStats, bnb_bfloat16* out, bnb_bfloat16* bias, int numRows, int numCols,
+    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

@@ -133,11 +133,15 @@ int igemmlt(
 void cutlass_igemm(
     bool transposeA, bool transposeB, int m, int n, int k, void* A, void* B, void* C, int lda, int ldb, int ldc
 );
+
+template <typename T>
 void dequant_mm_int32_fp16(
-    int* A, float* rowStats, float* colStats, half* out, half* bias, int numRows, int numCols, bnb_stream_t stream
+    int* A, float* rowStats, float* colStats, T* out, T* 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

@@ -551,13 +551,27 @@ int cigemmlt_8_rowscale(
 void cdequant_mm_int32_fp16(
     int* A, float* rowStats, float* colStats, half* out, half* bias, int numRows, int numCols, cudaStream_t stream
 ) {
-    dequant_mm_int32_fp16(A, rowStats, colStats, out, bias, numRows, numCols, stream);
+    dequant_mm_int32_fp16<half>(A, rowStats, colStats, out, bias, numRows, numCols, stream);
+}
+
+void cdequant_mm_int32_bf16(
+    int* A, float* rowStats, float* colStats, __nv_bfloat16* out, __nv_bfloat16* bias, int numRows, int numCols,
+    cudaStream_t stream
+) {
+    dequant_mm_int32_fp16<__nv_bfloat16>(A, rowStats, colStats, out, bias, numRows, numCols, stream);
 }
 
 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) {