[XPU] Fix int32 overflow in adaptive_avg_pool2d backward for large tensors

[sgurwinderr]

Summary

adaptive_avg_pool2d backward silently produces an all-zero (unwritten) grad_input for tensors with numel > 2^31 on XPU, because the kernel computes the element count in 32-bit and overflows.

Root cause

In AdaptiveAvgPool2dBwdKernelFunctor (and the SLM variant), the input dims ib_/ic_/ih_/iw_/oh_/ow_ and global_range_ were declared int. numel_ = ib_ * ic_ * ih_ * iw_ is therefore evaluated in 32-bit and stored into the int64_t numel_ after it has already overflowed. For a tensor whose true numel exceeds 2^31, the wrapped product can be negative, so the grad loop

for (int64_t i = gi; i < numel_; i += global_range_)

runs zero iterations and grad_input is never written. (The accessors are already PackedTensorAccessor64, so there is no out-of-bounds — the output is simply left uninitialized/zero.)

Fix

Widen the dimension and range members to int64_t in both backward functors so the numel_ product is computed in 64-bit. Mirrors the fix pattern already used for max_pool3d (#3632). No addressing changes needed (accessors are already 64-bit).

Test plan

Verified on Arc Pro B60 (from-source XPU build) with input (1, 4, 32768, 16385) — numel = 2,147,614,720 (> 2^31):

• Before: numel_ wraps to -2,147,352,576; backward writes 0 / numel elements (grad_input all zeros).
• After: all elements written correctly (fraction written 1.000), matching the CPU reference.

The failing test test_adaptive_avg_pool2d_backward_large_index_offsets_xpu (issue #4327) exercises exactly this large-index path.

Risk

XPU-only; touches only the two adaptive_avg_pool2d backward functors. Small tensors are unaffected (identical arithmetic, now in 64-bit). No API/behavior change beyond correctly writing the full output.

Fixes #4327.