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.

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 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 at a null base plus a
field offset, reported as InvalidArgument from execute().

The cached load path already returns NotFound for the same condition, so this
makes the other path agree. A missing blob is harmless when a model has no
constants to bind, 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.
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 f94f8ad with merge base c27baa8 (image):

FLAKY - The following job failed but was likely due to flakiness present on trunk:

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

Copy link
Copy Markdown

CLA Missing ID

  • ❌ The email address for the commit (f94f8ad) 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
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