Skip to content

Commit 2a8e5cf

Browse files
kolkovkolkov
authored
fix: SPIR-V compute shader bug fixes (v0.14.3)
* fix: SPIR-V deferred store with multiple call results in init expression findCallResultInTree returned only the first ExprCallResult in an expression tree. When a local var init contained two different call results (e.g. span(a,b) + span(c,d)), the deferred store triggered on the first call — before the second result was cached. Renamed to findLastCallResultInTree: explores all branches and returns the highest-handle CallResult. Since StmtCalls emit in handle order, the last call result is the last to complete, ensuring all earlier results are already in callResultIDs. Fixes path_count.wgsl compilation failure in Vello compute pipeline. * fix: SPIR-V deferred store for var x = atomicOp() initialization ExprAtomicResult was not handled in the local variable init loop, causing 'atomic result expression not found' when compiling patterns like var seg_ix = atomicAdd(&bump.counter, 1u). Fix: treat ExprAtomicResult the same as ExprCallResult — defer the store until emitAtomic caches the result, then execute it via processDeferredStores. Also added ExprAtomicResult to the renamed findLastDeferredResultInTree for complex init expressions. Added 7 Vello workaround validation tests confirming all 6 previously discovered SPIR-V codegen bugs are now fixed. * fix: SPIR-V OpLogicalEqual for bool comparisons, transitive deferred stores - Bool == / != now emits OpLogicalEqual/OpLogicalNotEqual (was OpIEqual) - Bool & / | now emits OpLogicalAnd/OpLogicalOr (was OpBitwiseAnd/Or) - Transitive deferred store: var X = Y correctly deferred when Y depends on a function call result (3-pass: classify, propagate, emit) - Fixes span() function call producing wrong results in compute shaders * fix: SPIR-V atomic result type for atomic<i32> struct fields - resolveAtomicScalarKind now handles AtomicType returned directly by ResolveExpressionType (not wrapped in PointerType) - Fixes spirv-val error: AtomicIAdd expected Pointer to point to value of type Result Type - Affects struct field atomics: tiles[i].backdrop (atomic<i32>) * fix: split var init to StmtStore when referencing local variables (NAGA-SPV-007) Frontend init splitting in WGSL lowerer: var inits that reference local variables are emitted as LocalVariable{Init: nil} + StmtStore at declaration position instead of being pre-computed in prologue. This prevents stale values when a var's init reads a local variable that is modified by preceding control flow (if/else blocks). Approach matches Rust naga (src/front/wgsl/lower/mod.rs:1737-1801). New: initReferencesLocalVariable() expression tree walker. Tests: 3 regression tests + 1 span() diagnostic test. * docs: update changelog for v0.14.3 * test: use t.TempDir() for SPIR-V test output files * test: skip spirv-val when not installed in CI
1 parent 7d94d37 commit 2a8e5cf

File tree

8 files changed

+1865
-47
lines changed

8 files changed

+1865
-47
lines changed

‎CHANGELOG.md‎

Lines changed: 15 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,17 @@ All notable changes to this project will be documented in this file.
55
The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.1.0/),
66
and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html).
77

8+
## [0.14.3] - 2026-02-25
9+
10+
### Fixed
11+
12+
#### SPIR-V Backend
13+
- **Deferred store for multiple call results** — Variables initialized from expressions containing multiple function call results now correctly emit deferred `OpStore` instructions for each intermediate result
14+
- **Deferred store for `var x = atomicOp()`** — Atomic operation results used in variable initialization now correctly generate deferred stores instead of losing the value (NAGA-SPV-006)
15+
- **`OpLogicalEqual` for bool comparisons** — Boolean equality expressions now emit correct `OpLogicalEqual` opcode; transitive deferred stores propagate through boolean comparison chains
16+
- **Atomic result type for `atomic<i32>` struct fields** — Atomic operations on signed integer struct members now use correct `OpTypeInt 32 1` result type instead of unsigned
17+
- **Prologue var init splitting** — Variable initializations that reference other local variables are now split from the function prologue into `StmtStore` at the declaration point, preventing use-before-definition in SPIR-V (NAGA-SPV-007)
18+
819
## [0.14.2] - 2026-02-22
920

1021
### Added
@@ -840,7 +851,10 @@ First stable release. Complete WGSL to SPIR-V compilation pipeline (~10K LOC).
840851

841852
---
842853

843-
[Unreleased]: https://github.com/gogpu/naga/compare/v0.13.1...HEAD
854+
[Unreleased]: https://github.com/gogpu/naga/compare/v0.14.3...HEAD
855+
[0.14.3]: https://github.com/gogpu/naga/compare/v0.14.2...v0.14.3
856+
[0.14.2]: https://github.com/gogpu/naga/compare/v0.14.1...v0.14.2
857+
[0.14.1]: https://github.com/gogpu/naga/compare/v0.14.0...v0.14.1
844858
[0.14.0]: https://github.com/gogpu/naga/compare/v0.13.1...v0.14.0
845859
[0.13.1]: https://github.com/gogpu/naga/compare/v0.13.0...v0.13.1
846860
[0.13.0]: https://github.com/gogpu/naga/compare/v0.12.1...v0.13.0

‎spirv/atomic_type_test.go‎

Lines changed: 116 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,116 @@
1+
package spirv
2+
3+
import (
4+
"testing"
5+
)
6+
7+
// TestAtomicI32ResultType verifies that atomicAdd on atomic<i32> emits
8+
// OpAtomicIAdd with int result type (not uint). SPIR-V spec requires the
9+
// result type to match the pointed-to scalar type.
10+
// Bug: NAGA-SPV-009 — resolveAtomicScalarKind returned ScalarUint for struct
11+
// field access (e.g., tiles[i].backdrop) because ResolveExpressionType returns
12+
// AtomicType directly, not wrapped in PointerType.
13+
func TestAtomicI32ResultType(t *testing.T) {
14+
const shader = `
15+
struct Tile {
16+
backdrop: atomic<i32>,
17+
seg_count: atomic<u32>,
18+
}
19+
20+
@group(0) @binding(0) var<storage, read_write> tiles: array<Tile>;
21+
22+
@compute @workgroup_size(1)
23+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
24+
let idx = gid.x;
25+
atomicAdd(&tiles[idx].backdrop, 1i);
26+
atomicAdd(&tiles[idx].seg_count, 1u);
27+
}
28+
`
29+
spirvBytes := compileWGSLToSPIRV(t, "AtomicI32", shader)
30+
instrs := decodeSPIRVInstructions(spirvBytes)
31+
32+
// Build type ID → name map from OpTypeInt declarations
33+
typeMap := make(map[uint32]string)
34+
for _, inst := range instrs {
35+
if inst.opcode == OpTypeInt && len(inst.words) >= 4 {
36+
id, width, signed := inst.words[1], inst.words[2], inst.words[3]
37+
if width == 32 && signed == 1 {
38+
typeMap[id] = "int"
39+
} else if width == 32 && signed == 0 {
40+
typeMap[id] = "uint"
41+
}
42+
}
43+
}
44+
45+
// Verify each OpAtomicIAdd has correct result type:
46+
// - atomic<i32> → int result type
47+
// - atomic<u32> → uint result type
48+
want := []string{"int", "uint"}
49+
atomicIdx := 0
50+
for _, inst := range instrs {
51+
if inst.opcode != OpAtomicIAdd || len(inst.words) < 3 {
52+
continue
53+
}
54+
resultTypeID := inst.words[1]
55+
got := typeMap[resultTypeID]
56+
if atomicIdx < len(want) {
57+
if got != want[atomicIdx] {
58+
t.Errorf("OpAtomicIAdd[%d]: result type = %q (ID %d), want %q",
59+
atomicIdx, got, resultTypeID, want[atomicIdx])
60+
}
61+
}
62+
atomicIdx++
63+
}
64+
if atomicIdx != len(want) {
65+
t.Errorf("found %d OpAtomicIAdd instructions, want %d", atomicIdx, len(want))
66+
}
67+
}
68+
69+
// TestAtomicI32PointerType verifies that the pointer operand of OpAtomicIAdd
70+
// points to the same type as the result type (int for i32, uint for u32).
71+
func TestAtomicI32PointerType(t *testing.T) {
72+
const shader = `
73+
struct Counter {
74+
signed_val: atomic<i32>,
75+
unsigned_val: atomic<u32>,
76+
}
77+
78+
@group(0) @binding(0) var<storage, read_write> counters: array<Counter>;
79+
80+
@compute @workgroup_size(1)
81+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
82+
atomicAdd(&counters[gid.x].signed_val, 1i);
83+
atomicAdd(&counters[gid.x].unsigned_val, 1u);
84+
}
85+
`
86+
spirvBytes := compileWGSLToSPIRV(t, "AtomicI32Pointer", shader)
87+
instrs := decodeSPIRVInstructions(spirvBytes)
88+
89+
// Build pointer type chain: OpTypePointer → base type, OpTypeInt → signed/unsigned
90+
typeNames := make(map[uint32]string) // type ID → "int" or "uint"
91+
for _, inst := range instrs {
92+
if inst.opcode == OpTypeInt && len(inst.words) >= 4 {
93+
id, width, signed := inst.words[1], inst.words[2], inst.words[3]
94+
if width == 32 && signed == 1 {
95+
typeNames[id] = "int"
96+
} else if width == 32 && signed == 0 {
97+
typeNames[id] = "uint"
98+
}
99+
}
100+
}
101+
102+
// For each OpAtomicIAdd, verify pointer base type matches result type
103+
atomicIdx := 0
104+
for _, inst := range instrs {
105+
if inst.opcode != OpAtomicIAdd || len(inst.words) < 4 {
106+
continue
107+
}
108+
resultTypeID := inst.words[1]
109+
resultTypeName := typeNames[resultTypeID]
110+
t.Logf("OpAtomicIAdd[%d]: result type ID=%d (%s)", atomicIdx, resultTypeID, resultTypeName)
111+
atomicIdx++
112+
}
113+
if atomicIdx != 2 {
114+
t.Errorf("found %d OpAtomicIAdd, want 2", atomicIdx)
115+
}
116+
}

0 commit comments

Comments
 (0)