Use conditional load/store algo for warp size compatibility

BLOCK_LOAD_WARP_TRANSPOSE requires threads >= warp_size. On CDNA
(warp=64), kQuantizeBlockwise with BLOCK_SIZE=64 has only 32
threads. Fall back to BLOCK_LOAD_DIRECT / BLOCK_STORE_DIRECT
when threads < BNB_WARP_SIZE. This avoids rocprim compilation
errors while keeping WARP_TRANSPOSE for larger block sizes.

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

Author: TimDettmers

csrc/kernels.cu

@@ -327,13 +327,22 @@ __global__ void kQuantizeBlockwise(
     float local_abs_max = 0.0f;
     int local_rand_idx = 0;
 
-    typedef bnb_cub::BlockLoad<T, BLOCK_SIZE / NUM_PER_TH, NUM_PER_TH, bnb_cub::BLOCK_LOAD_WARP_TRANSPOSE> LoadT;
+    // WARP_TRANSPOSE requires block_dim >= warp_size. On CDNA (warp=64),
+    // block_dim=32 (from BLOCK_SIZE=64/NUM_PER_TH=2) is too small. Fall back
+    // to DIRECT load/store in that case.
+    static constexpr int THREADS = BLOCK_SIZE / NUM_PER_TH;
+    static constexpr auto LOAD_ALGO = (THREADS >= BNB_WARP_SIZE)
+        ? bnb_cub::BLOCK_LOAD_WARP_TRANSPOSE : bnb_cub::BLOCK_LOAD_DIRECT;
+    static constexpr auto STORE_ALGO = (THREADS >= BNB_WARP_SIZE)
+        ? bnb_cub::BLOCK_STORE_WARP_TRANSPOSE : bnb_cub::BLOCK_STORE_DIRECT;
+
+    typedef bnb_cub::BlockLoad<T, THREADS, NUM_PER_TH, LOAD_ALGO> LoadT;
     typedef bnb_cub::BlockStore<
-        unsigned char, BLOCK_SIZE / NUM_PER_TH, (DATA_TYPE > 0) ? NUM_PER_TH / 2 : NUM_PER_TH,
-        bnb_cub::BLOCK_STORE_WARP_TRANSPOSE>
+        unsigned char, THREADS, (DATA_TYPE > 0) ? NUM_PER_TH / 2 : NUM_PER_TH,
+        STORE_ALGO>
         StoreChar;
-    typedef bnb_cub::BlockReduce<float, BLOCK_SIZE / NUM_PER_TH> BlockReduce;
-    typedef bnb_cub::BlockLoad<float, BLOCK_SIZE / NUM_PER_TH, NUM_PER_TH, bnb_cub::BLOCK_LOAD_WARP_TRANSPOSE>
+    typedef bnb_cub::BlockReduce<float, THREADS> BlockReduce;
+    typedef bnb_cub::BlockLoad<float, THREADS, NUM_PER_TH, LOAD_ALGO>
         LoadFloat;
 
     __shared__ typename LoadT::TempStorage loadt;