Distributed neighborhood-attention upsampling + consistent CUDA launch checking

[azrael417]

Summary

Adds the upsample direction to DistributedNeighborhoodAttentionS2 and hardens CUDA launch error checking across the attention and DISCO extensions.

1. Distributed upsampling (nlon_out % nlon_in == 0): new forward and backward upsample (scatter) ring-step CUDA kernels, matching the serial upsample attention. K/V rotate around the azimuth
   ring while queries and the softmax state stay local. All three directions — self-attention, downsampling, upsampling — are now supported by the distributed layer. Previously the distributed
   path asserted nlon_in % nlon_out == 0 and upsampling was unsupported.
2. Consistent CUDA launch checking: all attention and DISCO CUDA host wrappers now call C10_CUDA_KERNEL_LAUNCH_CHECK() and explicitly #include <c10/cuda/CUDAException.h> instead of relying on
   transitive includes, so a bad launch surfaces at the launch site rather than at a later, unrelated sync.

Version bumped to 0.9.2c.

Motivation

The serial NeighborhoodAttentionS2 already supported all three resolution directions; the distributed layer only did self-attention and downsampling. Upsampling needs an input-keyed scatter
(each input cell contributes to multiple output cells) rather than the output-keyed gather used by the other directions, which is why it gets its own ring-step kernels and psi convention
(_build_local_psi_upsample) rather than reusing the existing ones.

Changes

• attention/optimized/kernels_cuda/attention_cuda_{fwd,bwd}_ring_upsample.cu: new ring-step upsample kernels (forward; backward pass-1 stats + pass-2 scatter).
• attention_interface.cpp, attention_optimized.py, attention_cuda.cuh: three new ops registered with schemas + register_fake — forward_ring_step_upsample, backward_ring_step_upsample_pass1,
  backward_ring_step_upsample_pass2.
• distributed/distributed_attention.py: upsample ring path and _build_local_psi_upsample (input-keyed local psi, no halo padding at az=polar=1).
• *.cu / *.cuh (attention + DISCO, incl. disco_cuda_{fwd,bwd}.cu and the disco_cuda_fwd_dense_kpacked_sm{90,100}.cu paths): C10_CUDA_KERNEL_LAUNCH_CHECK() + explicit CUDAException.h include.
• Changelog.md, pyproject.toml, init.py: v0.9.2c.

API / numerical behavior

• No breaking API changes. DistributedNeighborhoodAttentionS2 now accepts upsample shapes (nlon_out % nlon_in == 0) that previously raised; existing self/downsample usage is unchanged.
• Distributed output is validated against the serial reference within existing test tolerances.

Testing

CPU (CI equivalent):
python3 -m pytest tests/test_attention.py -x
PT2 contract for the new ops (opcheck, single-rank so no NCCL) — new test:
python3 -m pytest tests/test_attention.py -k ring_upsample_kernels_pt2_compatibility -x
Distributed (multi-GPU), including the new upsample rows (pscale_out = 2/3/4 and odd nlat_in→even nlat_out):
bash tests/run_tests.sh -d --grid_size_lat 2 --grid_size_lon 2

validated across 1x1, 1x2, 2x1, 2x2, 2x4, 4x2

All attention tests pass on CPU/CUDA; distributed attention tests pass across the split configs above, now covering the upsample direction.

Changelog

Changelog.md updated under v0.9.2c with both items.