WGSL: emit runtime-indexable static const arrays as var<private>#11628
Merged
Conversation
… in WGSL
A module-scope `static const` array (or matrix) lowered for WGSL was emitted as a
WGSL `const`. A WGSL `const` is a compile-time value that may only be indexed by a
const-expression, so indexing one by a runtime value (e.g. `positions[SV_VertexID]`)
is rejected by the WGSL validator ("The expression may only be indexed by a
constant"). Emit such module-scope constants as `var<private>` with their inline
const-expression initializer instead: a private module-scope variable holds the same
value but, being addressable, is runtime-indexable. Scalar/vector constants and
function-local constants are unaffected.
Co-authored-by: Harsh Aggarwal <haaggarwal@nvidia.com>
…on to arrays Address review of the static-const-array -> var<private> change: - Nested aggregates: a nested `static const` (e.g. `int g[2][3]`) lowered each inner MakeArray to a separate named module-scope decl, so converting the outer array to `var<private>` produced an initializer that referenced inner `var<private>`s -- illegal WGSL. WGSL's shouldFoldInstIntoUseSites now folds a module-scope MakeArray/MakeStruct/MakeArrayFromElement inline when it is only a nested constituent of another aggregate (all uses are aggregate-constructor operands), so the outermost runtime-indexed array stays a decl while its constituents inline into a self-contained const-expression initializer. - Narrow the conversion to array types only; a WGSL matrix value is itself runtime-indexable, so a `const` matrix needs no conversion. - Add a nested-aggregate regression test and a wgsl-spirv-asm (Tint) validation directive to the array tests; tighten the const-index test's anchors and add a negative guard; document the const vs var<private> split in the user guide. Co-authored-by: Harsh Aggarwal <haaggarwal@nvidia.com>
A reviewer-identified edge case: an aggregate that is both a nested constituent and used directly (a named `static const` array shared as an element of another `static const` array and also independently runtime-indexed) is not folded inline, so it stays a separate `var<private>` declaration that the enclosing converted array's initializer references by name -- invalid WGSL. The shape is not constructible from typical anonymous nested literals and is not a regression (such a constant was rejected before this change too, as a runtime-indexed `const`). Record it where the fold decision is made. Co-authored-by: Harsh Aggarwal <haaggarwal@nvidia.com>
Round-2 review follow-up (all non-blocking):
- Reconcile the matrix wording: matrices stay `const` because a matrix value,
like a vector value, is runtime-indexable in WGSL (only array values are
restricted to const-expression indices). Add tests/wgsl/static-const-matrix.slang
(runtime-indexed `static const` matrix stays an inline value, not `var<private>`)
with a wgsl-spirv-asm/Tint directive that validates it on CI.
- Correct the predicate comment: only the `GlobalParam` exclusion is load-bearing
(a module-scope GlobalParam array would otherwise get `<private>` in the
address-space chain); the GlobalVar/Var terms are defensive.
- Guard `getParent()` in the predicate, matching the sibling fold-site check.
- Tighten the const-index test's negative to `var<private>{{.*}}weights` so it is
bound to the array rather than a file-global `var<private>` assertion.
Co-authored-by: Harsh Aggarwal <haaggarwal@nvidia.com>
jkwak-work
reviewed
Jun 16, 2026
jkwak-work
left a comment
Collaborator
There was a problem hiding this comment.
Comments are too verbose overall
Address maintainer review on the static-const-array -> var<private> change: use the existing isStaticConst() helper for the module-scope array-constant predicate (keeping only the load-bearing != kIROp_GlobalParam guard so a descriptor/GlobalParam array keeps its own address space rather than <private>), inline the single-use aggregate-constituent helper into shouldFoldInstIntoUseSites, and trim the comments. Behavior is unchanged; tests/wgsl/ pass 49/49. Co-authored-by: Harsh Aggarwal <haaggarwal@nvidia.com>
jvepsalainen-nv
pushed a commit
that referenced
this pull request
Jun 22, 2026
) ## Motivation When targeting WGSL, a module-scope `static const` array is emitted as a WGSL `const`. But a WGSL `const` is a compile-time value (closer to C++ `constexpr` than to `const`): the WGSL spec permits a *value* of array type to be indexed only by a const-expression. So indexing such a global by a runtime value is rejected by the WGSL validator. Motivating case (from the issue): ```hlsl static const float2 positions[] = { float2(-1.0f, -1.0f), float2(1.0f, -1.0f), float2(1.0f, 1.0f), float2(-1.0f, 1.0f) }; [shader("vertex")] VertexOutput vertexMain(uint vertexID : SV_VertexID) { VertexOutput o; o.position = float4(positions[vertexID], 0.0f, 1.0f); // runtime index return o; } ``` emits `const positions_0 : array<vec2<f32>, i32(4)> = ...;` then `positions_0[vertexID_0]`, which naga/tint reject: *"The expression may only be indexed by a constant."* ## Proposed solution Emit module-scope `static const` globals of **array** type as WGSL `var<private>` (a private module-scope variable), keeping their inline const-expression initializer, instead of `const`. A `var<private>` holds the same value but is addressable, so it is runtime-indexable. After the fix the array emits `var<private> positions_0 : array<vec2<f32>, i32(4)> = array<vec2<f32>, i32(4)>( ... );` and `positions_0[vertexID_0]` is valid. Only arrays are converted. A WGSL matrix *value* is itself runtime-indexable, so a `const` matrix needs no conversion; scalars/vectors stay `const` (a vector value is also dynamically indexable), and function-local constants are unaffected. Nested aggregates need one more step. A nested `static const` (e.g. `int grid[2][3]`) lowers to nested `MakeArray` insts, each a module-scope array constant emitted as a *separate named* declaration. Converting all of them to `var<private>` would make the outer initializer reference inner `var<private>`s — illegal WGSL (a `var<private>` initializer must be a const-expression and cannot read another `var`). The fix folds the inner constituents inline so the outermost runtime-indexed array becomes a single `var<private>` with a self-contained initializer: ```wgsl var<private> grid_0 : array<array<i32, i32(3)>, i32(2)> = array<array<i32, i32(3)>, i32(2)>( array<i32, i32(3)>( i32(1), i32(2), i32(3) ), array<i32, i32(3)>( i32(4), i32(5), i32(6) ) ); ``` Alternatives ruled out: a `static`-only (non-`const`) global lowers its initializer into an init function, producing the broken `const _S3 = ...;` statements seen in the issue thread; keeping the initializer inline avoids that. A precise "is this global ever dynamically indexed" use-analysis is unnecessary — constant-indexed reads are already constant-folded before emit (see Process report). ## Change summary | File | Change | | --- | --- | | `source/slang/slang-emit-wgsl.cpp` | `emitVarKeywordImpl`: a module-scope **array** constant emits `var` + `<private>` instead of `const`. `shouldFoldInstIntoUseSites`: for WGSL, fold a module-scope `MakeArray`/`MakeStruct`/`MakeArrayFromElement` inline when it is only a nested constituent of another aggregate, so the outermost array's initializer stays self-contained. | | `tests/wgsl/static-const-array-indexing.slang` | Runtime-indexed array → `var<private>`; `+ wgsl-spirv-asm` (Tint) validation directive. | | `tests/wgsl/static-const-array-nested.slang` | New: nested 2D-array regression — one inline `var<private>`, no separate inner-array decls; `+ wgsl-spirv-asm` directive. | | `tests/wgsl/static-const-matrix.slang` | New: a runtime-indexed `static const` matrix stays an inline `const` value (not `var<private>`); `+ wgsl-spirv-asm` directive validates it through Tint on CI. | | `tests/wgsl/static-const-array-const-index.slang` | A constant-indexed element folds away and leaves no `const`/`var` reading another variable; tightened anchors + negative guard. | | `docs/user-guide/a2-03-wgsl-target-specific.md` | Document the module-scope `const` (scalar/vector/matrix) vs `var<private>` (array) split. | ## Concepts and vocabulary - **WGSL `const` vs `var<private>`** — a WGSL `const` is a compile-time value; a `const` value of array type may only be indexed by a const-expression. A `var<private>` is an addressable module-scope variable, indexable by a runtime value, whose initializer must still be a const-expression. - **`GlobalConstant` / `MakeArray`** — a `static const` array lowers to an `IRGlobalConstant` wrapping an `IRMakeArray`; nested arrays nest `MakeArray`s. - **`replaceGlobalConstants` / peephole fold** — `replaceGlobalConstants` inlines a global constant's value into its uses; the peephole pass folds `GetElement(MakeArray, constIndex)` to the element, leaving a runtime index intact. - **`shouldFoldInstIntoUseSites`** — decides whether an inst is emitted inline at its use or as its own declaration. WGSL can fold `MakeArray`/`MakeStruct` (constructor expressions, valid in any expression context) where the base class cannot (C/HLSL initializer lists). ## Process report **`emitVarKeywordImpl` (`source/slang/slang-emit-wgsl.cpp`).** The declared inst for a module-scope `static const` array is an ordinary module-scope instruction; `emitInstResultDecl` (`slang-emit-c-like.cpp`) emits `<keyword> <name> : <type> = <initializer>`, with the keyword from `emitVarKeywordImpl`. Previously the `default:` arm emitted `const` for any module-scope constant. The predicate `emitModuleScopeArrayConstAsPrivateVar` reuses the existing `isStaticConst` helper and is true when the inst is a module-scope constant (`isStaticConst(varDecl)`), is not a `GlobalParam`, and has type `kIROp_ArrayType`; it then emits `var` and `<private>` (joining the existing `kIROp_GlobalVar` storage-space branch, after the array-of-ConstantBuffer/structured-buffer branches, so resource arrays are unaffected). The `!= kIROp_GlobalParam` guard is load-bearing: the predicate is also read in the address-space chain (which runs for all ops), where a module-scope `GlobalParam` array — e.g. a descriptor array like `Texture2D t[8]`, whose element type isn't ConstantBuffer/structured so it falls through to the final branch — would otherwise wrongly get `<private>` instead of its handle address space. (`isStaticConst` is true for any module child, including `GlobalParam`/`GlobalVar`, so the guard is required; a local `Var` is not module-scope and is already excluded.) **Why a type-based conversion is safe (no const-expression chain breaks).** A constant index into a static-const array is constant-folded before emit, so no `const` is left whose initializer reads the converted `var`: `replaceGlobalConstants` inlines `GlobalConstant` values into uses before `simplifyIR`, then the peephole `kIROp_GetElement` case folds `GetElement(MakeArray, IRIntLit)` to the element while a non-constant (runtime) index hits `if (!index) break;` and is left intact. `tests/wgsl/static-const-array-const-index.slang` guards this. **Nested aggregates — `shouldFoldInstIntoUseSites`.** The base class never-folds `MakeArray`/`MakeStruct`/`MakeArrayFromElement` because in C/HLSL they lower to initializer lists, invalid in a general expression context. WGSL instead emits them as constructor expressions (`type(args...)`, `tryEmitInstExprImpl`), valid anywhere. So a nested `static const` otherwise emits each inner `MakeArray` as a separate named module-scope decl, and converting them all to `var<private>` yields an outer initializer that references inner `var<private>`s — illegal. The override folds a module-scope constituent inline when every use of it is an operand of another `MakeArray`/`MakeStruct`/`MakeArrayFromElement` (checked by a small inlined loop), so the outermost aggregate — used directly, e.g. by a runtime `GetElement` — stays a declaration eligible for `var<private>` while its constituents inline. Function-local aggregates are left to the base policy; resource/global-parameter arrays keep their existing declaration paths. `tests/wgsl/static-const-array-nested.slang` covers a runtime-indexed 2D array. *Known limitation:* an aggregate that is both a nested constituent and used directly — a named `static const` array shared as an element of another `static const` array *and* independently runtime-indexed — is not folded, so it stays a separate `var<private>` declaration that the enclosing converted array's initializer references by name (invalid WGSL). This is not constructible from typical anonymous nested literals, and is not a regression for accepted WGSL output (such a constant was a rejected runtime-indexed `const` before this change too, though the failure now takes a different form); fully handling it would require inlining a duplicate copy of the constituent into the enclosing initializer. **Input-shape check (right layer?).** Yes. The IR for a `static const` array is correct (`IRGlobalConstant` wrapping `IRMakeArray`); the defects are WGSL keyword selection and WGSL fold policy, both target-specific. The fix belongs in the WGSL emitter. **Verification.** Built a debug `slangc`: the flat and nested cases emit single self-contained `var<private>` declarations indexed by name; a `static const` matrix is emitted as an inline matrix value indexed in place (`mat2x2<…>(…)[i]`), not promoted to `var<private>`. The full `tests/wgsl/` suite passes (49/49). Note: `slangc` emitting WGSL is not the same as a validator accepting it — Tint is not available in the build environment, so the `wgsl-spirv-asm` directives (which round-trip the emitted WGSL through Tint) are ignored locally and run on CI. The matrix test's directive is what would confirm or refute that a runtime-indexed `const` matrix is valid WGSL. **Known limitation — arrays nested inside a `static const` struct.** A `static const` struct with an array field, runtime-indexed via `g.arr[i]`, is *not* addressed: the struct is not `ArrayType`, so it stays `const`, and it lowers through a constructor function (`g.arr[i]` → `Foo_init(...).arr[i]`) whose result is an array value runtime-indexed — the same #6747 rejection. This is a realistic lookup-table shape but a distinct (struct-valued) problem from the top-level arrays #6747 reports, and is left as a follow-up. Fixes #6747. --------- Co-authored-by: nv-slang-bot[bot] <274397474+nv-slang-bot[bot]@users.noreply.github.com> Co-authored-by: Harsh Aggarwal <haaggarwal@nvidia.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
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.
Motivation
When targeting WGSL, a module-scope
static constarray is emitted as a WGSLconst. But a WGSLconstis a compile-time value (closer to C++constexprthan toconst): the WGSL spec permits a value of array type to be indexed only by a const-expression. So indexing such a global by a runtime value is rejected by the WGSL validator.Motivating case (from the issue):
emits
const positions_0 : array<vec2<f32>, i32(4)> = ...;thenpositions_0[vertexID_0], which naga/tint reject: "The expression may only be indexed by a constant."Proposed solution
Emit module-scope
static constglobals of array type as WGSLvar<private>(a private module-scope variable), keeping their inline const-expression initializer, instead ofconst. Avar<private>holds the same value but is addressable, so it is runtime-indexable. After the fix the array emitsvar<private> positions_0 : array<vec2<f32>, i32(4)> = array<vec2<f32>, i32(4)>( ... );andpositions_0[vertexID_0]is valid.Only arrays are converted. A WGSL matrix value is itself runtime-indexable, so a
constmatrix needs no conversion; scalars/vectors stayconst(a vector value is also dynamically indexable), and function-local constants are unaffected.Nested aggregates need one more step. A nested
static const(e.g.int grid[2][3]) lowers to nestedMakeArrayinsts, each a module-scope array constant emitted as a separate named declaration. Converting all of them tovar<private>would make the outer initializer reference innervar<private>s — illegal WGSL (avar<private>initializer must be a const-expression and cannot read anothervar). The fix folds the inner constituents inline so the outermost runtime-indexed array becomes a singlevar<private>with a self-contained initializer:Alternatives ruled out: a
static-only (non-const) global lowers its initializer into an init function, producing the brokenconst _S3 = ...;statements seen in the issue thread; keeping the initializer inline avoids that. A precise "is this global ever dynamically indexed" use-analysis is unnecessary — constant-indexed reads are already constant-folded before emit (see Process report).Change summary
source/slang/slang-emit-wgsl.cppemitVarKeywordImpl: a module-scope array constant emitsvar+<private>instead ofconst.shouldFoldInstIntoUseSites: for WGSL, fold a module-scopeMakeArray/MakeStruct/MakeArrayFromElementinline when it is only a nested constituent of another aggregate, so the outermost array's initializer stays self-contained.tests/wgsl/static-const-array-indexing.slangvar<private>;+ wgsl-spirv-asm(Tint) validation directive.tests/wgsl/static-const-array-nested.slangvar<private>, no separate inner-array decls;+ wgsl-spirv-asmdirective.tests/wgsl/static-const-matrix.slangstatic constmatrix stays an inlineconstvalue (notvar<private>);+ wgsl-spirv-asmdirective validates it through Tint on CI.tests/wgsl/static-const-array-const-index.slangconst/varreading another variable; tightened anchors + negative guard.docs/user-guide/a2-03-wgsl-target-specific.mdconst(scalar/vector/matrix) vsvar<private>(array) split.Concepts and vocabulary
constvsvar<private>— a WGSLconstis a compile-time value; aconstvalue of array type may only be indexed by a const-expression. Avar<private>is an addressable module-scope variable, indexable by a runtime value, whose initializer must still be a const-expression.GlobalConstant/MakeArray— astatic constarray lowers to anIRGlobalConstantwrapping anIRMakeArray; nested arrays nestMakeArrays.replaceGlobalConstants/ peephole fold —replaceGlobalConstantsinlines a global constant's value into its uses; the peephole pass foldsGetElement(MakeArray, constIndex)to the element, leaving a runtime index intact.shouldFoldInstIntoUseSites— decides whether an inst is emitted inline at its use or as its own declaration. WGSL can foldMakeArray/MakeStruct(constructor expressions, valid in any expression context) where the base class cannot (C/HLSL initializer lists).Process report
emitVarKeywordImpl(source/slang/slang-emit-wgsl.cpp). The declared inst for a module-scopestatic constarray is an ordinary module-scope instruction;emitInstResultDecl(slang-emit-c-like.cpp) emits<keyword> <name> : <type> = <initializer>, with the keyword fromemitVarKeywordImpl. Previously thedefault:arm emittedconstfor any module-scope constant. The predicateemitModuleScopeArrayConstAsPrivateVarreuses the existingisStaticConsthelper and is true when the inst is a module-scope constant (isStaticConst(varDecl)), is not aGlobalParam, and has typekIROp_ArrayType; it then emitsvarand<private>(joining the existingkIROp_GlobalVarstorage-space branch, after the array-of-ConstantBuffer/structured-buffer branches, so resource arrays are unaffected). The!= kIROp_GlobalParamguard is load-bearing: the predicate is also read in the address-space chain (which runs for all ops), where a module-scopeGlobalParamarray — e.g. a descriptor array likeTexture2D t[8], whose element type isn't ConstantBuffer/structured so it falls through to the final branch — would otherwise wrongly get<private>instead of its handle address space. (isStaticConstis true for any module child, includingGlobalParam/GlobalVar, so the guard is required; a localVaris not module-scope and is already excluded.)Why a type-based conversion is safe (no const-expression chain breaks). A constant index into a static-const array is constant-folded before emit, so no
constis left whose initializer reads the convertedvar:replaceGlobalConstantsinlinesGlobalConstantvalues into uses beforesimplifyIR, then the peepholekIROp_GetElementcase foldsGetElement(MakeArray, IRIntLit)to the element while a non-constant (runtime) index hitsif (!index) break;and is left intact.tests/wgsl/static-const-array-const-index.slangguards this.Nested aggregates —
shouldFoldInstIntoUseSites. The base class never-foldsMakeArray/MakeStruct/MakeArrayFromElementbecause in C/HLSL they lower to initializer lists, invalid in a general expression context. WGSL instead emits them as constructor expressions (type(args...),tryEmitInstExprImpl), valid anywhere. So a nestedstatic constotherwise emits each innerMakeArrayas a separate named module-scope decl, and converting them all tovar<private>yields an outer initializer that references innervar<private>s — illegal. The override folds a module-scope constituent inline when every use of it is an operand of anotherMakeArray/MakeStruct/MakeArrayFromElement(checked by a small inlined loop), so the outermost aggregate — used directly, e.g. by a runtimeGetElement— stays a declaration eligible forvar<private>while its constituents inline. Function-local aggregates are left to the base policy; resource/global-parameter arrays keep their existing declaration paths.tests/wgsl/static-const-array-nested.slangcovers a runtime-indexed 2D array.Known limitation: an aggregate that is both a nested constituent and used directly — a named
static constarray shared as an element of anotherstatic constarray and independently runtime-indexed — is not folded, so it stays a separatevar<private>declaration that the enclosing converted array's initializer references by name (invalid WGSL). This is not constructible from typical anonymous nested literals, and is not a regression for accepted WGSL output (such a constant was a rejected runtime-indexedconstbefore this change too, though the failure now takes a different form); fully handling it would require inlining a duplicate copy of the constituent into the enclosing initializer.Input-shape check (right layer?). Yes. The IR for a
static constarray is correct (IRGlobalConstantwrappingIRMakeArray); the defects are WGSL keyword selection and WGSL fold policy, both target-specific. The fix belongs in the WGSL emitter.Verification. Built a debug
slangc: the flat and nested cases emit single self-containedvar<private>declarations indexed by name; astatic constmatrix is emitted as an inline matrix value indexed in place (mat2x2<…>(…)[i]), not promoted tovar<private>. The fulltests/wgsl/suite passes (49/49). Note:slangcemitting WGSL is not the same as a validator accepting it — Tint is not available in the build environment, so thewgsl-spirv-asmdirectives (which round-trip the emitted WGSL through Tint) are ignored locally and run on CI. The matrix test's directive is what would confirm or refute that a runtime-indexedconstmatrix is valid WGSL.Known limitation — arrays nested inside a
static conststruct. Astatic conststruct with an array field, runtime-indexed viag.arr[i], is not addressed: the struct is notArrayType, so it staysconst, and it lowers through a constructor function (g.arr[i]→Foo_init(...).arr[i]) whose result is an array value runtime-indexed — the same #6747 rejection. This is a realistic lookup-table shape but a distinct (struct-valued) problem from the top-level arrays #6747 reports, and is left as a follow-up.Fixes #6747.