Clamp the runtime start in slice and slice_update so the window stays in bounds#3917
Open
kapellirohith wants to merge 1 commit into
Open
Clamp the runtime start in slice and slice_update so the window stays in bounds#3917kapellirohith wants to merge 1 commit into
kapellirohith wants to merge 1 commit into
Conversation
slice and slice_update take the start position as an array. That start is turned into a flat offset of sum(start[i] * strides[axes[i]]) and applied straight to the data pointer, and nothing range checks it along the way: normalize_dynamic_slice_inputs validates the count, ndim, dtype and axes of the start but never its value, and neither compute_dynamic_offset nor the dynamic copy kernel bounds check the offset they produce. An out of range start therefore reads and writes outside the array, silently on the GPU backends and with a fault on the CPU. It reaches slice, slice_update, both vjps and the compiled path. Clamp each start index to [0, operand_dim - window_dim] where the offset is computed, which is what XLA's dynamic-slice does. The bounds go in next to the strides and axes that are already passed as constants, so this costs no extra kernel launches. In range starts are unaffected. The CPU dispatch also read signed start dtypes through unsigned pointers, so an int32 start of -1 arrived as 4294967295. Read each dtype through its own pointer type, and do the comparison in the index's own signedness so that neither a negative signed start nor a large unsigned one clamps to the wrong end of the axis.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
mx.sliceandmx.slice_updateaccept the start position as an array (theruntime form). That start becomes a flat offset
which is applied straight to the data pointer. Nothing range-checks it.
normalize_dynamic_slice_inputs(mlx/ops.cpp:744) validates the count, thendim, the dtype and the axes of
start— but never its value — andneither
compute_dynamic_offsetnor the dynamic copy kernel bounds-checks theoffset it produces:
// mlx/backend/metal/kernels/copy.h:218-229 dst[dst_idx + dst_offset] = src[src_idx + src_offset];An out-of-range start therefore reads and writes outside the array. It reaches
five entry points —
slice,slice_update, both gradients, and the compiledpath — and all three backends carry the same unchecked arithmetic.
This PR clamps the start to
[0, operand_dim - window_dim]in the CPU, Metaland CUDA offset paths, matching XLA's
dynamic-slice/dynamic-update-slice,which is the operation this API mirrors. It also fixes a signedness bug in the
CPU offset path (§4) found while validating the clamp.
In-range starts are unaffected.
1. Reproducer
Deterministic, three lines:
On
973e27f8this terminates the process with SIGBUS, 3 runs out of 3.The write side, on Metal, silent — the stored value is nowhere in the
destination, so the store went outside it:
No exception is raised, and Metal shader validation
(
MTL_DEBUG_LAYER=1 MTL_SHADER_VALIDATION=1) reports nothing, because the offsetstill lands inside the process's own Metal heap — just in a different tensor.
Whether an out-of-range access faults or silently returns data depends on what
happens to be mapped at the computed address. Both outcomes are observed below
(80 faults, all on CPU; Metal never faulted in 311 cases). I mention this so the
crash is not read as the defining symptom — the silent case is the more common
one and the more dangerous.
Illustrative only. The store can be observed landing in another allocation:
The specific numbers — start
19537, index712— are allocator-dependent andwill not reproduce byte-for-byte elsewhere. Treat this as a demonstration of
what the unclamped offset can reach, not as a test. The reproducible evidence is
the SIGBUS above, the "store lands outside
dst" probe, and the audit in §2.2. Coverage audit
I did not want to rely on a pass count, so the surface is enumerated as 311
independent cases per backend, each run in its own process so a fault cannot
abort the run. The oracle needs no knowledge of the implementation: a dynamic
slice with start
Smust equal the same slice taken atclamp(S, 0, dim - window), computed in NumPy — which is exactly the XLA rule.Cases cross: entry point × geometry × start kind × start dtype × operand layout.
slice,slice_update, gradient w.r.t. operand, gradient w.r.t. update,mx.compiledint8/16/32/64,uint8/16/32/64Result —
before.txtandafter.txthold the raw per-case output:973e27f8Failures by entry point (both backends):
slice92,slice_update92,gradient-w.r.t.-operand 56, gradient-w.r.t.-update 56, compiled 56.
Failures by start kind: one-past 114, −1 98, +100 70, −100 70, and in-range,
zero and max-edge: 0. The in-range cases are the no-regression controls and
pass on both builds.
Two things this audit turned up that I would otherwise have missed:
DynamicSlice::vjpandDynamicSliceUpdate::vjp(mlx/primitives.cpp:5202,:5278) rebuild aslice/slice_updatewith the same start, so the gradient took the sameout-of-bounds offset as the forward pass. Both gradient entry points fail on
pristine and pass with the patch, and the patch keeps forward and backward
clamping identically, so the gradient still lands on the window the forward
pass used.
constant cotangent that gradient is all ones wherever the window lands, so it
passed on pristine. It only became load-bearing after weighting the cotangent
by position. Noting it because a reviewer should not read "grad is covered" as
automatically meaning "grad is tested".
Not affected, checked and excluded:
slice_update_add/prod/max/minhave noarray-start overload (
mlx/ops.cpp— only three functions takeconst array& start), andvmapover the start indices already throws(
mlx/primitives.cpp:5184,:5241).3. In-tree regression test
test_dynamic_slicing_start_is_clampedinpython/tests/test_ops.py, next totest_dynamic_slicing. It names both streams explicitly rather than relyingon the ambient device, because
MLXTestCaseonly switches device through theDEVICEenvironment variable — a default run would otherwise have exercisedthe GPU only, while the faults and the signedness fix are on the CPU side. The
gpu leg is guarded by
mx.metal.is_available() or mx.cuda.is_available(), so itcovers whichever GPU backend is built (see §11).
It covers, per backend: out-of-range and negative starts on
slice; partialaxes; a window covering the whole axis (only valid start is 0); all eight start
dtypes, asserting signed-negative clamps to 0 and large-unsigned clamps to the
far end;
slice_updatelanding at the clamped position; a neighbouringallocation left bit-identical; and the gradient landing on the clamped window.
Before / after (
intree_before.txt,intree_after.txt):(signal 10) inside the CPU section. Because that kills the runner, the Metal
section never executes, so I ran it standalone on pristine as well — 5 of 5
Metal assertions fail (returns data outside
x; the update lands nowhere indst).4. Second defect: signed starts read through unsigned pointers
Independent of the clamp, and visible by inspection. The CPU offset dispatch
funnels signed and unsigned dtypes through the same unsigned pointer type
(
mlx/backend/cpu/primitives.cpp:56-72on973e27f8):An
int32start of-1is read as4294967295, so the offset becomes4294967295 * striderather than-stride. That is why the CPU faults on[-1, -1]instead of reading just before the buffer. It also means a clamp aloneis insufficient: without this fix a negative start clamps to the far end of the
axis instead of to zero. I hit exactly that in the first iteration of this patch,
and the dtype cases in the audit and the in-tree test now pin it.
Metal is unaffected — its generated kernel is typed with the real dtype.
5. Open question for maintainers: negative starts
The static form treats a negative start as counting from the end; this patch
clamps a negative runtime start to zero. Both are defensible and it is a one-line
change either way, so I would rather you pick than have me choose silently.
Measured on
973e27f8,x = arange(16).reshape(4, 4):x[-1:, -1:][[15]](counts from the end)x[-100:, :]x[99:, :](0,4)x[2:100, 2:100][[10,11],[14,15]]y[-1:, -1:] = 9yy[99:, :] = 9ySo the same logical operation is already safe when the start is a Python value
and unsafe when it is an array, and nothing in the docstrings marks the runtime
form as weaker.
dynamic-slice, so code ported from JAX behaves identically. Diverges from thestatic form's Python semantics.
start + dim), then clamp. Consistent with thestatic form and with Python indexing. Diverges from XLA.
Either is a strict improvement over the status quo, and neither changes in-range
behaviour. Say the word and I will switch it and update the tests.
6. Root cause
973e27f8)mlx/ops.cppnormalize_dynamic_slice_inputschecks count / ndim / dtype / axes ofstart, never its valuemlx/backend/metal/slicing.cppacc += indices[i] * strides[axes[i]]— offset built uncheckedmlx/backend/cpu/primitives.cppmlx/backend/cuda/slicing.cppmlx/backend/cpu/primitives.cppmlx/backend/metal/kernels/copy.hdst[dst_idx + dst_offset] = src[src_idx + src_offset]— offset applied with no bounds check7. The fix
Clamp each start index to
[0, operand_dim - window_dim]insidecompute_dynamic_offset, wherestridesandaxesare already passed asconstants — so there are no extra kernel launches on the hot path.
mlx/backend/metal/slicing.cpp— clamp in the generated offset kernel; thecomparison is emitted per dtype in the index's own signedness.
mlx/backend/cpu/primitives.cpp— same clamp, plus signed dtypes are now readthrough signed pointers (§4).
mlx/backend/cuda/slicing.cpp— the same clamp in the CUDA offset kernel.See §10 for what is and is not validated locally there.
mlx/backend/gpu/slicing.h,mlx/backend/gpu/primitives.cpp— pass theper-axis bounds:
in.shape(ax) - out.shape(ax)forDynamicSlice,out.shape(ax) - upd.shape(ax)forDynamicSliceUpdate.8. Why clamp rather than raise
The obvious alternative is to reject an out-of-range start with an exception.
That is not free here.
The start is a runtime device value. It is not known on the host at graph
construction time. Verified from the code: nothing in
mlx/ops.cppreads aninput array's values while building an op, and in the dynamic-slice path the
start is consumed only on the device (
mlx/backend/metal/slicing.cpplaunches akernel over the
indicesbuffer; the host never dereferences it). Raising fromslice/slice_updatewould require evaluating and reading backstartonevery call — a device→host synchronisation on the exact op whose purpose is to
avoid one. The runtime form exists so a KV-cache append can be issued without
syncing on the current sequence length.
A device-side check would report in the wrong place. Metal offers no
exception mechanism, and MLX has no kernel-authored error channel: the only error
path in the Metal backend is
Device::error_(mlx/backend/metal/device.h:122),which carries
MTLCommandBufferStatusErrorfrom a completed command buffer andis surfaced at synchronize. A kernel could write a flag to a buffer, but the host
would observe it only at the next synchronisation — detached from the call that
caused it, unless you sync eagerly, which reintroduces the cost above. I am not
claiming a device-side error is impossible, only that it is a larger change and
still pays a sync to be timely.
The status quo is not a loud error. Today an out-of-range start is a silent
out-of-bounds access or a process fault — worse than either clamping or throwing.
If you would rather have a validating error despite the sync, the natural place is
an opt-in debug check rather than the default path, and I am happy to add it.
9. Validation
test_dynamic_slicing_start_is_clamped(new)python/tests)DEVICE=cpu(mirrors the CI cpu leg)DEVICE=gpuDEVICE=cpuslice_updateKV-cache append ([1,8,4096,128]fp16, 300 iters)Perf is unchanged within run-to-run spread: the clamp is a few integer operations
inside the existing single-thread offset kernel plus one small constant buffer.
Everything in this table is CPU and Metal. The CUDA hunk is covered by §11.
10. Notes
so the CUDA hunk is not compiled or run locally. See §11.
Metal a dynamic
slice_updatecan read a stale start across a command-bufferboundary. I reproduced that on
973e27f8(17 of 20 fresh processes) andconfirmed the offset is exactly the recycled buffer's prior contents — sweeping
a poison value
Pmoved the write to flat index9Pon an 8×8 array, i.e.Σ start_i·stride_iverbatim. This PR does not fix [BUG] Metal slice_update can use a stale array-valued start across command buffers #3880; it bounds thedamage when a start is out of range for any reason. The two are independent.
slice_updateis the preallocate-and-update KV-cachepattern, where the start comes from sequence length, so an index exceeding the
preallocated capacity is a realistic way to reach this.
live in host-side C++ and in a scalar offset kernel, so A-series behaviour
should match, but that is reasoning rather than a measurement.
11. CUDA: fixed here, validated by CI rather than locally
I have no CUDA device, so the CUDA hunk is not compiled or run on my machine.
Leaving it out was not an option:
compute_dynamic_offsetis declared once inmlx/backend/gpu/slicing.handmlx/backend/gpu/primitives.cppcalls it forevery GPU backend, so adding the bounds argument without updating
mlx/backend/cuda/slicing.cppwould have left the CUDA builds unable to link.What I did to make the blind change defensible. The CUDA kernel is a JIT
template whose only include is
mlx/backend/cuda/device/utils.cuh, so I checkedeach thing it relies on against what that header actually provides:
max_startsmirrors the existingaxesparameter exactly — sameconst __grid_constant__ cuda::std::array<int, NIDX>type, sameargs.append(...)overload (std::vector<int>→SmallVector<int>→pointer,
jit_module.h:49-52), same lengthNIDX. Argument order inKernelArgsmatches the kernel signature.if constexpris already used in NVRTC-visible device headers(
device/binary_ops.cuh:19), so C++17 in this context is not an assumption.cuda::std::numeric_limits<T>::is_signedrather thanis_signed_v, becauseutils.cuhincludes<cuda/std/limits>directlywhile
<cuda/std/type_traits>only arrives transitively.Tis always a realfixed-width integer type here (
dtype_to_cuda_type,backend/cuda/utils.cpp:32-47),and the start dtype is validated as integer at the op level.
compiled it as host C++17 with
-Wall -Wextraagainst stubs for__global__,__grid_constant__,cuda::std::array,cuda::std::numeric_limitsandStrides, instantiated for all eight integerstart dtypes and NIDX of 1, 2 and 4: clean, no warnings. Running it then gave
the correct clamped offset on 12 probes, including a
uint64start of 2^64-1(clamps to the far end, not to zero). That covers syntax, types, the
if constexpr, and the arithmetic; it does not cover NVRTC itself or the__grid_constant__argument passing.R"( ... )"raw literal, so I checked the added linesintroduce no
)"sequence that would terminate it early.CUDA has no signedness bug of its own — unlike the CPU path, its kernel is
templated on the real dtype — so only the clamp was needed there.
What CI will prove. From
.github/workflows/build_and_test.yml, thebuild_and_testmatrix buildscuda-12.6,cuda-12.9andcuda-13.0on Linuxfor both architectures, and the test step runs when
matrix.os == 'Linux' && (matrix.toolkit == 'cpu' || matrix.arch == 'x86_64')(lines 71-72). So:
alone catches the link breakage described above and any host-side compile
error.
gpu-t4-4-core, a real T4. Thetest-linuxaction probes for a GPU and, when present, runs the Python suitewith
DEVICE=gpuand the C++ suite withDEVICE=gpu. That is where the newtest actually exercises the CUDA offset kernel, and where an NVRTC compile
error in the JIT source would surface — NVRTC compiles at runtime, so a
malformed kernel fails the test, not the build.
mac_buildplus a self-hostedmac_test, which is the Metal coverage Ialready ran locally.
What CI will not prove. Nothing checks the CUDA path before it runs, so if
the JIT source has a typo it shows up as a failing test rather than as a build
error, and I would need a maintainer's run to see it. If that happens the fix is
confined to one kernel body in
mlx/backend/cuda/slicing.cppand I will turn itaround quickly.
One correction worth flagging, since it changes what CI covers: I originally
guarded the gpu leg of the new test with
mx.metal.is_available(), which isfalse on a CUDA build — the gpu leg would have been silently skipped on the
T4 runners, so CI would have gone green whether or not CUDA was fixed. The guard
is now
mx.metal.is_available() or mx.cuda.is_available(), matching how the restof
python/testsguards CUDA, so the CUDA kernel is genuinely exercised.Files
0001-clamp-runtime-start.patch— the commit (6 files, +198/-19, includes theregression test), parented on
973e27f8case.py,drive.py— the coverage audit (311 cases per backend, one process each)before.txt,after.txt— raw per-case audit output on pristine and with the patchintree_before.txt,intree_after.txt— the in-tree test on pristine and with the patchrepro_probe.py,evidence_before.txt,evidence_after.txt— the earlier18-probe read/write matrix, kept as independent corroboration of §1