reactantXLAExec: address an offset argument view as a slice of its base buffer

[wsmoses]

Replaces #3240's host staging. Depends on the read-only-argument change (the PR this branch is stacked on).

A kernel argument may point into the middle of an allocation: MFEM's BlockVector blocks are MakeRef views into one buffer, and everything launched on a block (LinearForm::Assemble on rhs.GetBlock(1), the sub-operator Mults in BlockOperator::Mult) receives base + offset. Instead of copying the tail out to a temporary buffer and back around every launch, the executable takes the base buffer itself. Before the pipeline runs, the argument's type becomes the base buffer's type, every use of it reads stablehlo.slice %base[offset : size], and the value returned for it becomes stablehlo.dynamic_update_slice(%base, %result, offset). The optimizer then folds the update of an unchanged view back to the base argument, so a read-only view is recognized by identity like any other argument (not donated, not returned), and a written view donates its base buffer and returns the updated base, which the runtime stores back into the allocation. The offsets are part of the executable cache key.

No host round trip, the parent buffer keeps aliasing with its blocks, and XLA fuses the slice and the update into the kernel. A view that is not a whole number of elements into a 1-D buffer aborts with a message.

Validation on the MFEM GPU unit-test suite (BlockOperators is the test that exercises this) is running; will update.

🤖 Generated with Claude Code

https://claude.ai/code/session_016zErYp7upmqr4NHfhod9UD