`DeviceReduce::ArgMinMax`: ptxas miscompiles the generic warp-shuffle reduction for 12-byte accumulators with a sub-4-byte value field (surfaced by #10731)

(written by claude)

While investigating a flaky CI failure in the test added by #10731, I traced the root cause to an **nvcc/ptxas miscompilation** of CUB's generic warp-shuffle reduction (`WarpReduceShfl::ReduceStep`), *not* a logic bug in CUB or in the reduction operator.

The miscompile corrupts the tie-break of the reduction operator when **all** of the following hold:
1. the accumulator has a **sub-4-byte value field** (here `argminmax_accum_t` = 12 bytes, shuffled as three `uint` words via `reinterpret_cast`);
2. the reduction op uses a **custom comparator** whose magnitude is not representable in the value type — `cuda::uabs(INT16_MIN) == 32768` overflows signed `short`;
3. the value bit pattern `0x8000` flows through `__shfl_down_sync`.

The result: on ties, the reduction returns the wrong index (the tie-break is silently dropped).

The `ArgMinMax` path that #10731 introduced accumulates in the *input* type (`short`), which produces the 12-byte accumulator that triggers the miscompile. This issue tracks the underlying compiler bug, which affects any CUB warp-shuffle reduction meeting the conditions above.

## Environment

- CUDA Toolkit **13.3** (`nvcc`/`ptxas` `V13.3.73`, `cuda_13.3.r13.3/compiler.38244171_0`)
- Reproduced on **sm_86** and **sm_120**
- Original failing CI job: [`CUB nvcc GCC / [CTK13.3 GCC15 C++20] HostLaunch(amd64, RTXPRO6000)`](https://github.com/NVIDIA/cccl/actions/runs/32198759311/job/95925774661?pr=10731) (passed on re-run — see below)

## Reproduction via the CUB test

Against the `ArgMinMax` code as introduced by #10731:

```
# build the HostLaunch C++20 test
ninja -C build/cub-lid0-cpp20 cub.test.device.reduce_arg_minmax.lid_0

# short/int type combo (types_212 - 1), abs_less_t section, last-max
./bin/cub.test.device.reduce_arg_minmax.lid_0 '* - 1' --rng-seed 27
```

Fails at `catch2_test_device_reduce_arg_minmax.cu:338` (the `abs_less_t` section):

```
CATCH_REQUIRE( exp_max_index == d_max_index[0] )
with expansion:
1113075 (0x10fbf3) == 397727
with messages:
c2h::type_name() := "short"
c2h::type_name() := "int"
num_items := 1764559 (0x1aeccf)
```

It is seed-dependent (the input must contain `-32768` landing across a warp boundary), which is why the original CI run passed on re-run. Other failing seeds include `328`. The plain-`less` "default" section always passes; only the `abs_less_t` section fails.

## Minimal standalone reproducer (no CUB / no libcudacxx)

```cpp
// nvcc -std=c++20 -arch=sm_120 repro.cu -o repro && ./repro -> WRONG (max_index != 2)
// nvcc -std=c++20 -arch=sm_120 -Xptxas -O0 repro.cu -o repro && ./repro -> OK (max_index == 2)
#include
#include

struct accum_t { short min_value, max_value; int min_index, max_index; };

__device__ __forceinline__ unsigned short uabs16(short v) { // == cuda::uabs, no UB
return (v < 0) ? static_cast(~static_cast(v) + 1)
: static_cast(v);
}
__device__ __forceinline__ accum_t op(accum_t a, accum_t b) { // "last max" on |value|
accum_t r = a;
bool less_ab = uabs16(a.max_value) < uabs16(b.max_value);
bool less_ba = uabs16(b.max_value) < uabs16(a.max_value);
if (less_ab || (!less_ba && b.max_index > a.max_index)) { r.max_value = b.max_value; r.max_index = b.max_index; }
return r;
}
__device__ __forceinline__ accum_t shuffle_down(accum_t in, int off, unsigned m) { // CUB-style 3-word shuffle
accum_t out; auto* o = reinterpret_cast(&out); auto* i = reinterpret_cast(&in);
#pragma unroll
for (int w = 0; w < 3; ++w) o[w] = __shfl_down_sync(m, i[w], off);
return out;
}
__global__ void kernel(accum_t* out) {
int lane = threadIdx.x;
accum_t v{-32768, -32768, lane, lane}; // all three lanes: value -32768
accum_t acc = v;
for (int off = 1; off <= 2; off <<= 1) { // guarded butterfly over lanes 0,1,2
accum_t t = shuffle_down(acc, off, 0xffffffffu);
if (off + lane <= 2) acc = op(acc, t);
}
if (lane == 0) *out = acc;
}
int main() {
accum_t *d, h; cudaMalloc(&d, sizeof h);
kernel<<<1, 32>>>(d);
cudaMemcpy(&h, d, sizeof h, cudaMemcpyDeviceToHost); cudaDeviceSynchronize();
printf("max_index=%d (expected 2) %s\n", h.max_index, h.max_index == 2 ? "OK" : "*** WRONG ***");
return h.max_index == 2;
}
```

Observed:

```
default -O3: max_index=1 (expected 2) *** WRONG ***
-Xptxas -O0: max_index=2 (expected 2) OK
```

(The exact wrong index depends on the reduction shape; the point is it is not `2`.)

## Analysis

- It is a **codegen** bug: correct with `-G`, `-Xptxas -O0`, or `-Xcicc -O0`; wrong with `-Xptxas -O1`+. Since `-Xptxas -O0` alone (keeping `cicc -O3`) fixes it, **ptxas** is miscompiling the cicc-optimized PTX.
- **Not UB**: the comparator / `cuda::uabs` / `cuda::neg` are well-defined for `INT16_MIN` (the complement is computed on the unsigned value; all casts are in range). `clang`/`gcc -fsanitize=undefined` over all `short` values plus the op are clean, and the UB-free standalone above still reproduces.
- The SASS shows the abs comparison compiled **inconsistently between shuffle steps**: the offset-1 step sign-extends the 16-bit field (`PRMT …0x9910`) and tests it **signed** (`ISETP.LT Rx, RZ` = "value < 0"), while the offset-2 step zero-masks (`LOP3 …0xffff`) and compares **unsigned** (`ISETP.GE.U32`). For bit pattern `0x8000` the sign bit is set, so the two steps disagree and the tie-break is lost.

## Impact

Any CUB warp-shuffle reduction (`WarpReduce`/`BlockReduce` with `BLOCK_REDUCE_WARP_REDUCTIONS`) over a struct with a sub-4-byte field, driven by an operator that compares a value whose transformed magnitude leaves the type's signed range, can silently produce wrong results at default optimization.

Refs: #10731