Skip to content

deps: bump pliron to master (no_std/hashbrown) + AttributeDict encapsulation sweep - #175

Merged
nihalpasham merged 3 commits into
mainfrom
chore/pliron-nostd-bump
Jun 12, 2026
Merged

nihalpasham merged 3 commits into
mainfrom
chore/pliron-nostd-bump

Conversation

@nihalpasham

Copy link
Copy Markdown
Collaborator

deps: bump pliron to master af1b7d0 + AttributeDict encapsulation sweep

Tracks upstream pliron master at af1b7d07dda9, picking up:

Both the root workspace and the isolated rustc-codegen-cuda workspace move to the same rev (required, or pliron resolves to two distinct crates and types stop unifying). Both lockfiles updated together.

Why the bump broke 4 call sites

hashbrown's get<Q: Hash + Equivalent<K>> no longer pins the key type, so call sites that reached into AttributeDict's inner map through the public .0 field with "literal".try_into().unwrap() keys stopped inferring (E0283: type annotations needed). Exactly 4 sites broke, all in mir-lower/src/convert/ops/memory.rs.

Call sites that already used the wrapper API (dialect-nvvm debug ops, AlignmentAttr helpers in llvm-export) were untouched by the change. The break only bit code that bypassed the abstraction.

The encapsulation sweep (the actual point)

Rather than patching the 4 sites with turbofish, this PR routes every keyed AttributeDict access through the typed wrapper, so no future map-internals change upstream can reach us:

Before: op.attributes.0.insert(key, attr.into())          // manual AttrObj boxing
After:  op.attributes.set(key, attr)                      // typed

Before: op.attributes.0.get(&key) + downcast_ref::<T>()   // two-step, two error paths
After:  op.attributes.get::<T>(&key)                      // downcasts for you

Before: op.attributes.0.contains_key(&gpu_kernel)
After:  op.attributes.get::<StringAttr>(&gpu_kernel).is_some()

Swept sites: callee/mutable stamps in the mir-importer translator, gpu_kernel + cluster-dim + launch-bound stamps and propagation (mir-importer body.rs, mir-lower lowering.rs), fastmath/nneg/overflow flag stamps (mir-lower arithmetic/cast/call), the nneg read and kernel metadata reads in the llvm-export emitter.

Whole-map operations keep using the public field because the wrapper has no equivalent yet: the MirFuncOp printer in dialect-mir (retain + is_empty on a clone_skip_outlined copy) and a keys() introspection in llvm-export's ops_test. Upstream wrapper candidates: remove/retain and key iteration.

One additional fallout fix: upstream relaxed set_operand_segment_sizes to take &Context; a mir-lower test passed &mut ctx and tripped clippy::unnecessary_mut_passed under -D warnings.

Verification

  • 4 inference breaks reproduced pre-fix, workspace check clean post-fix
  • All 5 compiler-crate test suites pass (101 tests, 0 failed); workspace test binaries compile
  • fmt, clippy -D warnings (--all-targets), doc -D warnings, doctests: clean
  • Backend rebuilt against the new rev; compile-only builds verified: vecadd (sm_90), tcgen05 (sm_100, 102 tcgen05 instructions in PTX), wgmma (sm_90a)
  • Rewritten attribute paths verified end-to-end in artifacts: cluster emits cluster_dim_x/y/z nvvm.annotations; debug emits maxntidx/minctasm metadata and .maxntid in PTX

Labels: build-related, refactor, IR-lowering, llvm-export

GPU verification (done)

Full smoketest on a 2x RTX 5090 box (sm_120): compile-only 92/92, execution 91/92. The one execution failure is cluster Test 2 (cluster sync), which fails identically on current main on this hardware and is expected for now per the maintainer; not related to this bump.

Upstream pliron PR #114 makes the framework no_std-capable and switches
its maps to hashbrown. hashbrown's get<Q: Hash + Equivalent<K>> no
longer pins the key type, so the four call sites that reached into
AttributeDict's inner map with `"literal".try_into().unwrap()` keys
stopped inferring. Route them through AttributeDict::get::<T>, whose
&Identifier parameter pins the key type and downcasts in one step.

Both lockfiles move together: the isolated rustc-codegen-cuda
workspace must stay on the same pliron rev as the root, or types stop
unifying across the boundary.
The hashbrown inference break in the previous commit only bit call
sites that bypassed AttributeDict's typed wrapper and poked at the
public .0 field. Sweep every remaining keyed access onto the wrapper
so no future map-internals change can reach us:

- .0.insert(key, attr.into()) becomes .set(key, attr); set takes the
  typed attribute directly, so the AttrObj boxing disappears from
  call sites (mir-importer callee/mutable/gpu_kernel/cluster/launch
  bound stamps, mir-lower fastmath/nneg/overflow flag stamps).
- .0.get + manual downcast_ref becomes .get::<ConcreteAttr>, which
  downcasts for us (llvm-export nneg read, kernel attr propagation).
- .0.contains_key(&gpu_kernel) becomes .get::<StringAttr>(..)
  .is_some(), matching the type the importer actually stores.

Whole-map operations stay on the public field because the wrapper has
no equivalent yet: the print-side clone_skip_outlined + retain +
is_empty in dialect-mir's MirFuncOp printer, and the keys()
introspection in llvm-export's ops_test.
…akes &Context

Upstream pliron relaxed the receiver context to a shared borrow;
clippy's unnecessary_mut_passed flags the old call shape under
-D warnings.
@nihalpasham nihalpasham added build-related Build scripts, toolchain discovery, or compilation setup llvm-export Textual LLVM/NVVM IR exporter (llvm-export crate) refactor Internal restructuring with no behavior change IR-lowering Lowering between dialects (dialect-mir to LLVM dialect) labels Jun 12, 2026
@nihalpasham
nihalpasham merged commit 540b960 into main Jun 12, 2026
24 checks passed
@nihalpasham
nihalpasham deleted the chore/pliron-nostd-bump branch June 19, 2026 11:30
meefs pushed a commit to meefs/cuda-oxide that referenced this pull request Sep 30, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

build-related Build scripts, toolchain discovery, or compilation setup IR-lowering Lowering between dialects (dialect-mir to LLVM dialect) llvm-export Textual LLVM/NVVM IR exporter (llvm-export crate) refactor Internal restructuring with no behavior change

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant