1:1 from upstream_ref/sgemm_cuda/10_kernel_warptiling.cuh. 3 changes: WARPSIZE 32→64, float→__half, FP32 accumulator. Launch config (confirmed by probe_warp64.sh): NUM_THREADS=128, 2 warps of 64 BM=128 BN=128 BK=16 WM=64 WN=128 WNITER=4 TM=4 TN=4 WMITER=2, WSUBM=32, WSUBN=32, threads_per_warp=64 ✓