perf: Reduce register pressure (N_TILES=2, 16 warps/block)

NCU showed 80 regs/thread from N_TILES_PER_WARP=4 accumulators.
Reduce to 2 N-tiles per warp (8 acc floats instead of 16),
increase to 4×4=16 warps per block (m64×n64 tile, 512 threads).
Target 2 blocks/SM for better occupancy.

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

Author: TimDettmers

csrc/kernels_nvfp4_sm120.cu

@@ -86,14 +86,15 @@ __device__ __forceinline__ uint32_t pack_8_nibbles_slow(const unsigned char* dat
 //   Shared memory: A tile (512 bytes) + SFA (64 bytes) per K-step
 // ============================================================================
 
-// N-tiles per warp: each warp computes m16 x (N_TILE_PER_WARP * 8)
-#define N_TILES_PER_WARP 4
+// N-tiles per warp: each warp computes m16 x (N_TILES_PER_WARP * 8)
+#define N_TILES_PER_WARP 2
 // Block config: M_WARPS x N_WARPS warps per block
 // M_WARPS groups along M (each m16), N_WARPS groups along N (each handles N_TILES_PER_WARP n8-tiles)
-#define M_WARPS 2
+#define M_WARPS 4
 #define N_WARPS 4
-#define WARPS_PER_BLOCK (M_WARPS * N_WARPS) // 8
+#define WARPS_PER_BLOCK (M_WARPS * N_WARPS) // 16
 
+// 512 threads, target 2 blocks/SM for good occupancy
 __global__ __launch_bounds__(WARPS_PER_BLOCK * 32, 2) void kGemmNVFP4_opt(
     const unsigned char* __restrict__ A,   // M x K/2 packed FP4 (row-major)
     const unsigned char* __restrict__ B,   // N x K/2 packed FP4 (B transposed, row-major)