Skip to content

Report NotFound when a CUDA delegate's weights blob is missing - #22311

Open
shoumikhin wants to merge 1 commit into
pytorch:mainfrom
shoumikhin:fix/cuda-missing-weights-blob
Open

Report NotFound when a CUDA delegate's weights blob is missing#22311
shoumikhin wants to merge 1 commit into
pytorch:mainfrom
shoumikhin:fix/cuda-missing-weights-blob

Conversation

@shoumikhin

Copy link
Copy Markdown
Contributor

The CUDA backend moves a model's constants into a separate weights blob and binds them while loading the method. When that blob is not supplied, the load path logged a message at Info level and returned Ok, so the model finished initialising with its constant pointers still null. The program then failed much later, inside a generated kernel, as an illegal memory access:

Invalid __global__ read of size 4 bytes
  at triton_poi_fused_mul_0+0x170
  Address 0x40 is out of bounds

That address is a null base plus a field offset, and the reported error was 0x12 (InvalidArgument) from execute(), which points nowhere near the real cause.

The cached load path already returns NotFound for the same condition, so this makes the other path agree.

A missing blob is genuinely harmless when a model has no constants to bind, and that case is common, so the check keys on the constant count rather than on the blob being absent. get_num_constants is an optional symbol; when it is not available the two cases cannot be told apart, so the previous permissive behaviour is kept.

Before, on a model with three buffers whose blob was not supplied:

RESULT execute RAISED RuntimeError: method->execute() failed with error 0x12

After:

RESULT init RAISED RuntimeError: Failed to load method forward, error: 0x:20

0x20 is Error::NotFound.

Test plan

Added backends/cuda/tests/test_missing_weights_blob.py. It lowers a module with three buffers, checks the externalized blob holds real data rather than being an empty placeholder, then loads the program without it and requires the load to fail.

Verified both directions against a build of this branch:

  • without the change: 1 failed, AssertionError: RuntimeError not raised, and the log shows initialisation completing (container handle and CUDA stream both created)
  • with the change: 1 passed

Built and ran on an aarch64 CUDA machine with GCC 13.3. clang-format reports no changes for the modified file, and ufmt reports the new test already formatted.

Copilot AI lite review requested due to automatic review settings August 29, 2026 15:35
@pytorch-bot

pytorch-bot Bot commented Aug 29, 2026

Copy link
Copy Markdown

🔗 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.

✅ You can merge normally! (1 Unrelated Failure)

As of commit 9c5070b with merge base c27baa8 (image):

BROKEN TRUNK - The following job failed but were present on the merge base:

👉 Rebase onto the `viable/strict` branch to avoid these failures

This comment was automatically generated by Dr. CI and updates every 15 minutes.

@meta-cla meta-cla Bot added the CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed. label Aug 29, 2026

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

@linux-foundation-easycla

linux-foundation-easycla Bot commented Aug 29, 2026

Copy link
Copy Markdown

CLA Missing ID

  • ❌ The email address for the commit (9c5070b) is not linked to the GitHub account, preventing the EasyCLA check. Consult this Help Article and GitHub Help to resolve. (To view the commit's email address, add .patch at the end of this PR page's URL.) For further assistance with EasyCLA, please visit our EasyCLA portal and chat with our support bot.

@shoumikhin shoumikhin added the release notes: runtime Changes related to the core runtime which loads the program methods, initializes delegates, and runs label Aug 29, 2026
Copilot AI review requested due to automatic review settings August 30, 2026 18:07
@shoumikhin
shoumikhin force-pushed the fix/cuda-missing-weights-blob branch from f94f8ad to 04a3828 Compare August 30, 2026 18:07

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

@shoumikhin
shoumikhin force-pushed the fix/cuda-missing-weights-blob branch from 04a3828 to 04a4055 Compare August 31, 2026 04:35
Copilot AI review requested due to automatic review settings August 31, 2026 04:35

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

@shoumikhin
shoumikhin force-pushed the fix/cuda-missing-weights-blob branch from 04a4055 to 26befe8 Compare August 31, 2026 04:52
Copilot AI review requested due to automatic review settings August 31, 2026 04:52

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

A CUDA delegate keeps 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. Forgetting the data file is an easy mistake, and the error gave no
hint of it.

The load now fails, naming the blob and how many constants the model expected to
bind, and reporting why the blob was unavailable rather than assuming it was
absent: a corrupt or unreadable sidecar returns a different status, and saying "is
missing" 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 no function to bind one, so nothing here applies to it whether or
not a file was supplied, and such a model still loads. That case is checked first,
and it no longer claims to know why: the exported symbol tracks the torch version,
not where the constants live, so it cannot tell an embedded-constants library from
an external one.

The constant count query's status is checked too, 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.

Test plan:

backends/cuda/tests/test_missing_weights_blob.py 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.

The rewrite reads the payload from the program rather than scanning for it, since a
program with several delegates holds several payloads and a scan cannot tell where
one ends. The test then asserts the metadata payload is gone, because otherwise a
change in file layout would leave it exercising the weight cache path instead,
which reports the same error number for the same program.

Verified in both directions on Linux aarch64. Without the change the load succeeds
and the test fails. With it the load reports

  weights_blob '..._weights_blob' is unavailable (0x20), but the model has 2
  constant(s) to bind

and fails, and the test passes.

The test skips on ROCm, where the Python runtime is not built.

Not covered: this hardens the older two-key payload, which the current exporter no
longer emits, so it applies to older artifacts. The Metal backend uses that payload
for every export and has the same gap, which is left for its own change.
Copilot AI review requested due to automatic review settings August 31, 2026 05:13
@shoumikhin
shoumikhin force-pushed the fix/cuda-missing-weights-blob branch from 26befe8 to 9c5070b Compare August 31, 2026 05:13

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed. release notes: runtime Changes related to the core runtime which loads the program methods, initializes delegates, and runs

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants