Fix lint and formatting issues from CI pre-commit checks

Apply ruff lint fix (unused variable), ruff format, and clang-format
to pass CI pre-commit hooks.

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

Author: TimDettmers

bitsandbytes/functional.py

@@ -1031,8 +1031,7 @@ def create_normal_float_codebook(k: int, device=None) -> torch.Tensor:
         from scipy.stats import norm
     except ImportError as ie:
         raise ImportError(
-            "Scipy is required for `create_normal_float_codebook`. "
-            "Install `bitsandbytes` with the `[test]` extra.",
+            "Scipy is required for `create_normal_float_codebook`. Install `bitsandbytes` with the `[test]` extra.",
         ) from ie
 
     if device is None:

csrc/ops.cu

@@ -656,23 +656,22 @@ template void percentileClipping(half* g, float* gnorm_vec, int step, const int
 // ---- Device helpers ----
 
 __device__ __forceinline__ float warp_reduce_absmax_kbit(float val) {
-    #pragma unroll
+#pragma unroll
     for (int offset = 16; offset > 0; offset >>= 1)
         val = fmaxf(val, __shfl_down_sync(0xFFFFFFFF, val, offset));
     return __shfl_sync(0xFFFFFFFF, val, 0);
 }
 
-template <int K>
-__device__ __forceinline__ void pack_kbit_warp(unsigned char qval, unsigned int* packed_words) {
-    #pragma unroll
+template <int K> __device__ __forceinline__ void pack_kbit_warp(unsigned char qval, unsigned int* packed_words) {
+#pragma unroll
     for (int bit = 0; bit < K; bit++)
         packed_words[bit] = __ballot_sync(0xFFFFFFFF, (qval >> bit) & 1);
 }
 
 template <int K>
 __device__ __forceinline__ unsigned char unpack_kbit_warp(const unsigned int* packed_words, int lane_id) {
     unsigned char val = 0;
-    #pragma unroll
+#pragma unroll
     for (int bit = 0; bit < K; bit++)
         val |= ((packed_words[bit] >> lane_id) & 1) << bit;
     return val;
@@ -682,25 +681,24 @@ __device__ __forceinline__ unsigned char unpack_kbit_warp(const unsigned int* pa
 
 template <typename T, int K>
 __global__ void kQuantizeBlockwise_kbit(
-    const float* __restrict__ codebook,
-    const T* __restrict__ A,
-    float* __restrict__ absmax,
-    unsigned int* __restrict__ packed_out,
-    const int n
+    const float* __restrict__ codebook, const T* __restrict__ A, float* __restrict__ absmax,
+    unsigned int* __restrict__ packed_out, const int n
 ) {
     const int warp_id = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
     const int lane_id = threadIdx.x % 32;
     const int block_start = warp_id * 32;
-    if (block_start >= n) return;
+    if (block_start >= n)
+        return;
     float val = (block_start + lane_id < n) ? (float)A[block_start + lane_id] : 0.0f;
     float amax = warp_reduce_absmax_kbit(fabsf(val));
     float amax_safe = fmaxf(amax, 1e-8f);
-    if (lane_id == 0) absmax[warp_id] = amax;
+    if (lane_id == 0)
+        absmax[warp_id] = amax;
     float normalized = val / amax_safe;
     float cb = (lane_id < (1 << K)) ? codebook[lane_id] : 0.0f;
     unsigned char best_idx = 0;
     float best_dist = 1e10f;
-    #pragma unroll
+#pragma unroll
     for (int i = 0; i < (1 << K); i++) {
         float cb_val = __shfl_sync(0xFFFFFFFF, cb, i);
         float dist = fabsf(normalized - cb_val);
@@ -722,7 +720,8 @@ __global__ void kQuantizeBlockwise_kbit(
 constexpr int E4M4_BIAS = 11;
 
 __device__ __forceinline__ float decode_e4m4_absmax(unsigned char raw) {
-    if (raw == 0) return 0.0f;
+    if (raw == 0)
+        return 0.0f;
     int e = raw >> 4;
     int m = raw & 0xF;
     if (e == 0) {
@@ -738,13 +737,11 @@ __device__ __forceinline__ float decode_e4m4_absmax(unsigned char raw) {
 
 // Template helper: convert ABSMAX_T to float.
 // Specialization for unsigned char uses E4M4 decode.
-template <typename ABSMAX_T>
-__device__ __forceinline__ float load_absmax(const ABSMAX_T* absmax, int idx) {
+template <typename ABSMAX_T> __device__ __forceinline__ float load_absmax(const ABSMAX_T* absmax, int idx) {
     return (float)absmax[idx];
 }
 
-template <>
-__device__ __forceinline__ float load_absmax<unsigned char>(const unsigned char* absmax, int idx) {
+template <> __device__ __forceinline__ float load_absmax<unsigned char>(const unsigned char* absmax, int idx) {
     return decode_e4m4_absmax(absmax[idx]);
 }
 
@@ -755,30 +752,29 @@ __device__ __forceinline__ float load_absmax<unsigned char>(const unsigned char*
 // Templated on T (output type) and ABSMAX_T (absmax format).
 template <typename T, int K, int BLOCKS_PER_WARP, typename ABSMAX_T>
 __global__ void kDequantizeBlockwise_kbit_vec(
-    const unsigned int* __restrict__ packed_in,
-    const float* __restrict__ codebook,
-    const ABSMAX_T* __restrict__ absmax,
-    T* __restrict__ out,
-    const int n
+    const unsigned int* __restrict__ packed_in, const float* __restrict__ codebook, const ABSMAX_T* __restrict__ absmax,
+    T* __restrict__ out, const int n
 ) {
     const int warp_id = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
     const int lane_id = threadIdx.x % 32;
     const int base_block = warp_id * BLOCKS_PER_WARP;
 
-    if (base_block * 32 >= n) return;
+    if (base_block * 32 >= n)
+        return;
 
     // Load codebook into lane registers (one-time, amortized across BLOCKS_PER_WARP blocks)
     float cb = (lane_id < (1 << K)) ? codebook[lane_id] : 0.0f;
 
-    #pragma unroll
+#pragma unroll
     for (int b = 0; b < BLOCKS_PER_WARP; b++) {
         const int block_id = base_block + b;
         const int block_start = block_id * 32;
-        if (block_start >= n) break;
+        if (block_start >= n)
+            break;
 
         float amax = load_absmax(absmax, block_id);
         unsigned int packed[K];
-        #pragma unroll
+#pragma unroll
         for (int bit = 0; bit < K; bit++) {
             unsigned int word = (lane_id == bit) ? packed_in[block_id * K + bit] : 0;
             packed[bit] = __shfl_sync(0xFFFFFFFF, word, bit);
@@ -794,14 +790,12 @@ __global__ void kDequantizeBlockwise_kbit_vec(
 // ---- Launch wrappers ----
 
 #define KBIT_WARPS_PER_BLOCK 8
-#define KBIT_THREADS_PER_BLOCK (KBIT_WARPS_PER_BLOCK * 32)  // 256
+#define KBIT_THREADS_PER_BLOCK (KBIT_WARPS_PER_BLOCK * 32) // 256
 
 // ---- Production kernel launchers (Stage 4-5) ----
 
 template <typename T, int K>
-void quantizeBlockwise_kbit(
-    const float* codebook, const T* A, float* absmax, unsigned int* packed_out, int n
-) {
+void quantizeBlockwise_kbit(const float* codebook, const T* A, float* absmax, unsigned int* packed_out, int n) {
     int num_blocks_quant = (n + 31) / 32;
     int num_cuda_blocks = (num_blocks_quant + KBIT_WARPS_PER_BLOCK - 1) / KBIT_WARPS_PER_BLOCK;
     kQuantizeBlockwise_kbit<T, K><<<num_cuda_blocks, KBIT_THREADS_PER_BLOCK>>>(codebook, A, absmax, packed_out, n);
@@ -811,23 +805,21 @@ void quantizeBlockwise_kbit(
 // Generic dequant launcher: supports all output types and absmax formats.
 template <typename T, int K, typename ABSMAX_T>
 void dequantizeBlockwise_kbit(
-    const unsigned int* packed_in, const float* codebook, const ABSMAX_T* absmax,
-    T* out, int n, cudaStream_t stream
+    const unsigned int* packed_in, const float* codebook, const ABSMAX_T* absmax, T* out, int n, cudaStream_t stream
 ) {
-    constexpr int BPW = 4;  // blocks per warp
+    constexpr int BPW = 4; // blocks per warp
     int num_blocks_quant = (n + 31) / 32;
     int num_warps = (num_blocks_quant + BPW - 1) / BPW;
     int num_cuda_blocks = (num_warps + KBIT_WARPS_PER_BLOCK - 1) / KBIT_WARPS_PER_BLOCK;
-    kDequantizeBlockwise_kbit_vec<T, K, BPW, ABSMAX_T><<<num_cuda_blocks, KBIT_THREADS_PER_BLOCK, 0, stream>>>(
-        packed_in, codebook, absmax, out, n);
+    kDequantizeBlockwise_kbit_vec<T, K, BPW, ABSMAX_T>
+        <<<num_cuda_blocks, KBIT_THREADS_PER_BLOCK, 0, stream>>>(packed_in, codebook, absmax, out, n);
     CUDA_CHECK_RETURN(cudaPeekAtLastError());
 }
 
 // ---- Template instantiations ----
 
-#define INSTANTIATE_KBIT_QUANT(T, K) \
-    template void quantizeBlockwise_kbit<T, K>( \
-        const float*, const T*, float*, unsigned int*, int);
+#define INSTANTIATE_KBIT_QUANT(T, K)                                                                                   \
+    template void quantizeBlockwise_kbit<T, K>(const float*, const T*, float*, unsigned int*, int);
 
 INSTANTIATE_KBIT_QUANT(half, 2)
 INSTANTIATE_KBIT_QUANT(half, 3)
@@ -843,9 +835,10 @@ INSTANTIATE_KBIT_QUANT(float, 4)
 INSTANTIATE_KBIT_QUANT(float, 5)
 
 // Dequant instantiations: all output types × absmax types × K values
-#define INSTANTIATE_KBIT_DEQUANT(T, K, ABSMAX_T) \
-    template void dequantizeBlockwise_kbit<T, K, ABSMAX_T>( \
-        const unsigned int*, const float*, const ABSMAX_T*, T*, int, cudaStream_t);
+#define INSTANTIATE_KBIT_DEQUANT(T, K, ABSMAX_T)                                                                       \
+    template void dequantizeBlockwise_kbit<T, K, ABSMAX_T>(                                                            \
+        const unsigned int*, const float*, const ABSMAX_T*, T*, int, cudaStream_t                                      \
+    );
 
 // uint8 E4M4 absmax (default)
 INSTANTIATE_KBIT_DEQUANT(half, 2, unsigned char)

csrc/pythonInterface.cpp

@@ -391,12 +391,16 @@ void gemv_4bit_inference_fp32(
 
 // Forward declarations of ops.cu template functions
 template <typename T, int K> void quantizeBlockwise_kbit(const float*, const T*, float*, unsigned int*, int);
-template <typename T, int K, typename ABSMAX_T> void dequantizeBlockwise_kbit(const unsigned int*, const float*, const ABSMAX_T*, T*, int, cudaStream_t);
+template <typename T, int K, typename ABSMAX_T>
+void dequantizeBlockwise_kbit(const unsigned int*, const float*, const ABSMAX_T*, T*, int, cudaStream_t);
 
 // Unmangled quantize wrappers
-#define MAKE_KBIT_QUANT(tname, T, K) \
-    void quantize_kbit_##tname##_k##K(const float* codebook, const T* A, float* absmax, unsigned int* packed_out, int n) { \
-        quantizeBlockwise_kbit<T, K>(codebook, A, absmax, packed_out, n); }
+#define MAKE_KBIT_QUANT(tname, T, K)                                                                                   \
+    void quantize_kbit_##tname##_k##K(                                                                                 \
+        const float* codebook, const T* A, float* absmax, unsigned int* packed_out, int n                              \
+    ) {                                                                                                                \
+        quantizeBlockwise_kbit<T, K>(codebook, A, absmax, packed_out, n);                                              \
+    }
 
 MAKE_KBIT_QUANT(fp16, half, 2)
 MAKE_KBIT_QUANT(fp16, half, 3)
@@ -412,10 +416,13 @@ MAKE_KBIT_QUANT(fp32, float, 4)
 MAKE_KBIT_QUANT(fp32, float, 5)
 
 // Unmangled dequant wrappers: output type × absmax type × K
-#define MAKE_KBIT_DEQUANT(tname, T, aname, ABSMAX_T, K) \
-    void dequantize_kbit_##tname##_##aname##_k##K(const unsigned int* packed_in, const float* codebook, \
-                                                    const ABSMAX_T* absmax, T* out, int n, cudaStream_t stream) { \
-        dequantizeBlockwise_kbit<T, K, ABSMAX_T>(packed_in, codebook, absmax, out, n, stream); }
+#define MAKE_KBIT_DEQUANT(tname, T, aname, ABSMAX_T, K)                                                                \
+    void dequantize_kbit_##tname##_##aname##_k##K(                                                                     \
+        const unsigned int* packed_in, const float* codebook, const ABSMAX_T* absmax, T* out, int n,                   \
+        cudaStream_t stream                                                                                            \
+    ) {                                                                                                                \
+        dequantizeBlockwise_kbit<T, K, ABSMAX_T>(packed_in, codebook, absmax, out, n, stream);                         \
+    }
 
 // uint8 E4M4 absmax (default) - all output types
 MAKE_KBIT_DEQUANT(fp16, half, u8abs, unsigned char, 2)
@@ -958,10 +965,12 @@ bool has_avx512bf16_cpu() { return has_avx512bf16(); }
 #if BUILD_CUDA || BUILD_HIP
 
 // Production kernels (Stage 4-5) - quantize only
-#define MAKE_CKBIT(tname, T, K) \
-    void cquantize_kbit_##tname##_k##K(const float* codebook, const T* A, float* absmax, \
-                                        unsigned int* packed_out, int n) { \
-        quantize_kbit_##tname##_k##K(codebook, A, absmax, packed_out, n); }
+#define MAKE_CKBIT(tname, T, K)                                                                                        \
+    void cquantize_kbit_##tname##_k##K(                                                                                \
+        const float* codebook, const T* A, float* absmax, unsigned int* packed_out, int n                              \
+    ) {                                                                                                                \
+        quantize_kbit_##tname##_k##K(codebook, A, absmax, packed_out, n);                                              \
+    }
 
 MAKE_CKBIT(fp16, half, 2)
 MAKE_CKBIT(fp16, half, 3)
@@ -977,10 +986,13 @@ MAKE_CKBIT(fp32, float, 4)
 MAKE_CKBIT(fp32, float, 5)
 
 // Dequant extern C wrappers: output type × absmax type × K
-#define MAKE_CKBIT_DEQUANT(tname, T, aname, ABSMAX_T, K) \
-    void cdequantize_kbit_##tname##_##aname##_k##K(const unsigned int* packed_in, const float* codebook, \
-                                                     const ABSMAX_T* absmax, T* out, int n, cudaStream_t stream) { \
-        dequantize_kbit_##tname##_##aname##_k##K(packed_in, codebook, absmax, out, n, stream); }
+#define MAKE_CKBIT_DEQUANT(tname, T, aname, ABSMAX_T, K)                                                               \
+    void cdequantize_kbit_##tname##_##aname##_k##K(                                                                    \
+        const unsigned int* packed_in, const float* codebook, const ABSMAX_T* absmax, T* out, int n,                   \
+        cudaStream_t stream                                                                                            \
+    ) {                                                                                                                \
+        dequantize_kbit_##tname##_##aname##_k##K(packed_in, codebook, absmax, out, n, stream);                         \
+    }
 
 // uint8 E4M4 absmax - all output types
 MAKE_CKBIT_DEQUANT(fp16, half, u8abs, unsigned char, 2)