Make kernel configurations symetric between ROCm and CUDA for bs 64

Author: sstamenk

csrc/kernels.cu

@@ -1783,7 +1783,7 @@ MAKE_kQuantizeBlockwiseSmall(half, 32, NF4)
 MAKE_kQuantizeBlockwiseSmall(float, 32, NF4)
 MAKE_kQuantizeBlockwiseSmall(bnb_bfloat16, 32, NF4)
 
-// QBLOCK_SIZE=64 instantiations (used on HIP for blocksize=64)
+// QBLOCK_SIZE=64 instantiations (blocksize=64, 4-bit)
 MAKE_kQuantizeBlockwiseSmall(half, 64, FP4)
 MAKE_kQuantizeBlockwiseSmall(float, 64, FP4)
 MAKE_kQuantizeBlockwiseSmall(bnb_bfloat16, 64, FP4)

csrc/ops.cu

@@ -51,7 +51,6 @@ void quantizeBlockwise(
     else if (blocksize == 128)
         kQuantizeBlockwise<T, 128, 2, 0, DATA_TYPE><<<num_blocks, 64>>>(code, A, absmax, out, rand, rand_offset, n);
     else if (blocksize == 64) {
-#if BNB_HIP
         if constexpr (DATA_TYPE > 0) {
             const int ws = bnb_host_warp_size();
             const int num_qb = ws / (64 / 2);
@@ -61,9 +60,6 @@ void quantizeBlockwise(
         } else {
             kQuantizeBlockwise<T, 64, 2, 0, DATA_TYPE><<<num_blocks, 32>>>(code, A, absmax, out, rand, rand_offset, n);
         }
-#else
-        kQuantizeBlockwise<T, 64, 2, 0, DATA_TYPE><<<num_blocks, 32>>>(code, A, absmax, out, rand, rand_offset, n);
-#endif
     } else if (blocksize == 32) {
         if constexpr (DATA_TYPE > 0) {
             const int ws = bnb_host_warp_size();