fix(cuda-device): select tcgen05 alloc warp by linear thread id

[YeonwooSung]

Summary

TmemGuard::alloc_by and dealloc_by selected the warp that runs tcgen05.alloc / tcgen05.dealloc with warp::warp_id(), which is threadIdx.x / 32 and is documented as valid only for 1D blocks. In a 2D block every row whose threadIdx.x / 32 matched the target id executed the warp-collective allocation.

Those two methods now use a block-linear warp id:

linear_tid = threadIdx.x + blockDim.x * (threadIdx.y + blockDim.y * threadIdx.z)
warp       = linear_tid / 32

alloc() / alloc_by(0) is still the warp that contains linear tids 0..31. In a 1D block that is the same as threadIdx.x / 32. In a 32x4 block, warp 0 is only the y = 0 row.

warp::warp_id() is unchanged and still documented as 1D-only. The pure block_linear_warp_id helper is what the host tests cover, because threadIdx_* is not available on the host.

Test plan

• Host test written first; cargo test -p cuda-device --lib warp::tests::block_linear_warp_id failed to compile because block_linear_warp_id did not exist.
• After the implementation, cargo test -p cuda-device passed (106 lib tests, including 1D 256-thread warps, 2D 32x4 one row per warp with exactly 32 threads each, and a 3D 8x4x2 partition, plus integration and compile-fail tests).

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.