Repository navigation
Derive intrinsic PTX floors through the existing target vocabulary owner #1213
Description
What problem does this solve?
minimum_ptx is written in both the intrinsic overlay and Rust recipe. Validation only compares the copies (validation). Nothing downstream checks it. ptx_isa_requirement_for_floor returns Default for every encoded floor <= 60 (codegen), so a wrong floor through PTX 6.0 changes only the generated reference.
The #1138 review named this shape:
The recorded libNVVM evidence was a hand-built fixture rather than real lowering output, which is why nothing caught it.
Where minimum_sm is a hardware requirement, the effective floor can be derived instead of copied.
Proposed API / behavior
AS-IS TO-BE
recipe "1.0" --. minimum_sm
overlay "1.0" --+ ensure!(equal) |
| v
| no external TargetPtxFloor (in-tree)
| floor input |
v v
catalog.json max(target default,
/ \ instruction requirement)
encoded<=60 verbatim |
| | v
Default "1.0" in catalog.json
1.0 = ... = 6.0 reference
TargetPtxFloor / RECORDED_PTX_FLOORS already store the per-target floor (cuda-target-spec). #1105 introduced them, and #1184 routed target parsing and PTX spelling through them. Intrinsic minimum_ptx still does not use them.
# before
minimum_ptx = "1.0"
minimum_sm = "sm_20"
# after — floor derived as PTX 2.0
minimum_sm = "sm_20"This rule applies only when minimum_sm is a hardware requirement. Legacy rows such as sm_20 -> 2.0 would be recorded as pinned ISA history. Instructions needing more than their target floor would keep a written value.
Alternatives considered
Correcting each floor separately leaves the self-check unchanged. #519 fixed trap/breakpoint/pmevent; the 1.0 / 3.2 / sm_20 shape remains in counted bar.arrive (overlay), from merged #721.
Additional context
Of 130 explicit overlay routes, 104 meet or exceed the target floor. The 26 below are lib_nvvm at sm_75: 16 at PTX 6.0, 7 at 4.2, 2 at 1.0, and 1 at 6.2. Here sm_75 records evidence collection, not instruction hardware. The #519 thread explains:
backend-codegen evidence must sit exactly at the hardware floor and CUDA 13 cannot probe older targets; the PTX floor alone was the blocker.
minimum_sm has two route meanings. That distinction must be resolved before or with floor derivation; no fix is proposed here.
Does anything depend on the per-entry declaration that I have not seen? Should this cover merged declarations or new entries only? #1189 records a second consumer left outside the same owner.
#1212 opened this as a PR before any discussion and was closed for that; this issue is that discussion.
Activity
nihalpasham commentedon Sep 3, 2026
Where I stand: the defect is real but small, and it is catalog hygiene, not a codegen gap. Every target cuda-oxide emits for is sm_70+ and llc prints that target's default
.version(7.0 at sm_80); a floor only reaches the PTX when it is above that default. A wrong floor at or below 6.0 changes the reference and the generated tables, not any kernel.Deriving from
RECORDED_PTX_FLOORSis out. Its doc says the entries are LLVM 23 defaults, "not backend-independent CUDA facts", and the ISA keeps adding instructions to an existing SM later:The
llvm_nvptxroute already holds what derivation would give: its floor must equal the.versionLLVM emitted at the evidence target (3.2 at sm_20), checked inresolve/evidence.rs.Your two questions:
validate_selected_target_predicatesrequires it to equal LLVM'sgetPTXVersion() >= Ngate when the imported record has one (the counted barrier records have none, which is why 1.0 passed), evidence validation compares against it, and codegen falls back to it when a backend has no route.minimum_sm) can cover every entry, but know its yield. Today it flags 39 places, none a hardware floor: 32 libNVVM routes atsm_75(the target the CUDA 13 proof was compiled for; the route field is documented as the backend-profile floor, leave it) and 7 packed f16x2 ops atsm_70(my oldest supported GPU;native_minimum_sm = 53holds the hardware fact, compare against that). It would not have caughtbar.arriveeither:targets = "all"and nominimum_sm, so only the ISA text catches that one.What I would take as one PR:
bar.sync/bar.arriveto PTX 2.0 / sm_20 per the ISA, overlay plus theexecution_control.rsrecipe, regenerate (feat(intrinsics): add CTA barrier reductions #1138 is closed, socounted_barrier.tomlis free)check) and one controljust checkcannot pass inside the recommendednix developshell: the probe pinsllcby SHA-256, nixpkgs rewrites it #1211) costs more than it returnsBar for reference: https://discord.com/channels/1515530041767759993/1515561621156663487/1540970771323551775
Sequencing: this one comes after #1211 and #1189 (one series at a time; all three touch the target vocabulary or the catalog artifacts). Does that scope work for you?
feat(target-spec): record the PTX version that first named each target
nihalpasham commentedon Sep 23, 2026
Merged #1261 after #1256. This completes the scope we agreed here: counted barrier floors are corrected to PTX 2.0 / sm_20, authored hardware floors are checked against the separate SM-introduction table, and the negative/control tests and regenerated catalog artifacts are in place. Closing this as completed.