Promote small non-escaping fixed local arrays after specialization and loop unrolling

[tdavidovicNV]

Problem Description

After generic specialization, inlining, and [ForceUnroll], a small fixed-size
local array can have only constant-index element accesses while remaining in
addressable aggregate form in Slang IR and generated target code.

This prevents individual elements from becoming ordinary SSA values, which in
turn limits per-element constant propagation, dead-code elimination, common
subexpression elimination, and lifetime shortening. It also makes the result
depend on whether each downstream compiler can reconstruct the lost scalar
form from an addressable array.

The requested optimization is the standard, narrow form of scalar replacement
of aggregates (SROA): promote small, non-escaping fixed local arrays, or array
fields of local aggregates, when every remaining access chain has a constant
index after the existing specialization, inlining, and loop-unrolling passes.

Minimal Reproducer

Save as repro.slang:

StructuredBuffer<float3> inputData;
RWStructuredBuffer<float3> outputData;

struct InlineArray<let N : int>
{
    float3 data[N];
}

[ForceInline]
InlineArray<N> prepare<let N : int>(uint baseIndex)
{
    InlineArray<N> result;
    [ForceUnroll]
    for (int i = 0; i < N; ++i)
        result.data[i] = inputData[baseIndex + i];
    return result;
}

[ForceInline]
float3 consume<let N : int>(InlineArray<N> values)
{
    float3 result = 0.0;
    [ForceUnroll]
    for (int i = 0; i < N; ++i)
        result += values.data[i] * float(i + 1);
    return result;
}

[shader("compute")]
[numthreads(1, 1, 1)]
void computeMain(uint3 tid : SV_DispatchThreadID)
{
    let values = prepare<11>(tid.x * 11);
    outputData[tid.x] = consume(values);
}

Compile with:

slangc repro.slang -O2 -target cuda -entry computeMain -stage compute -o repro.cu
slangc repro.slang -O2 -target spirv-asm -entry computeMain -stage compute -o repro.spvasm
slangc repro.slang -O2 -target ptx -entry computeMain -stage compute -o repro.ptx

With the v2026.16 release compiler, the optimized CUDA output retains the
132-byte InlineArray local and all eleven stores:

struct InlineArray_0
{
    FixedArray<float3, 11> data_0;
};

InlineArray_0 result_0;
float3 _S6 = slang_ldg_0(...);
(&result_0)->data_0[int(0)] = _S6;
// ...ten more literal-index stores...

The final reduction does not read result_0; it directly uses _S6 through
_S16. Thus the complete local and its stores are dead, but Slang's optimized
CUDA output still contains them.

The optimized Slang IR immediately before CUDA emission likewise retains:

%result = var InlineArray<float3, 11>
%data = get_field_addr %result, data
%element0 = getElementPtr %data, 0
store %element0, %loaded0
// ...literal indices 1 through 10...

In contrast, direct SPIR-V contains no function-local array, and Slang's PTX
target (through NVRTC) contains no .local, ld.local, st.local, or calls.
The small case is therefore recovered downstream, but Slang emits target- and
program-size-dependent dead addressable work after it already has enough
information to remove it.

Expected Behavior

After specialization and forced unrolling, the eleven accesses are all at
literal indices. The local array or its independently accessed elements should
be eligible for SSA promotion, subject to a conservative size/profitability
limit.

An IR regression test should be able to verify that an eligible local no
longer contains address-based partial element stores (getElementPtr +
store) after the relevant optimization stage.

Proposed Initial Scope

Only consider storage where all of the following hold:

• The array size is statically known and small.
• The storage is function-local and does not escape.
• Every remaining access chain has a constant index.
• Specialization, inlining, constant propagation, and requested loop unrolling
  have already exposed those indices.
• The element type can safely participate in ordinary SSA construction.

Explicit non-goals for an initial implementation:

• Do not transform genuinely dynamic array[index] operations into switch or
  select trees.
• Do not unroll additional loops merely to create an SROA opportunity.
• Do not scalarize resource, uniform, shared, global, ABI-visible, or escaping
  arrays.
• Do not scalarize large arrays or promise that promotion is always profitable
  under register pressure.
• Do not add any MaterialX-specific transformation.

Why Downstream Optimization Is Not Reliable

A directly authored small array is often easy for NVRTC to recover. Generated
programs can present a substantially harder analysis problem even when their
closure or slot identities were statically known at the source level.

For example, a generated MaterialX CUDA translation reaches NVRTC in this
shape (names shortened, but the representation is unchanged):

struct MxStackData
{
    FixedArray<float3, 11> bsdf_weights;
    FixedArray<float3, 11> bsdf_transmission_scales;
    FixedArray<float3, 10> layer_weights;
    FixedArray<float3, 10> layer_transmission_scales;
    FixedArray<float3, 10> mix_values;
    FixedArray<float3, 11> bsdf_n;
    FixedArray<float3, 11> bsdf_t;
    FixedArray<float, 11> transmits;
    // Eleven concrete BSDF fields follow.
};

__device__ void set_bsdf_weight(
    MxStackData* data, int bsdf_index, float3 weight)
{
    data->bsdf_weights[bsdf_index] = weight;
}

// The slot identity originated from generated graph structure, but is
// transported through a generic BSDF value before the indexed store.
BSDF closure = generated_bsdf_node(...);
set_bsdf_weight(&stack, closure.index, weight);

The generic consumer later constructs and copies more array-valued state:

FixedArray<float3, 11> wi_local;
// Eleven literal-index writes after Slang force-unrolling.

FixedArray<float3, 11> arg0 = wi_local;
FixedArray<float3, 11> arg1 = arg0;
MxWeights weights = calculate_weights(&material, &arg1, front_facing, context);

Recovering independent scalars now requires a successful chain of inlining,
alias analysis, propagation through index-carrying aggregates, constant
propagation, and SROA. If any step fails or is rejected by the downstream
inliner's cost model, the complete array remains addressable. In the captured
vanilla sample PTX, calculate_weights also remains out of line and returns a
144-byte aggregate twice.

There is also a later, genuinely dynamic operation:

float weight = selection.bsdf_weight[sampled_lobe].x;

That array is deliberately outside this proposal. The candidate addressed by
this issue is the earlier state whose indices become constant after Slang's
existing specialization and forced unrolling.

MaterialX Evidence

An isolated generator experiment replaced only generated array-backed stack
storage with named fields while retaining the generic handwritten material
consumer, the same eleven BSDFs, the same ten combiners, and zero closure
pruning.

With CUDA 13.3 targeting compute_120:

Operation	Vanilla ptxas stack	Named-field ptxas stack	Runtime change
sample	4,624 B	1,768 B	-3.2%
eval	4,496 B	2,088 B	-10.3%

Kernel spills were removed in both operations. NVRTC-produced PTX, before
ptxas, also reduced the primary local depots from 3,712 to 1,680 bytes for
sample and from 3,680 to 1,640 bytes for eval.

The same change was nearly neutral on Vulkan and D3D12, which is useful scope
information: this is not claimed to be a universal runtime win. It is evidence
that source/IR aggregate representation can materially determine whether the
CUDA downstream compiler recovers scalar state in a large specialized program.

The experiment did not produce the much larger performance gain of a fully
specialized material implementation. SROA is one bounded missing optimization,
not a request for the compiler to rediscover an entire MaterialX-specific
schedule.

Broader Relevance

This pattern is used by Slang's own neural standard module. For example,
InlineVector<T, N> stores T[N] data, and its constructors and operations
populate the array using [ForceUnroll] loops:

slang/source/standard-modules/neural/inline-vector.slang

Lines 3 to 43 in 400ef93

	/**
	Concrete implementation of IVector storing elements inline (on stack/registers).
	InlineVector stores all elements in a fixed-size array, making it suitable for
	small vectors that can fit in registers or stack memory. Supports automatic differentiation
	for gradient computation in neural networks.
	@param T The element type
	@param N The vector size (compile-time constant).
	@remarks Type constraints:
	- `T` must conform to `__BuiltinFloatingPointType` (float, double, half, etc.)
	- `T.Differential` must conform to `__BuiltinFloatingPointType` for automatic differentiation
	@category neural
	*/
	public struct InlineVector<T, int N> : IVector<T>
	where T : __BuiltinFloatingPointType
	where T.Differential == T
	{
	/// The differential type for automatic differentiation.
	public typealias Differential = InlineVector<T, N>;
	/// The compile-time size of the vector.
	public static const int Size = N;
	public int getCount() {return N;}
	/**
	Internal storage for vector elements.
	@remarks Marked as derivative member to enable automatic differentiation.
	*/
	// [DerivativeMember(Differential.data)]
	internal T[N] data;
	/// Default constructor - initializes all elements to zero.
	public __init() { data = {}; }
	/**
	Scalar broadcast constructor - fills all elements with the same value.
	@param[in] value The value to broadcast to all elements.
	*/
	public __init(T value) {
	[ForceUnroll]
	for (int i = 0; i < N; i++)
	data[i] = value;

An audit of the production neural standard module found 34 fixed-array
declarations, of which 25 are plausible local/register aggregate candidates.
Similar patterns occur in generic and autodiff tests, cooperative-matrix code,
and generated derivative code. Not every occurrence satisfies the proposed
constant-index/non-escaping restrictions, but the programming idiom is not
specific to MaterialX.

As a downstream control, NVRTC completely scalarizes a small InlineVector
test. That reinforces the intended claim: the optimization opportunity is
standard and sometimes recovered downstream, but larger generated programs and
other targets should not have to depend on reconstructing it late.

Existing Compiler Machinery

There appears to be a narrow implementation path already present in Slang:

• constructSSA documents that it does not handle partial assignments through
  getElementPtr/field addresses.
• eliminateAddressInsts rewrites partial address-based assignments into
  aggregate updates so SSA construction can process them.
• eliminateAddressInsts is currently called only from forward-autodiff
  preparation.
• Buffer-element lowering already uses a 32-element threshold when choosing
  direct array construction instead of an addressable temporary, providing a
  precedent for a conservative expansion limit.

This is only an implementation lead. Enabling the existing transformation
globally without candidate analysis and profitability controls may cause code
expansion or register-pressure regressions.

Related Issues That Do Not Cover This Case

• Slang fails to output a static variable in form that is optimized by downstream compiler #4722 addressed static globals and SPIR-V global-variable scalar
  replacement, not function-local partially stored fixed arrays.
• Metal AD synthesis materializes large backward-context copies that NVCC optimizes away #11085 tracks reverse-AD context copying, primarily on Metal; its CUDA
  control is successfully optimized by NVCC.
• [CUDA] runtime-indexed resource arrays are ~10-30x slower than they need to be (kept in .param instead of __constant__) #11774 concerns runtime-indexed resource arrays and CUDA parameter
  ABI/routing, and explicitly identifies routing rather than a missing IR pass
  as its cause.
• CUDA: fixed-width vector operator overloads use dynamic-index loops #12401/Spell CUDA vector components directly in fixed-width binary and comparison operators #12410 changed fixed-width CUDA vector operators to spell constant
  components directly. That is useful precedent, but deliberately leaves
  genuine dynamic indexing and general local-array promotion out of scope.

Environment

• Slang source inspected through 2026-08-26.
• MaterialX CUDA evidence: CUDA 13.3, compute_120, Windows.
• The MaterialX comparison retained all eleven authored BSDF occurrences and
  performed no zero-weight closure pruning.

[nv-slang-bot]

Triage update — nv-slang-bot

Design review complete. This remains a well-scoped, actionable feature; no code has been written yet because the mechanism is gated on a small set of maintainer design decisions (below). Research firsthand against top-of-tree strengthened the plan on three points.

• Status: Triaged → planned. Classification unchanged: feature-request (narrow SROA) · medium · P2 · IR (SSA construction / optimization pipeline) · Issue Type: Feature. Recommended path Approach A — gated reuse of eliminateAddressInsts' value-rewrite + the existing constructSSA/peephole pipeline — confirmed as the right first cut (no new SROA pass needed).
• Link: Three findings that de-risk Approach A, all verified in-source: (1) the scalarization tail is already free — source/slang/slang-ir-peephole.cpp:957-1018 folds a constant-index updateElement over makeArray/makeArrayFromElement (and whole updateElement chains) straight to scalars, so your acceptance criterion (no residual getElementPtr+store) is met with no new folding code; (2) the non-escaping analysis already exists, inlined in isPromotableVar (slang-ir-ssa.cpp:478-573), so the gate mirrors it rather than adding an escape pass; (3) only partial stores are the gap — reads-through-chains are already promoted generally, so "constant-index" is a profitability/scope gate, not a correctness requirement.
• Verdict: The real work is the two gating tasks already identified — a conservative candidate filter, and decoupling the rewrite from its autodiff-only diagnostic (slang-ir-autodiff-fwd.cpp:2433 is currently eliminateAddressInsts' only caller) so general use skips-not-errors. No correctness blockers found.
• Next-action: Four design-checkpoint questions for maintainers before implementation is finalized: (a) profitability threshold + target-gating — the 32-element precedent (kMaxArraySizeToUnroll) as the cap, and whether to gate to CUDA/targets that don't recover downstream (Vulkan/D3D12 were near-neutral in the reporter's data); (b) the shape of the autodiff decouple; (c) pipeline placement of the gated rewrite (post-specialization / post-force-inline, before a constructSSA run); (d) arrays-only first, or also array-typed struct fields. We can spike the mechanism (repro test + whole-function dry-run) first if you'd prefer to de-risk before deciding.
• Blocker: Maintainer decision on the four questions above. Nothing else outstanding; ready to implement once the scope/threshold is set.

_{🤖 Generated by an automated Slang coworker — may be inaccurate. A human maintainer should verify.}