Eager race detector crashes on tl.max / tl.min of a loaded tile #506

The eager race detector from the race-detector port (#504, also on #505) crashes when a kernel reduces a loaded tile with tl.max / tl.min and uses the result as an atomic's value. #361's detector crashes the same way; the eager sanitizer on main is not affected.

Repro (Triton 3.8, CPU)

import torch, triton, triton.language as tl
import tilelens
from tilelens.clients import RaceDetector

@triton.jit
def k(x, h, B: tl.constexpr):
    v = tl.load(x + tl.arange(0, B))
    tl.atomic_max(h + 4, tl.max(v))

x = torch.arange(8, dtype=torch.int32)
h = torch.zeros(8, dtype=torch.int32)
tilelens.trace(client=RaceDetector())(k)[(1,)](x, h, B=8)

  File ".../tilelens/clients/symbolic_engine.py", line 1783, in _reduce_max
    return reduce(lambda a, b: If(a >= b, a, b), arr), constraints
TypeError: reduce() arg 2 must support iteration

Plain TRITON_INTERPRET=1 runs the kernel fine (h[4] == 7).

Cause

ReduceSymbolicExpr._reduce_max / _reduce_min (and _reduce_xor_sum) fold self.input._to_z3() as a list of lane values. The race detector installs a load-value provider, so a block load lowers to one Z3 expression over the tile's symbolic lane variable rather than a list, and reduce() gets a non-iterable. The eager sanitizer does not reach this path because it concretizes data-dependent loads first.

tl.sum of the same tile used in an address is already refused cleanly ("data-dependent memory address through sum is unsupported"). The value path of tl.max / tl.min (here, an atomic's value) has no such guard.

Expected

The launch is reported unsupported (the reduction of a symbolic tile is not modelled), or the reduction is modelled; the detector never raises out of the kernel launch.