feat(cuda-device): support u64 reductions and scans in cooperative_groups #1497

What problem does this solve?

warp_reduce, warp_scan, block_reduce, and block_scan do not accept u64 directly. Built-in support would simplify aggregating 64-bit counters with these APIs, without requiring callers to define a scalar wrapper and its operations.

Proposed API / behavior

Implement WarpShuffle for u64 and ReduceOp<u64> for the existing Sum, Min, Max, BitAnd, BitOr, and BitXor operations. Preserve group masks, tile boundaries, and wrapping integer-sum semantics.

The following should compile using the existing API:

use cuda_device::cooperative_groups::{WarpTile, ops::Sum, warp_reduce};

pub fn total(tile: &WarpTile<32>, value: u64) -> u64 {
    warp_reduce::<u64, Sum, 32>(tile, value)
}

It currently produces E0277 for both u64: WarpShuffle and Sum: ReduceOp<u64>. The same bounds prevent using u64 with the other three functions.

Tests should compare GPU results against a CPU oracle, including tiles of 1/2/4/8/16/32 lanes, multiple tiles, high-word values, wrapping overflow, and block operations with partial final warps.

Alternatives considered

A local Copy newtype around u64 could implement WarpShuffle using two 32-bit exchanges through WarpCollective, alongside local reduction operations. This is a viable extension point; the proposal moves that boilerplate and its tests into the library. Such a wrapper was not implemented or GPU-tested during this investigation.

The low-level warp::shuffle_*_u64 functions also allow custom algorithms. Defining only a local ReduceOp<u64> is insufficient: u64: WarpShuffle is still missing, and downstream crates cannot implement that external trait for the primitive type.

Additional context

Validated against upstream commit d141118c05f705fcf932a8e9077dc714477edc4d, using nightly-2026-10-11:

• cargo check accepted all four functions with u32; equivalent u64 probes failed with the bounds above. A local sum operation removed only the ReduceOp<u64> error.
• The existing shuffle_64 example passed its CPU comparisons on an RTX 5060 Ti (sm_120); Compute Sanitizer memcheck and synccheck both reported zero errors. This validates the existing low-level shuffles, not the proposed collective implementations.
• PR #278 supplied the low-level 64-bit shuffles. PR #1305 addresses the u32/i32 redux fast path in the same file, without adding u64 support.