Report NotFound when a CUDA delegate's weights blob is missing - #22311
Conversation
🔗 Helpful Links🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/executorch/22311
Note: Links to docs will display an error until the docs builds have been completed. ❌ 3 New Failures, 2 Unrelated FailuresAs of commit 84d60b6 with merge base 2b3a32d ( NEW FAILURES - The following jobs have failed:
FLAKY - The following jobs failed but were likely due to flakiness present on trunk:
This comment was automatically generated by Dr. CI and updates every 15 minutes. |
meta-cla
Bot
added
the
CLA Signed
|
shoumikhin
added
the
release notes: runtime
f94f8ad to
04a3828
Compare
04a3828 to
04a4055
Compare
04a4055 to
26befe8
Compare
26befe8 to
9c5070b
Compare
| ET_CUDA_CHECK_OR_RETURN_ERROR(cudaDeviceSynchronize()); | ||
| buffer_res->Free(); | ||
| } else { | ||
| // The container expects a blob and did not get one, so its constant |
There was a problem hiding this comment.
i don't think it is useful to codify into comment a failure mode that is no longer applicable
There was a problem hiding this comment.
just in general very verbose
There was a problem hiding this comment.
Both addressed. Rather than only trimming the comment I removed the case it was arguing for: the count symbol has shipped since torch 2.6 and the bind symbol since 2.9, so a library with bind and no count is not something a released toolchain produces, and the NotSupported arm was carrying five lines of rationale for it. The block is now 5 comment lines instead of 10, and it explains only why a missing blob fails.
One correction in the other direction, from the Copilot comment on the same function: I had first folded the count into the early return, which made a library that can bind a supplied blob return Ok without binding it. That is the silent failure this change exists to prevent, so the early return now keys only on the bind function and the count is checked where it is actually needed.
There was a problem hiding this comment.
Trimmed. That block is two lines now and only says the part that is not visible in the code: the fetch is deferred because a file-backed data map reads the whole segment, which is gigabytes on a large model. The paragraph arguing about the older failure mode is gone, along with the case it was arguing for.
9c5070b to
5f71d6e
Compare
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.
Suppressed comments (1)
backends/cuda/runtime/cuda_backend.cpp:1273
- If
get_num_constantsis unavailable (it’s loaded as an optional symbol), the missing-blob path unconditionally calls it. After loosening the early-return to only depend onupdate_constants_from_blob, this becomes a potential null dereference.
Add a null-check here and conservatively fail the load when the blob is unavailable and the constant count can’t be queried.
// Without the blob the container's constant pointers stay null and the
// failure resurfaces much later as an illegal access inside a generated
// kernel, so report it while the cause is still identifiable. A model
// with no constants needs nothing bound and stays valid.
size_t num_constants = 0;
ET_CHECK_OK_OR_RETURN_ERROR(
handle->get_num_constants(handle->container_handle, &num_constants),
"Failed to enumerate CUDA AOTI constants");
| // A library built before external weights exports neither symbol, so there | ||
| // is nothing to do with a blob whether or not one was supplied. | ||
| if (handle->update_constants_from_blob == nullptr || | ||
| handle->get_num_constants == nullptr) { | ||
| ET_LOG( | ||
| Info, | ||
| "weights_blob '%s' is not used: this library cannot bind one", | ||
| weights_blob_key.c_str()); | ||
| return Error::Ok; | ||
| } |
There was a problem hiding this comment.
Good catch, this was a real regression and it is fixed.
The early return now keys only on update_constants_from_blob, so a library that can bind a supplied blob always binds it regardless of whether the count symbol is present. The count is checked in the missing-blob path where it is actually needed, with the null guard your suppressed comment asked for, so that path refuses with NotSupported rather than dereferencing a null or binding nothing.
Both of the cases you raised trace correctly now:
- bind present, count absent, blob present -> binds the blob
- bind present, count absent, blob absent -> refuses, no null dereference
5f71d6e to
da9eea9
Compare
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.
Suppressed comments (1)
backends/cuda/runtime/cuda_backend.cpp:1260
- If
update_constants_from_blobfails,buffer_res->Free()is not called before returning, which can leak the named-data buffer/mapping. Free the buffer on the error path as well.
auto update_err = handle->update_constants_from_blob(
handle->container_handle, static_cast<const uint8_t*>(weights_blob));
if (update_err != Error::Ok) {
ET_LOG(Error, "update_constants_from_blob failed");
return update_err;
| blobs = [f for f in os.listdir(outdir) if f.endswith(".ptd")] | ||
| self.assertTrue(blobs, "expected an externalized weights blob") | ||
| # A blob holding no constants is a bare header, and a missing blob | ||
| # really is harmless then, so the model has to carry real data for | ||
| # this to be testing anything. | ||
| largest = max(os.path.getsize(os.path.join(outdir, f)) for f in blobs) | ||
| self.assertGreater(largest, 256, "expected non-empty constants") | ||
|
|
||
| # A positive control first: the same program loads when its blob is | ||
| # supplied, so a later failure is about the missing blob rather than the | ||
| # program or the backend registration. The rewrite is covered by the | ||
| # payload check below it, which runs after the rewrite has happened. | ||
| sidecar = os.path.join(outdir, blobs[0]) | ||
| Runtime.get().load_program(path, data_path=sidecar).load_method("forward") |
There was a problem hiding this comment.
Good catch, fixed. The control now uses the largest blob and reuses that same path for the size assertion and for data_path, so os.listdir order cannot make it load a different file than the one it checked.
A CUDA delegate can keep its constants in a sidecar file that the caller supplies alongside the program. When that file is absent the container was created with null constant pointers and the load still succeeded, so the first execute died inside a generated kernel with an illegal memory access and nothing in the message naming the blob. The load now fails, naming the blob and how many constants the model expected to bind, and returning the status the data map gave rather than assuming the key was absent: a corrupt or unreadable sidecar reports a different error, and saying "not found" would send someone looking for a file they are holding. Only when the container can actually use a blob. A library built before external weights exports neither the bind function nor the constant count, so nothing here applies to it and such a model still loads. A model whose constant count is zero also still loads, since it has nothing to bind. The constant count query's status is checked, since a failed query would otherwise read as a real zero and let the load through with unbound constants, which is the case this exists to prevent. The cached loader a few lines above made the same unchecked call and now checks it too. Test plan: Two tests in backends/cuda/tests/test_missing_weights_blob.py, both run on Linux aarch64 and on an H100. The first exports a model with constants, loads it once with its sidecar to prove the program and the backend are fine, then rewrites the delegate payload into the older two-key form and loads it again without the sidecar. It fails on the merge base with "RuntimeError not raised" and passes here. The rewrite reads the payload from the program rather than scanning for it, since a program with several delegates holds several payloads. The test then asserts the metadata payload is gone, because otherwise a change in file layout would leave it exercising the weight cache path, which reports the same error number for the same program. Its weights key is derived from the shared library key, so it carries the library's hash and would not resolve even if a sidecar were supplied; that is deliberate, since the point is that an unresolvable key now fails the load. The second covers the branch that must keep loading: a model with no parameters or buffers loads without any sidecar and returns the right answer. That export already emits the two-key payload, so it needs no rewrite. Nothing else covers that branch. Both skip on ROCm, where the Python runtime is not built. Scope: this hardens the loader for artifacts whose payload names one blob. The current exporter emits per-constant metadata instead, so a model exported from main takes the weight cache path, which already reports a missing storage. A current export without its sidecar therefore already fails on the merge base, and the crash above is reachable only from older artifacts. The Metal backend emits the one-blob payload on every export and has the same gap in its loader, which is left for its own change. Not covered: the branch for a library that exports no bind function, which needs an artifact older than torch 2.9 to reach honestly. An error return from init leaves the extracted shared library on disk, which is true of every early return in init on the merge base too and belongs in its own commit.
da9eea9 to
d8f1dbd
Compare
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.
Suppressed comments (1)
backends/cuda/tests/test_missing_weights_blob.py:147
blobs[0]depends onos.listdir()ordering (not guaranteed) and can select the wrong.ptdif multiple sidecar blobs exist, making the positive-control load flaky. Select the largest (and tie-break deterministically) instead.
# supplied, so a later failure is about the missing blob rather than the
# program or the backend registration. The rewrite is covered by the
# payload check below it, which runs after the rewrite has happened.
sidecar = os.path.join(outdir, blobs[0])
Runtime.get().load_program(path, data_path=sidecar).load_method("forward")
| ET_CHECK_OK_OR_RETURN_ERROR( | ||
| handle->get_num_constants(handle->container_handle, &num_constants), | ||
| "Failed to enumerate CUDA AOTI constants"); |
There was a problem hiding this comment.
Done. The message now reads "Failed to enumerate CUDA AOTI constants for method '%s'", matching the log line right below it. The legacy path has the same message but no method name in scope, so that one is unchanged.
Two review points. The positive control took the first entry from `os.listdir`, which promises no order, so with more than one blob present the test could load a different file than the one it had just size-checked. It now picks the largest, which is the one the size assertion was about, and uses that same path for both. And the enumerate failure said only that it failed, while every other message in that function names the method. With several methods in a program that made the log hard to place, so the method name is in the text now. The legacy path has the same message but no method name in scope, so it is left as it is.
A CUDA delegate can keep its constants in a sidecar file that the caller supplies
alongside the program. When that file is absent the container was created with null
constant pointers and the load still succeeded, so the first execute died inside a
generated kernel with an illegal memory access and nothing in the message naming the
blob.
The load now fails, naming the blob and how many constants the model expected to bind,
and returning the status the data map gave rather than assuming the key was absent: a
corrupt or unreadable sidecar reports a different error, and saying "not found" would
send someone looking for a file they are holding.
Only when the container can actually use a blob. A library built before external
weights exports neither the bind function nor the constant count, so nothing here
applies to it and such a model still loads. A model whose constant count is zero also
still loads, since it has nothing to bind.
The constant count query's status is checked, since a failed query would otherwise
read as a real zero and let the load through with unbound constants, which is the case
this exists to prevent. The cached loader a few lines above made the same unchecked
call and now checks it too.
Test plan:
Two tests in backends/cuda/tests/test_missing_weights_blob.py, both run on Linux
aarch64 and on an H100.
The first exports a model with constants, loads it once with its sidecar to prove the
program and the backend are fine, then rewrites the delegate payload into the older
two-key form and loads it again without the sidecar. It fails on the merge base with
"RuntimeError not raised" and passes here. The rewrite reads the payload from the
program rather than scanning for it, since a program with several delegates holds
several payloads. The test then asserts the metadata payload is gone, because
otherwise a change in file layout would leave it exercising the weight cache path,
which reports the same error number for the same program. Its weights key is derived
from the shared library key, so it carries the library's hash and would not resolve
even if a sidecar were supplied; that is deliberate, since the point is that an
unresolvable key now fails the load.
The second covers the branch that must keep loading: a model with no parameters or
buffers loads without any sidecar and returns the right answer. That export already
emits the two-key payload, so it needs no rewrite. Nothing else covers that branch.
Both skip on ROCm, where the Python runtime is not built.
Scope: this hardens the loader for artifacts whose payload names one blob. The current
exporter emits per-constant metadata instead, so a model exported from main takes the
weight cache path, which already reports a missing storage. A current export without
its sidecar therefore already fails on the merge base, and the crash above is reachable
only from older artifacts. The Metal backend emits the one-blob payload on every
export and has the same gap in its loader, which is left for its own change.
Not covered: the branch for a library that exports no bind function, which needs an
artifact older than torch 2.9 to reach honestly. An error return from init leaves the
extracted shared library on disk, which is true of every early return in init on the
merge base too and belongs in its own commit.