Cache the ROCm warp size per device instead of for device 0

[IxMxAMAR]

Follow-up to #2059.

That one fixed the compile-time BNB_WARP_SIZE so all of gfx9 gets 64. The runtime side in ops.cu is still off though: bnb_host_warp_size() asks for hipDeviceAttributeWarpSize on device 0 once and caches it for the whole process. The comment directly above it already says that's wrong on mixed RDNA+CDNA systems, so this is just closing that out.

It's not cosmetic, because the two sides have to agree. kQuantizeBlockwiseSmall derives NUM_QB from the compile-time BNB_WARP_SIZE, and ops.cu derives num_qb, the grid and the block size from the runtime value. When those disagree, the kernel's NUM_QB no longer matches how the launch was sized, and everything indexed by qb_id goes with it: warp_reduce[] and smem_absmax_value[] are NUM_QB elements, and the write is absmax[blockIdx.x * NUM_QB + qb_id].

It breaks differently depending on the direction:

• wave64 device, 32-wide compile-time size: the kernel computes NUM_QB = 1 so a block covers 64 values, but the host sizes the grid for 2 quantization blocks per block, so the grid is half of what it needs to be and half the input never gets quantized. qb_id reaches 1 against one-element arrays.
• the other way round: NUM_QB = 2 so a block covers 128 values, but the grid is sized for 1 per block, so it walks twice as far as it should and the absmax writes run past the buffer.

Fix is to query the current device and cache per device, capped at 16. A failed query falls back to the compile-time BNB_WARP_SIZE rather than to 32, because the kernel's NUM_QB is derived from that same constant, so it is the only value guaranteed to agree with the launch. On gfx9 that constant is 64, and falling back to 32 there would have reproduced the mismatch above instead of avoiding it.

Checked on an RX 9070 XT (gfx1201, ROCm 7.2, clang 23, MSVC 18), building this branch on its own:

• -DCOMPUTE_BACKEND=hip -DBNB_ROCM_ARCH=gfx1201 builds clean
• test_ops.py, test_linear4bit.py, test_linear8bitlt.py, test_autograd.py: 1342 passed, 149 skipped, 30 xfailed
• -DBNB_ROCM_ARCH=gfx906 cross-compiles clean. It did not before [ROCm] Restore Wave64 warp size for all gfx9 targets #2059: at 3fe95b5 it fails instantiating BlockLoad<__half, 32, 2, BLOCK_LOAD_WARP_TRANSPOSE> in kQuantizeBlockwise, with rocPRIM's static assertion failed ... 'BlockSize % ::rocprim::arch::wavefront::min_size() == 0': BlockSize must be a multiple of hardware warpsize. That's THREADS=32 picking BLOCK_LOAD_WARP_TRANSPOSE on a 64-wide wave, so the gfx906 build break was real and [ROCm] Restore Wave64 warp size for all gfx9 targets #2059 closed it. This PR just keeps the host side consistent with that.

What I could not check is the mixed wave32/wave64 case itself. I don't have a machine with both, and on a single-device system this returns the same value as before, so it's a no-op there. That's also why there's no new test: reproducing it needs two devices with different wave sizes, and I'd rather not add a test that only runs on hardware neither of us has.

One choice in here worth a second opinion: the 16-device cap is arbitrary. Past the cap the compile-time width is used, which is right for the arch the library was built for but wrong if someone genuinely has 17+ devices with mixed wave sizes, so a mutex-guarded vector is fine too if you'd rather not have a ceiling. That's easy to change.