Fix dropped alignment operand in byteAddressBufferLoad specialization - #1
Closed
stramit wants to merge 1 commit into
Closed
Fix dropped alignment operand in byteAddressBufferLoad specialization#1stramit wants to merge 1 commit into
stramit wants to merge 1 commit into
Conversation
specializeFuncsForBufferLoadArgs lumped kIROp_ByteAddressBufferLoad in
with the generic two-operand "element access" instructions (GetElement,
StructuredBufferLoad, etc.) in both getCallInfoForArg and
getSpecializedValueForArg. byteAddressBufferLoad actually has three
operands -- (buffer, offset, alignment) -- so the rewritten load inside
the specialized callee was constructed with only two operands.
Downstream, legalizeByteAddressBufferOps::processLoad unconditionally
reads operand index 2 to recover the alignment, hitting
assert failure: slang-ir.h(710): index < getOperandCount()
during SPIR-V emission. A minimal reproducer is a struct large enough
to trip the buffer-load specialization threshold (e.g. two float4x4s
plus a padded float3) loaded via Buf.Load<T>() and passed by value into
a helper function.
Give byteAddressBufferLoad dedicated handling in both functions so the
alignment operand survives specialization, and add a regression test
under tests/optimization that exercises the failing pattern on SPIR-V.
stramit
force-pushed
the
claude/sweet-darwin-AoOG8
branch
from
August 21, 2026 06:21
00d1206 to
e70a4a0
Compare
stramit
pushed a commit
that referenced
this pull request
Aug 21, 2026
…-casting (shader-slang#11446) ## Summary Fixes the undefined behaviour surfaced by the **Nightly Sanitizer** run [26861035639](https://github.com/shader-slang/slang/actions/runs/26861035639): ``` source/core/slang-string-util.cpp:550:29: runtime error: member call on address … which does not point to an object of type 'ISlangBlob' note: object is of type 'hlsl::InternalDxcBlobEncoding_Impl<hlsl::DxcBlobUtf8_Impl>' #0 StringUtil::getSlice(ISlangBlob*) #1 _getSlice(IDxcBlob*) source/compiler-core/slang-dxc-compiler.cpp:68 #2 _handleOperationResult(…) shader-slang#3 DXCDownstreamCompiler::compile(…) ``` ## Root cause DXC's `IDxcBlob` and Slang's `ISlangBlob` are *distinct* COM interfaces that happen to be ABI/layout-compatible (`IUnknown` + a buffer pointer/size getter at the same vtable slots). `slang-dxc-compiler.cpp` reinterpret-casts `IDxcBlob*` → `ISlangBlob*` and calls `ISlangBlob`'s virtuals on the result. It works at runtime, but it is UB — the object isn't actually an `ISlangBlob` — and UBSan's `-fsanitize=vptr` flags it. **Why it started failing now:** shader-slang#10935 ("Build DXC from source when system GLIBC does not match") means DXC is now compiled inside the sanitizer-instrumented build, so UBSan has RTTI for the DXC blob types and can finally see the type mismatch that was invisible while DXC was a prebuilt binary. The CI triage confirms it: *"1 PR-related, 0 pre-existing."* ## Fix Add a small `DxcBlob` adapter — a genuine `ISlangBlob` (deriving from `BlobBase`) that holds a ref to the `IDxcBlob` and forwards `getBufferPointer`/`getBufferSize` to its `GetBufferPointer`/`GetBufferSize`. Use it at the flagged sites: `_getSlice` and the three `addRepresentationUnknown` calls. Holding the ref also keeps the DXC blob's data alive for stored artifact representations (the old pun relied on the caller's `ComPtr`). ## Scope / follow-up The `IDxcLinker::RegisterLibrary` path (`libraryBlobs` … `(IDxcBlob*)libraryBlobs[i].get()`) round-trips blobs `ISlangBlob* ↔ IDxcBlob*` in both directions. De-punning it requires creating a real DXC blob via `IDxcLibrary` (not in scope in that function), so it's left as a follow-up. ## Validation I could not build/run this locally (macOS, no DXC/DXIL toolchain). I verified the *mechanism* with a standalone repro: the reinterpret-pun trips `-fsanitize=vptr` with the identical diagnostic, while the adapter form is clean. **Relying on this PR's sanitizer CI to validate the actual build** — the branch was pushed to the main repo so the Nightly Sanitizer can run against it.
stramit
pushed a commit
that referenced
this pull request
Aug 21, 2026
…1656) ## Motivation When the compiler cannot specialize a generic for a given call, it almost always reports the catch-all `E39999` — *"could not specialize generic for arguments of type ..."* — which tells the user *that* specialization failed but not *why*. PR shader-slang#11571 introduced a *specialization-failure-reason* mechanism (a tagged union, `GenericArgumentInferenceFailure`) and wired up exactly one focused reason (variadic pack-count mismatch). This PR extends that mechanism to the common cases that currently surface the confusing `E39999`, so the diagnostic names the offending counts, parameter, constraint, or conflicting types instead. ```slang void f<T>(T x) {} f(1, 2); // too many call arguments T pick<T>(int which) { ... } // T appears only in return position let v = pick(0); // T cannot be inferred from the call void need<T : IFoo>(T x) {} struct S {} // S does not conform to IFoo S s; need(s); // unsatisfied interface conformance void h<T>(T x) where T == int {} float f; h(f); // float does not satisfy `T == int` struct Foo<T, U> {} Foo<int> b; // explicit arg list: 1 given, 2 expected void g<T>(T a, T b) {} A a; B b; g(a, b); // T must be both A and B (no common type) ``` Before this change each of these reports a generic catch-all — `E39999` *"could not specialize generic for arguments of type ..."*, or the bare `E30075` *"cannot specialize generic ... with the provided arguments"* for the explicit argument-list case — with no specifics. After it each reports a focused diagnostic. ## Proposed solution Follow the design the maintainer settled on the issue: **keep the tagged-union infrastructure and add one union case per inference-failure reason, capturing only cheap structured fields at the failure site and deferring all diagnostic *formatting* to `CompleteOverloadCandidate`** (which runs only when the failed candidate is actually selected for emission). Most specialization failures happen on speculative overload candidates that never surface a diagnostic, so the capture must stay cheap and side-effect-free; the union exists precisely so that formatting is paid for only when needed. Each inference-failure site records its fields guarded by `failure && kind == None` (first recorded reason wins, so the existing variadic path is never clobbered), and `CompleteOverloadCandidate` switches on `kind` to format the focused diagnostic + a `GenericSignatureTried` note, falling back to `E39999` for `Kind::None`. The reasons captured into the union: 1. **Call-argument arity mismatch** (`GenericArityMismatch` → `E30438`) — the call supplies more value arguments than the generic's parameter list can match. Captured when `matchArgumentsToParams` fails (which it does for an over-supply; an under-supply under-fills and instead surfaces as the not-inferred reason below). 2. **Ordinary type/value parameter not inferred** (`OrdinaryGenericParamNotInferred` → `E30439`) — including a type `T` that appears only in return position and a value `N` not mentioned in any parameter. Captured at **two** sites: `areOrdinaryConstraintsSatisfied` and the ordinary-parameter branch of `areFinalArgsValid`. Per maintainer feedback it now stores the parameter **declaration** (not just its name), so the diagnostic has the full decl. 3. **Interface conformance not satisfied** (`InterfaceConformanceNotSatisfied` → reuses `E38029`) — a required `T : IFoo` subtype witness cannot be proven. 4. **Any non-conformance source constraint not satisfied** (`GenericConstraintNotSatisfied` → new `E30440` + `see generic constraint declaration` note) — the general fallback for *every* constraint kind the witness solver handles other than interface conformance: equality `where T == int`, type coercion `where int(T)`, non-empty pack `where nonempty(T)`, etc. Captured uniformly at the readiness-confirmed witness-failure point (not per-kind), and the constraint is rendered adaptively by a new `getGenericConstraintFailureString` printing utility, so the message reflects the actual constraint — *"could not satisfy the generic constraint 'T == int'"* / `'T -> int'` / `'nonempty(T)'` — rather than assuming an equality. 5. **Generic-parameter unification conflict** (`GenericParamUnificationConflict` → new `E30442`) — a single generic parameter is constrained to two conflicting arguments. Covers **type** parameters with no common type (`g<T>(T, T)` with unrelated `A`/`B`) **and** value parameters forced to two different values (`bar<let N:int>(int[N], int[N])` with mismatched array sizes → `N` both `3` and `4`). The payload stores both candidates as `Val*` (a `Type` for type params, an `IntVal` for value params); reports both. One more focused diagnostic is emitted **outside** the union, at the explicit-generic-argument-list site: 6. **Explicit generic-argument-list arity** (new `E30441`) — an explicit list such as `Foo<int>` for `Foo<T, U>` supplies the wrong number of arguments. This is decided in `TryCheckGenericOverloadCandidateTypes` from the generic's parameter-type count, *before* the inference-failure union is involved, so it is reported there directly rather than threaded through the union. ## Change summary | File | What changed | | --- | --- | | `source/slang/slang-check-impl.h` | Add five `Kind` enumerators and payload structs (`GenericArityMismatch`, `OrdinaryGenericParamNotInferred` storing a `Decl*`, `InterfaceConformanceNotSatisfied`, `GenericConstraintNotSatisfied`, `GenericParamUnificationConflict`) and their union members. Hand-written copy semantics route through `copyActiveMemberFrom()` (`switch(kind)`), one `set*()` placement-new helper per variant, and a `static_assert` that every payload is trivially destructible. | | `source/slang/slang-check-overload.cpp` | Capture the call-arity reason in `inferGenericArguments`; `CompleteOverloadCandidate` switches on `kind` to emit each focused diagnostic + `GenericSignatureTried`, with `None` → `E39999`. In `TryCheckGenericOverloadCandidateTypes`, an explicit-arg-list count mismatch (non-variadic) now emits `E30441` instead of the bare `cannot specialize generic` (`E30075`); in practice this is reached only for under-supplied lists. | | `source/slang/slang-check-constraint.cpp` | Add the `getGenericConstraintFailureString` printing utility. Capture the not-inferred reason in `areOrdinaryConstraintsSatisfied` / `areFinalArgsValid` (recording the `Decl*`); capture the conformance reason in `trySolveSubtypeWitnessForConstraint`; capture the **general** unsatisfied-constraint reason uniformly in `solveWitnessConstraint` at the readiness-confirmed witness-failure point; capture the unification conflict in **both** `solveTypeParamConstraint` (type params) and `solveValueParamConstraint` (value params). | | `source/slang/slang-diagnostics.lua` | Add `generic-specialization-arity-mismatch` (30438), `generic-parameter-could-not-be-inferred` (30439), `generic-argument-does-not-satisfy-constraint` (30440, message `"could not satisfy the generic constraint '~constraint:String'"`) + `see-generic-constraint-declaration` note, `generic-argument-list-arity-mismatch` (30441), `generic-parameter-unification-conflict` (30442). The conformance case reuses existing `type-argument-does-not-conform-to-interface` (38029). | | `tests/diagnostics/` + `tests/language-feature/` | Nine focused tests (arity, value `N`, return-only `T`, conformance, equality-constraint, coercion-constraint, nonempty-constraint, explicit-arg-list arity, type+value unification conflict); nine pre-existing tests updated from `E39999`/`cannot specialize` to the focused diagnostics (strictly-better output on already-invalid programs — see process report). | ## Concepts and vocabulary - **`GenericArgumentInferenceFailure`** — the tagged union from PR shader-slang#11571: a `Kind` enum + a payload `union` + hand-written copy semantics, threaded through inference as `GenericInferenceContext::failure`. - **Deferred formatting** — capture sites store only fields; the `Diagnostics::*` sink calls live solely in `CompleteOverloadCandidate`, so speculative candidates pay nothing. - **`solveWitnessConstraint`** — the work-list handler for a source generic constraint. It returns `Blocked` (via `hasUnreadyDependenciesForWitnessConstraint`) while the constraint's inputs are not yet ready, and only calls `trySolveWitnessForConstraint` once they are; so a null witness there is a *readiness-confirmed* genuine failure — the uniform point to record the general unsatisfied-constraint reason for every constraint kind. - **`trySolveSubtypeWitnessForConstraint`** — solves a single `GenericTypeConstraintDecl` (conformance `T : IFoo` or equality `where T == X`, flagged by `isEqualityConstraint`); only the conformance case records its own specific reason, the rest defer to the general capture above. - **`getGenericConstraintFailureString`** — renders a source constraint declaration to a short readable form per kind (`T : I`, `T == X`, coercion `T -> U`, `nonempty(P)`), so the diagnostic adapts to the actual constraint. - **`mergeTypeConstraint`** — joins two inferred candidate types for the same type parameter; returns `false` (without mutating the first candidate) when they have no common type. - **`hasUnreadyDependenciesForVal`** — true while a `Val` still mentions a free generic parameter of the generic being solved (i.e. not yet concrete). - **`GenericSignatureTried`** — the existing note pointing at the generic declaration that was being specialized. ## Process report **Call-arity capture (`inferGenericArguments`, `slang-check-overload.cpp`).** `matchArgumentsToParams` is where value-argument-to-parameter matching fails before inference runs; today that lands at `E39999`. We record expected-vs-supplied counts and emit `E30438`. It replaces the `E39999` fallback rather than duplicating the non-generic *"too many arguments to call"* path (a separate code path). Input-shape check: a genuine arg/param-count mismatch on a generic `CallableDecl`; the count is known here and nowhere more specific. **Ordinary-parameter capture — two sites, now storing the `Decl`.** A parameter that forms an ordinary solver constraint fails in `areOrdinaryConstraintsSatisfied`; a type `T` only in return position, or a value `N` in no parameter, forms no ordinary constraint and fails later in `areFinalArgsValid`. Both sites store the parameter `Decl*` (guarded only on a non-null decl — the `getName()` check was dropped per maintainer feedback) so the diagnostic has the full declaration, and emit `E30439`. Input-shape check: both are real "parameter stayed unsolved" states; neither masks malformed input — they refine where the existing `GenericArgumentInferenceFailed` status was already set. **Constraint-failure capture — specific conformance, general everything-else (maintainer asks #1 + shader-slang#4).** csyonghe asked for the unsatisfied-constraint message to *not* hard-code an equality form and to be the general fallback for every constraint kind, with a printing utility and coverage for coercion/non-empty constraints. The capture is now split across two layers: - **conformance**, captured *specifically* in `trySolveSubtypeWitnessForConstraint` (`!isEqualityConstraint`): records the substituted sub/interface and reuses `E38029` (*"does not conform to interface"*), since that message already matches and a second code would split one concept. Gated on **concreteness** (`!hasUnreadyDependenciesForVal`) because that routine can run while the subject is still a free generic parameter. - **everything else** (equality, coercion `where U(T)`, non-empty `where nonempty(P)`, ...), captured *uniformly* in `solveWitnessConstraint` at the point where `trySolveWitnessForConstraint` returns null. This is the authoritative, readiness-confirmed failure point: the `hasUnreadyDependenciesForWitnessConstraint` check just above returns `Blocked` for any not-yet-ready constraint, so a null witness here means a genuine, fully-substituted failure — no separate concreteness guard or masking risk (a `Failed` return rejects the candidate immediately; there is no later retry-to-success). It records only the source `constraintDecl` (kept cheap; `kind == None` so the specific conformance reason, or a variadic pack-count reason, still wins) and `CompleteOverloadCandidate` renders it via `getGenericConstraintFailureString`, so the message adapts: `'T == int'`, `'T -> int'`, `'nonempty(T)'`. This is why coercion and non-empty constraints now get a focused message on the inference path (their dedicated `E38043`/`E30414` still fire on the explicit-argument path, which this PR leaves unchanged — verified: `diagnose-coerce-constraint-missing` / `diagnose-empty-pack-nonempty-constraint` are unaffected). Input-shape check: reaching the conformance site means the abstract-self-type and optional-constraint cases have already returned, so this is a genuine unsatisfiable constraint reported at the right layer, not a guard around malformed input. **Unification-conflict capture — type *and* value params (maintainer ask #2).** For a **type** parameter, two inferred candidates with no common type make `mergeTypeConstraint` return `false` *before* mutating the first candidate (`solveTypeParamConstraint`), so `type`/`cType` are the two candidates. csyonghe asked for the same on **value** parameters: `solveValueParamConstraint` rejects a candidate when two same-priority constraints disagree (`!val->equals(cVal)`, e.g. `N` required to be both `4` and `8`), so the conflict is captured there too. The `GenericParamUnificationConflict` payload now stores the parameter `Decl` (not just its name, per the same ask) and both candidates as `Val*` — a `Type` for type params, an `IntVal` for value params — and `CompleteOverloadCandidate` renders each candidate with `toText`. Input-shape check: both are real conflicting-binding states on a live parameter, the natural place both candidates are in hand. **Explicit generic-argument-list arity (`TryCheckGenericOverloadCandidateTypes`).** The general error reporter here previously emitted only the bare `cannot specialize generic`. It now compares the provided argument count against the generic's collected parameter-type count and, when they differ, emits `E30441` naming both counts; otherwise it keeps the `cannot specialize` fallback. A variadic generic (a trailing type/value pack, detected via `isPackType`) has no single expected count, so it is excluded from the arity message and keeps the general fallback. In practice this fires for *under*-supplied explicit lists (`Foo<int>`, `Buffer<>`, `RWTex<float3>`); an *over*-supplied list is intercepted earlier by a check that already prints its own counted *"too many arguments to call"* message, so `E30441` does not duplicate it. Input-shape check: the expected count is known here from the generic's own parameter list, so this is the correct layer for an explicit-list arity error. **Union member lifetimes (`slang-check-impl.h`).** Each payload embeds `SourceLoc`, which has a user-provided copy/default constructor and is therefore not trivially copyable; writing into a union member whose lifetime has not begun is undefined. `copyActiveMemberFrom` (used by copy-ctor/assign) and one `set*()` helper per variant both placement-new the selected member and keep `kind` in lockstep, so every capture site begins the member's lifetime before writing it and the invariant "`kind` names the live union member" holds by construction. A `static_assert` proves every payload is trivially destructible, which is what makes destroying the previously-active member unnecessary before the placement-new; `copyActiveMemberFrom`'s `switch` has no `default` (only an explicit `case Kind::None`) so a future enumerator triggers `-Wswitch` rather than being silently dropped. **Pre-existing tests now assert focused diagnostics.** Nine already-invalid programs previously locked in `E39999` / `cannot specialize generic` and now assert the focused diagnostics this PR adds — `E30438`/`E30439` (e.g. `lambda-and-function-type-generic-inference-2.slang`, `generic-function-default-param-errors.slang`, which carried a `// TODO` asking for exactly this), `E30440` (`extension-with-where-clause-1.slang`'s `where A == B` equality failure, and `variadic-pack-query-nonempty-constraint.slang`'s `nonempty(D)` on an empty pack), and `E30441` (`invalid-buffer.slang`'s `Buffer<>`, `shader-slanggh-4150.slang`'s `RWTex<float3>`). These are strictly-better messages on programs that were already errors; no valid program changes behaviour. One judgment call flagged in the initial round still stands: `extension-visibility.slang` / `generic-default-solver-failure.slang` surface the focused diagnostics on their (invalid) default-argument-substitution paths as well as inferred-argument paths; this is a genuine improvement, but the guards can be narrowed to inferred-argument failures only if preferred. Out of scope (left for follow-up, per the issue): the remaining null-`DeclRef` sites (pack-unification mismatch, work-list non-convergence, nested-generic missing-arg-array, etc.). The autodiff caller `inferGenericArguments(... outFailure=nullptr ...)` still drops failures and is untouched (it never reads the new union state). Fixes shader-slang#11643 --------- Co-authored-by: nv-slang-bot[bot] <274397474+nv-slang-bot[bot]@users.noreply.github.com>
stramit
pushed a commit
that referenced
this pull request
Aug 21, 2026
…r is present (shader-slang#11920) ## Motivation `linkAndOptimizeIR` (`source/slang/slang-emit.cpp`) runs many backend IR passes. Some are gated behind `if (requiredLoweringPassSet.<flag>)`, but many run **unconditionally** as full-module walks even when the IR they act on is absent — wasted compile time with no IR change. Issue shader-slang#11917 asks to bring the `RequiredLoweringPassSet` mechanism "up to date" so such passes are skipped when they cannot apply. This is an **incremental epic**, not a one-shot change: mechanically batch-gating all remaining unconditional passes is a correctness hazard. A flag that is *stale-FALSE* (feature present, flag not set) makes a needed pass skip → **miscompile**. The autodiff case (shader-slang#11474 / open PR shader-slang#11476) is the canonical trap — `DifferentialPair`/`no_diff` appear without `fwd_diff`/`bwd_diff`, so its predicate must be a *safe superset* of everything the passes strip. This PR is the **first, deliberately-small A slice**: it gates exactly one pass whose trigger is *structurally* immune to stale-FALSE, so the direction can be reviewed before expanding. Consider a compute shader that uses no append/consume buffers (the overwhelmingly common case): ```hlsl RWStructuredBuffer<int> inBuf; RWStructuredBuffer<int> outBuf; [numthreads(4,1,1)] void computeMain(uint i : SV_GroupIndex) { outBuf[i] = inBuf[i] * 2; } ``` For every non-HLSL target, `lowerAppendConsumeStructuredBuffers` still walks all global insts of this module looking for `HLSLAppendStructuredBufferType` / `HLSLConsumeStructuredBufferType`, finds none, and does nothing. ## Proposed solution Approach **A — incremental per-pass gating** (recommended): add a flag to `RequiredLoweringPassSet`, detect the triggering op(s) in `calcRequiredLoweringPassSet` as a *safe superset* of what the pass consumes, and wrap the pass call in `if (flag)`. This is the pattern established by PR shader-slang#11476. For pass #1, `lowerAppendConsumeStructuredBuffers` is chosen because a stale-FALSE flag is **structurally impossible**: - Its whole effect keys on two IR types, `HLSLAppendStructuredBufferType` and `HLSLConsumeStructuredBufferType` (see `slang-ir-lower-append-consume-structured-buffer.cpp` — it iterates `module->getGlobalInsts()` and only acts on those two ops; when neither is present the pass is a pure no-op). - Those types are produced by the front-end (from user `AppendStructuredBuffer<T>` / `ConsumeStructuredBuffer<T>` declarations). A source search found **no IR-builder construction** of `kIROp_HLSLAppend/ConsumeStructuredBufferType` anywhere in `source/`; every other reference (SPIR-V legalization name hints, buffer-element-type layout `switch`, emitters, reflection, ir-util) is read-only and op-preserving — i.e. no IR pass synthesizes them. - `calcRequiredLoweringPassSet` flags **accumulate** across the post-link and post-specialization scans (they are not reset between them). Since these types are never synthesized by an IR pass, none is created after the last (post-specialization) scan — so any instance present at the gate was recorded by a scan and set the flag, and the flag can never be a **false-negative** (skip a needed lowering). It can only be stale-*true* (e.g. an unused buffer dead-code-eliminated after a scan), which costs a harmless no-op walk. Therefore gating is a behavior-preserving skip: the pass fires whenever it would have done work, and when it runs on a module with no such type (whether via the gate or a stale-true flag) it is a no-op producing identical output. ### Alternatives considered - **B — fix staleness (rescan / incremental maintenance):** rescan before gate clusters that follow IR-generating passes (each rescan is itself a full-module walk → partly self-defeating), or have mutating passes report added/removed features (correct + fast, but invasive across many pass signatures). Needed only for passes whose triggers *are* synthesized late — not this one. - **C — shared opcode-presence index / cheap per-pass early-out:** one bitset per module, O(1) queries, refreshed at defined points. Localizes correctness to each pass but still has a refresh-cadence question. **Recommendation: A**, one profiled pass at a time — and, for the first PRs, restricted to passes where stale-FALSE is structurally impossible (as here). @pdeayton-nv / maintainers: this PR is intended to validate that direction before the next slices. Deferring B/C to passes whose triggers are actually synthesized after the last scan. **On the byte-identical evidence standard (maintainer input welcome).** A committed CI test builds exactly one compiler, so it cannot compare gated-vs-ungated output; and a behavior-preserving skip is byte-indistinguishable from a no-op run in emitted code — that indistinguishability *is* the correctness claim, so there is nothing for a single-build test to detect. "Byte-identical vs ungated" is therefore inherently the two-build revert-drill above, not a single-build regression. The committed tests instead cover the CI-observable failure mode: the feature-present test fails if the gate ever breaks in the dangerous stale-FALSE direction (flag wrongly false → pass skipped → buffer unlowered → emitter error). If maintainers prefer a stronger committed proxy than this pair + the documented revert-drill, please say so on this draft. ## Change summary | File | Change | |---|---| | `source/slang/slang-code-gen.h` | Add `bool appendConsumeStructuredBuffer;` to `RequiredLoweringPassSet` (appended; internal struct, no ABI surface). | | `source/slang/slang-emit.cpp` — `calcRequiredLoweringPassSet` | Set the flag on `kIROp_HLSLAppendStructuredBufferType` / `kIROp_HLSLConsumeStructuredBufferType` (detection is exactly the ops the pass consumes). | | `source/slang/slang-emit.cpp` — gate site | `if (target != CodeGenTarget::HLSL && requiredLoweringPassSet.appendConsumeStructuredBuffer)`, with a comment explaining the false-negative-safety. | | `tests/hlsl/append-consume-buffer-lowering-gated.slang` | Feature-present regression: asserts the pass fires (types lowered to synthesized accessor functions) on `-target glsl`. | | `tests/hlsl/append-consume-buffer-lowering-skip.slang` | Feature-absent baseline: asserts a plain-buffer module emits no Append/Consume accessor functions. | ## Concepts and vocabulary - **`RequiredLoweringPassSet`** — a set of bool flags recording which optional lowering passes a module needs, populated by `calcRequiredLoweringPassSet` (a whole-module opcode walk) and read in `linkAndOptimizeIR`. - **Scan points** — `calcRequiredLoweringPassSet` runs twice: once post-link (which first resets the set), once post-specialization (which accumulates onto it, no reset). The gate added here reads the accumulated flag. - **Stale-FALSE vs stale-TRUE** — stale-FALSE = a flag left `false` while its feature is present, which would wrongly skip a needed pass (a miscompile); shown impossible for this flag. Stale-TRUE = flag `true` after the feature was removed, which only runs the pass as a no-op (benign). ## Process report **Why the change, and why this layer.** The pass is unconditional today; the only reason to run it on a module with no append/consume buffer is the absence of a gate. Adding the gate at the pass call site (not inside the pass) matches every other gated pass and keeps the pass body unchanged. **Input-shape check.** The triggering shapes are `HLSLAppendStructuredBufferType` / `HLSLConsumeStructuredBufferType` global type insts. These are *correct and principled* front-end output (they model the user's `AppendStructuredBuffer<T>` / `ConsumeStructuredBuffer<T>` resource declarations) — there is no accidental alternative spelling and no producer to fix upstream. The detection predicate equals the set of ops the pass consumes, so it is a safe superset by construction. No new helper, fallback, or graph-walk is introduced. **Correctness evidence (revert-drill, byte-identical).** Built Debug, then emitted a feature-present shader (Append + Consume + GetDimensions) and a feature-absent shader on `-target glsl`, `-target cpp`, and `-target spirv-asm`, with the gate in place; reverted only the gate condition to unconditional, rebuilt, and re-emitted. **All outputs were byte-identical** in both directions — the gate fires and produces identical output when the types are present, and is a no-op skip when they are absent. (`-target cpp` on the present shader fails identically with and without the gate — a pre-existing CPU-emit limitation, unrelated to this change.) To reproduce: Two scratch shaders — `present.slang` (uses Append/Consume) and `absent.slang` (neither): ```hlsl // present.slang RWStructuredBuffer<int> inBuf; AppendStructuredBuffer<int> appendBuf; ConsumeStructuredBuffer<int> consumeBuf; RWStructuredBuffer<int> outBuf; [numthreads(4,1,1)] void computeMain(uint i : SV_GroupIndex) { appendBuf.Append(inBuf[i]); uint n, s; appendBuf.GetDimensions(n, s); outBuf[i] = consumeBuf.Consume() + int(n); } // absent.slang RWStructuredBuffer<int> inBuf; RWStructuredBuffer<int> outBuf; [numthreads(4,1,1)] void computeMain(uint i : SV_GroupIndex) { outBuf[i] = inBuf[i] * 2; } ``` ```sh # Gated build = this branch at commit 8541c38. git checkout 8541c38 && cmake --workflow --preset debug mkdir -p gated ungated for sh in present absent; do for t in glsl cpp spirv-asm; do build/Debug/bin/slangc $sh.slang -target $t -entry computeMain -stage compute \ -o gated/$sh.$t 2>gated/$sh.$t.err; done; done # Ungated baseline = revert ONLY the gate condition to unconditional, rebuild slangc: sed -i 's/if (target != CodeGenTarget::HLSL && requiredLoweringPassSet.appendConsumeStructuredBuffer)/if (target != CodeGenTarget::HLSL)/' source/slang/slang-emit.cpp cmake --build --preset debug --target slangc for sh in present absent; do for t in glsl cpp spirv-asm; do build/Debug/bin/slangc $sh.slang -target $t -entry computeMain -stage compute \ -o ungated/$sh.$t 2>ungated/$sh.$t.err; done; done diff -r gated ungated # empty: byte-identical (present.cpp fails identically both ways — pre-existing CPU limit) ``` **Tests.** Both new tests pass locally (slang-test + bundled FileCheck). The feature-present test is the false-negative guard: if detection were too narrow / the gate inverted, the un-lowered Append type would reach the GLSL emitter and error, so the CHECK would fail. The feature-absent test is a baseline: a plain-buffer module emits no Append/Consume accessor functions. **Byte-identical gated-vs-ungated output is verified by a manual revert-drill, not by a committed test** — the skip is behavior-preserving, so no single-build output test can distinguish "skipped" from "ran as a no-op"; proving equivalence inherently requires building both variants and diffing (done here on `-target glsl`/`cpp`/`spirv-asm` for present and absent shaders — all byte-identical). Existing `tests/hlsl/append-structured-buffer.slang` and `consume-structured-buffer.slang` still pass on `-vk` and `-cuda` (no regression; runtime compute correctness for Append/Consume remains covered there). `Addresses shader-slang#11917` (non-closing — this is one slice of the epic). <sub>🤖 Generated by an automated Slang coworker — may be inaccurate. A human maintainer should verify.</sub> --------- Co-authored-by: nv-slang-bot[bot] <274397474+nv-slang-bot[bot]@users.noreply.github.com>
stramit
pushed a commit
that referenced
this pull request
Aug 21, 2026
…r-slang#11957) Fixes shader-slang#11956. ## Motivation `prelude/slang-cuda-prelude.h` provides 1-wide vector "make" helpers for the extended float types. Consider generated CUDA code constructing a 1-wide bfloat16 vector: ```cpp __nv_bfloat161 v = make___nv_bfloat161(x); ``` In the non-NVRTC path, `make___nv_bfloat161` is declared to return the *scalar* `__nv_bfloat16` (and returns `__nv_bfloat16{x}`), so nvcc rejects this with "no suitable user-defined conversion from `__nv_bfloat16` to `__nv_bfloat161`". The two FP8 helpers (`make___nv_fp8_e4m31`, `make___nv_fp8_e5m21`) carry the identical defect. A downstream nvcc-based build (slangtorch-generated kernels) hit this and had to patch the header locally. ## Proposed solution Make the three helpers return the `1`-suffixed vector struct, constructed as `T1{x}` — which is what the codebase already defines as the contract in two places: - the NVRTC branch of `SLANG_MAKE_VECTOR_FROM_SCALAR` generates `T##1 make_##T##1(T x) { return T##1{x}; }` for every type it covers, and - the adjacent non-NVRTC `make___half1` implements exactly that form for `__half`. The non-NVRTC helpers exist separately from the macro because the macro's non-RTC branch omits the 1-wide form (for the standard CUDA vector types, the SDK's own `make_*` functions occupy those names); each extended type therefore supplies its explicit `make_T1` under `#if !SLANG_CUDA_RTC`. The bf16/FP8 ones diverged from the pattern by copy-paste. The fix aligns them with the existing contract — no new representation, no behavior change for NVRTC builds (which already got the correct macro-generated definitions). ## Change summary | File | Change | | --- | --- | | `prelude/slang-cuda-prelude.h` | `make___nv_bfloat161`, `make___nv_fp8_e4m31`, and `make___nv_fp8_e5m21` return their `1`-suffixed vector structs (`T1{x}`), matching `make___half1` and the NVRTC macro. | | `tests/cuda/cuda-prelude-vec1-make.cu` | Compile-only regression fixture that assigns all four extended-type helper results (the three fixed ones plus the known-good `make___half1`) to their corresponding wrapper structs under offline nvcc. | | `.github/workflows/ci-slang-test.yml` | Compiles the fixture directly with nvcc on Windows CUDA test runners. | ## Concepts and vocabulary - **NVRTC vs nvcc prelude paths** — the prelude compiles under both NVRTC (runtime compilation, `SLANG_CUDA_RTC`) and offline nvcc. The 1-wide make helpers have per-path definitions: macro-generated under NVRTC, explicit functions otherwise. - **`T1` structs** — Slang-defined single-member aggregates (`struct __nv_bfloat161 { __nv_bfloat16 x; };`) standing in for the 1-wide vector types the CUDA SDK does not provide. ## Process report - **Input-shape check**: the correct producer-side contract already exists in this same header (macro RTC branch + `make___half1`); the three fixed helpers are accidental divergences from it, not an alternative intended form. The fix eliminates the divergent spelling rather than teaching any consumer to tolerate it. - **Regression coverage**: `tests/cuda/cuda-prelude-vec1-make.cu` enables half, bf16, and FP8 support, asserts that it is compiling the non-NVRTC path, and assigns each helper result to its `T1` output type. The old scalar-return signatures fail those assignments. The fixture deliberately has no slang-test directive because slang-test's CUDA pass-through uses NVRTC and would exercise the already-correct macro branch. Instead, the test workflow invokes `nvcc -std=c++17 -c` directly on the `full-gpu-tests` Windows lanes, whose self-hosted runners carry the CUDA toolkit; the step is gated on that input (not bare `os == 'windows'`) so CPU-only Windows lanes added later cannot break on a missing nvcc. - The three edits are mechanical and symmetric; both FP8 helpers are included (not just the reported bf16 one) because they are the same copy-paste instance and fixing one while leaving the twins would preserve a known-broken shape. **Suggested reviewers**: @haaggarwal-nvidia (most recent prelude authorship, incl. the bf16/FP8 additions this fixes), @kaizhangNV (CUDA target maintenance, second-most prelude changes).
stramit
pushed a commit
that referenced
this pull request
Aug 21, 2026
…er-slang#9382) (shader-slang#12133) ## Motivation A `Gather` with a **compile-time-constant** offset over-declares a SPIR-V capability. Consider: ```hlsl Texture2D<float> tex; SamplerState samp; float4 main(float2 uv) : SV_Target { return tex.Gather(samp, uv, int2(2, 1)); // constant offset } ``` Compiling to SPIR-V emits the `Offset` image operand (`ImageOperandsOffsetMask`, `0x10`) on the `OpImageGather`, which requires `OpCapability ImageGatherExtended`. Per the SPIR-V spec a *constant* offset should use `ConstOffset` (`0x08`), which requires **no** capability. The needless capability can break consumers/devices that do not advertise `ImageGatherExtended`. The over-declaration was baked into the core module: the `__texture_gather_offset` intrinsics emitted their SPIR-V from a `spirv_asm` block that hard-coded both `OpCapability ImageGatherExtended;` and `... Offset $offset`, regardless of the offset's constness. At `-O0` this is worse than a spare capability — `spirv-opt` only rewrites `Offset`→`ConstOffset` for a constant at `-O1+`, so at `-O0` a constant offset stays `Offset` and needs the capability. Simply dropping the hard-coded capability while still authoring `Offset` would therefore emit `Offset`-without-capability at `-O0` = invalid SPIR-V. The constant case genuinely needs the `ConstOffset` **operand**, chosen at emit time. ## Proposed solution Slang does not iterate instructions inside a `spirv_asm` block, and cannot overload a function on `constexpr`, so the block itself cannot branch on constness. Following maintainer direction (shader-slang#9382), the principled fix moves the gather-offset SPIR-V emission out of the frontend `spirv_asm` block and onto the backend as a new IR op, where the offset's constness is available: - A new IR op `kIROp_ImageGatherOffset` carries `{sampledImage, location, component, offset}`. - Its SPIR-V emitter inspects the offset operand: - **constant** offset → `OpImageGather ... ConstOffset`, **no** capability. - **runtime** offset → `OpImageGather ... Offset` + `requireSPIRVCapability(ImageGatherExtended)`. The runtime path must stay reachable: GLSL's `textureGatherOffset` uniquely permits a *variable* offset and routes through this same intrinsic (merged PR shader-slang#5426 / shader-slang#5339 deliberately switched to `Offset` for exactly this). Forcing `ConstOffset` on a runtime value is invalid SPIR-V ("Expected Image Operand ConstOffset to be a const object", the shader-slang#5339 CTS failure), so the branch is load-bearing. This supersedes the two earlier stale/conflicting drafts on this issue — **shader-slang#9741** (jkwak-work; IR-legalization pass iterating the `spirv_asm` block) and **shader-slang#9410** (copilot; emit-time recursive check). It does not close shader-slang#9741 (a maintainer's own draft). ## Change summary | File | Change | | --- | --- | | `source/slang/slang-ir-insts.lua` | Add IR op `imageGatherOffset` (operands `sampledImage, location, component, offset`). | | `source/slang/slang-ir-insts-stable-names.lua` | Append stable id `898` for the new op. | | `source/slang/hlsl.meta.slang` | New `__intrinsic_op($(kIROp_ImageGatherOffset))` helper `__spirvImageGatherOffset`; route both singular `__texture_gather_offset` `spirv` cases to it, dropping the two hard-coded `OpCapability ImageGatherExtended;` + `Offset $offset` `spirv_asm` blocks. | | `source/slang/slang-emit-spirv.cpp` | `emitImageGatherOffset` (assembles `OpImageGather` with the chosen image-operand mask) + `isConstantGatherOffset` classifier. | | `tests/bugs/shader-slanggh-9382-gather-constoffset.slang` | Regression test: constant, constant-splat, and runtime offsets at `-O0`. | ## Concepts and vocabulary - **`ConstOffset` vs `Offset` image operands** — SPIR-V `OpImageGather` optional operands. `ConstOffset` (`0x08`) requires the offset be a constant object and needs no capability; `Offset` (`0x10`) allows a runtime offset but requires `OpCapability ImageGatherExtended`. - **`__intrinsic_op($(kIROp_X))`** — a `.meta.slang` function-level attribute that lowers the whole function to IR op `X`; the call arguments become the op's operands, via the generic `emitIntrinsicInst` path in `slang-lower-to-ir.cpp` (no bespoke `IRBuilder` method needed — mirrors `kIROp_MakeCombinedTextureSampler`). - **Combined sampled image** — the `OpSampledImage` value produced by `__makeCombinedTextureSampler` (a `_Texture<...,isCombined=1,...>`), fed as the image operand of `OpImageGather`. ## Process report **New IR op `kIROp_ImageGatherOffset` (lua + stable name).** Adding a backend IR op is the sanctioned way (CLAUDE.md) to move a decision that a `spirv_asm` block cannot make onto the backend. The op has no flags, matching its sibling texture ops (`imageLoad`/`imageStore`/`imageSubscript`/`sample`), and is only ever produced by the SPIR-V `case` of the two singular `__texture_gather_offset` defs. It needs no manual `IRInst` subclass or `IRBuilder` method: `emitCallToDeclRef`'s intrinsic-op path (`slang-lower-to-ir.cpp`) maps the call's arguments to operands generically, exactly as `kIROp_MakeCombinedTextureSampler` does. Clone/serialize are generic (operand-count loop + stable-name indirection); `mightHaveSideEffects` defaults to conservative-true for an unlisted op, so it is not wrongly DCE'd or reordered. The stable id `898` is appended (previous max `897`); the mid-file lua position shifts raw `kIROp_` integers for later ops but that is not an ABI concern — serialized IR encodes opcodes by stable name, decoupled from lua order (e.g. `imageStore`=256 but `ImageTexelPointer`=862 despite being lua-adjacent). No `include/` public enum is touched. **Frontend routing (`hlsl.meta.slang`).** The two singular `__texture_gather_offset` defs differ only in how they obtain the combined sampled image: def #1 takes a separate texture+sampler and builds one via `__makeCombinedTextureSampler(texture, s)` (kept outside any `spirv_asm` block, as before, so `NonUniform` still propagates); def #2's `sampler` parameter is already `isCombined=1`. Both now call `__spirvImageGatherOffset(sampledImage, location, component, offset)`, whose parameter order the emitter reads positionally. The removed `spirv_asm` blocks' hard-coded `OpCapability ImageGatherExtended;` lines are gone with them. The plural `__texture_gather_offsets` (all-`constexpr` offsets → `ConstOffsets`, which legitimately requires the capability) and the shadow `gatherCmp` intrinsics are untouched. **Constness classifier — is the input shape principled?** `isConstantGatherOffset` accepts an `IRConstant` leaf, or a `MakeVector` / `MakeVectorFromScalar` whose every operand is an `IRConstant`. This is exactly the flat integer-vector shape a gather offset takes, and such an inst is `@ConstExpr`, hoisted to the `ConstantsAndTypes` section, where `emitCompositeConstruct` renders it as `OpConstantComposite` — a valid `ConstOffset` operand by construction. Iterating operands directly is correct for these two ops specifically because their constructed IR has no leading type operand (the result type is the inst's data type, not an operand): `emitMakeVectorFromScalar` builds a single scalar operand, `emitMakeVector` builds element-value operands. The classifier deliberately does **not** broaden to `MakeArray`/`Struct`/`Matrix`/nested exprs — a broad, un-hoisted composite could be a runtime `OpCompositeConstruct`, and emitting `ConstOffset` on a non-constant id is invalid SPIR-V. A misclassification is only ever safe in the `Offset` direction (an unproven-constant offset falls back to `Offset` + capability, which is valid), never toward an invalid `ConstOffset`. **Emitter (`slang-emit-spirv.cpp`).** `emitImageGatherOffset` mirrors `emitImageLoad`: it uses `emitInstCustomOperandFunc` to assemble `OpImageGather` with the SPIR-V operand order `result-type, result-id, sampled-image, coordinate, component, image-operands-mask, offset`. The mask is `ConstOffset` (constant) or `Offset` (runtime); only the runtime branch calls `requireSPIRVCapability(SpvCapabilityImageGatherExtended)`. **Verification.** Built `slangc`/`slang-test` (Debug). All three test variants pass, and — because a `SIMPLE(filecheck=...)` test is silently skipped when FileCheck is absent — the SPIR-V was also confirmed directly with the built `slangc`: - `int2(2,1)` → `OpImageGather ... ConstOffset`, **0** `ImageGatherExtended` capabilities. - `int2(1)` (splat) → `OpImageGather ... ConstOffset`, **0** capabilities. - runtime `int2` → `OpImageGather ... Offset` + **1** `ImageGatherExtended` capability. All three validate under `SLANG_RUN_SPIRV_VALIDATION=1 -target spirv -O0`. The runtime-offset regression guard `tests/bugs/shader-slanggh-5339.slang` still passes, as do the broader gather tests (`glsl-texturegather`, `texture-2d-gather`, `gather-texture2darray`, `wgsl/texture-gather`). `NonUniformResourceIndex` propagation was verified end-to-end: the decoration reaches the `OpSampledImage` feeding the gather — the IR-op path is strictly better than the removed `spirv_asm` form, whose own comment warned that `NonUniform` does not propagate to asm-local ids. **Out of scope.** `__texture_gatherCmp_offset` (singular) emits `ConstOffset` unconditionally today, a latent mirror bug for a *runtime* comparison-gather offset. It is a separate defect, not what shader-slang#9382 reports; noting it here as a follow-up rather than expanding scope. Closes shader-slang#9382. --------- Co-authored-by: nv-slang-bot[bot] <274397474+nv-slang-bot[bot]@users.noreply.github.com> Co-authored-by: Jay Kwak <82421531+jkwak-work@users.noreply.github.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Superseded by #2 (rebranched off
fix/byteaddressbufferload-specialize-alignment).