Skip to content

fix(mir-importer): dispatch sync_threads via its canonical thread:: path#360

Open
haixuanTao wants to merge 1 commit into
NVlabs:mainfrom
haixuanTao:fix/thread-sync-threads-dispatch-alias
Open

fix(mir-importer): dispatch sync_threads via its canonical thread:: path#360
haixuanTao wants to merge 1 commit into
NVlabs:mainfrom
haixuanTao:fix/thread-sync-threads-dispatch-alias

Conversation

@haixuanTao

Copy link
Copy Markdown
Contributor

Summary

Every sreg accessor arm in try_dispatch_intrinsic matches both the crate-root re-export path and the canonical module path ("cuda_device::threadIdx_x" | "cuda_device::thread::threadIdx_x"), but the barrier arm only matches "cuda_device::sync_threads". A call that resolves through the canonical cuda_device::thread::sync_threads path — e.g. a dependency crate calling the API by full path — misses intrinsic dispatch and falls through to a plain function call against a symbol no definition provides, failing module verification with "Symbol cuda_device__thread__sync_threads not found". (Found while porting a multi-crate shader codebase whose support crate calls cuda_device::thread::sync_threads() directly; mentioned in #359.)

Changes

  • Add the missing "cuda_device::thread::sync_threads" alias to the barrier dispatch arm, matching the sreg arms' pattern.
  • Swept the other dispatch arms against the cuda_device module tree: sync_threads was the only function with a root-path arm but no canonical-path alias.

Testing

  • cargo test -p mir-importer — all green.

  • Downstream validation: a support crate calling cuda_device::thread::sync_threads() by full path fails module verification without this change and compiles + runs correctly (barrier lowers to bar.sync) with it.

  • cargo test passes

  • New example — not applicable (one-line dispatch alias)

Checklist

  • All commits signed off (git commit -s)
  • SPDX headers on new source files (no new files)

🤖 Generated with Claude Code

Every sreg accessor arm matches both the crate-root re-export path and the
canonical module path ("cuda_device::threadIdx_x" |
"cuda_device::thread::threadIdx_x"), but the barrier arm only matched
"cuda_device::sync_threads". A call that resolves through the canonical
`cuda_device::thread::sync_threads` path (e.g. from a dependency crate
calling the API by full path) missed intrinsic dispatch and fell through to
a plain function call against a symbol no definition provides, failing
module verification with "Symbol cuda_device__thread__sync_threads not
found". Add the missing alias, matching the sreg arms.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: haixuantao <tao.xavier@1ms.ai>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant