diff --git a/CHANGELOG.md b/CHANGELOG.md index c780923..1a23e7c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -7,14 +7,22 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ## [Unreleased] +## [0.17.16] - 2026-07-26 + ### Fixed (MSL) -- **64-bit atomic capability validation** ([#79](https://github.com/gogpu/naga/issues/79)) — +- **64-bit atomic capability validation** ([#79](https://github.com/gogpu/naga/issues/79), PR #82, @besmpl) — `msl.Compile` now returns a descriptive error for 64-bit atomic operations that Metal cannot represent instead of emitting invalid intrinsics. Matching Rust Naga, result-discarded `min`/`max` in storage address space remain supported; loads, stores, full operations, result-producing min/max, and workgroup min/max are rejected. +- **MSL: swizzle of binary expression missing parentheses** ([#83](https://github.com/gogpu/naga/issues/83)) — + `writeSwizzle` now wraps binary/select sub-expressions in parentheses, + matching Rust naga's `is_scoped=false` semantics. Without this, C++ member + access (`.`) bound tighter than arithmetic operators, causing + `mat * vec.xyz` instead of `(mat * vec).xyz` — a type error that Metal + compiler rejected. Discovered via g3d `standard.wgsl` shader on Apple M5. ## [0.17.15] - 2026-06-15 diff --git a/ROADMAP.md b/ROADMAP.md index 6eb2e01..7ae67e0 100644 --- a/ROADMAP.md +++ b/ROADMAP.md @@ -2,7 +2,7 @@ > **Pure Go Shader Compiler — WGSL to SPIR-V, MSL, GLSL, HLSL, and DXIL. Zero CGO.** > -> Current: **v0.17.15** (June 2026) · Target: **v1.0.0** (December 2026) +> Current: **v0.17.16** (July 2026) · Target: **v1.0.0** (December 2026) --- @@ -32,7 +32,7 @@ Our goal: **the most complete, most tested, most portable shader compiler availa --- -## Where We Are (v0.17.15) +## Where We Are (v0.17.16) **~323K LOC, 6 backends, 100% Rust parity on all text backends, 172/172 SPIR-V validation.** @@ -189,6 +189,7 @@ Our goal: **the most complete, most tested, most portable shader compiler availa | Version | Date | Highlights | |---------|------|------------| +| **v0.17.16** | 2026-07 | MSL 64-bit atomic validation (PR #82, @besmpl), swizzle parenthesization fix (#83) | | **v0.17.15** | 2026-06 | MSL function-scope workgroup vars (PR #77, @georgebuilds), per-EP zero-init filtering | | **v0.17.14** | 2026-06 | GLSL version-aware binding, UniformInfo reflection, Codecov OIDC (BUG-GLES-005) | | **v0.17.13** | 2026-05 | DXIL PHI node ordering fix, coverage waves 3-4, ~60% overall | diff --git a/msl/internal/codegen/expressions.go b/msl/internal/codegen/expressions.go index 83f79f7..4a886e6 100644 --- a/msl/internal/codegen/expressions.go +++ b/msl/internal/codegen/expressions.go @@ -1095,9 +1095,20 @@ func (w *Writer) writeSwizzle(swizzle ir.ExprSwizzle) error { w.write("%sfloat3(", Namespace) } } + // Wrap in parens if vector is a binary/select expression. + // C++ member access (.) binds tighter than arithmetic operators, + // so (mat * vec).xyz must not be emitted as mat * vec.xyz. + // Matches Rust naga is_scoped=false in put_expression → put_binop. + needParens := !needsUnpack && w.needsParensInContext(swizzle.Vector) + if needParens { + w.write("(") + } if err := w.writeExpression(swizzle.Vector); err != nil { return err } + if needParens { + w.write(")") + } if needsUnpack { w.write(")") } diff --git a/msl/internal/codegen/int64_atomics_test.go b/msl/internal/codegen/int64_atomics_test.go index 0edd167..93c402c 100644 --- a/msl/internal/codegen/int64_atomics_test.go +++ b/msl/internal/codegen/int64_atomics_test.go @@ -227,6 +227,16 @@ var value: atomic; @compute @workgroup_size(1) fn main() { atomicMax(&value, 1lu); } +`, + }, + { + name: "signed add", + operation: "add", + source: ` +@group(0) @binding(0) var value: atomic; +@compute @workgroup_size(1) fn main() { + atomicAdd(&value, 1li); +} `, }, } diff --git a/msl/internal/codegen/msl_integration5_test.go b/msl/internal/codegen/msl_integration5_test.go index 72e505b..1c27c84 100644 --- a/msl/internal/codegen/msl_integration5_test.go +++ b/msl/internal/codegen/msl_integration5_test.go @@ -581,6 +581,66 @@ fn main() { mustContainMSL(t, code, ".yx") } +func TestIntegration_SwizzleOfBinaryNeedsParens(t *testing.T) { + src := ` +struct Uniforms { model: mat4x4 }; +@group(0) @binding(0) var u: Uniforms; +struct In { pos: vec4 }; +@group(0) @binding(1) var inp: In; +struct Out { v: vec3 }; +@group(0) @binding(2) var out: Out; +@compute @workgroup_size(1) +fn main() { + out.v = (u.model * inp.pos).xyz; +} +` + code := compileWGSL(t, src) + if strings.Contains(code, "* ") && strings.Contains(code, ".xyz") { + lines := strings.Split(code, "\n") + for _, line := range lines { + if strings.Contains(line, ".xyz") && strings.Contains(line, "*") { + if !strings.Contains(line, ").xyz") { + t.Errorf("swizzle of binary missing parens — C++ . binds tighter than *:\n%s", line) + } + } + } + } + mustContainMSL(t, code, ").xyz") + mustNotContainMSL(t, code, "* metal::float4(") +} + +func TestIntegration_SwizzleOfAddNeedsParens(t *testing.T) { + src := ` +struct In { a: vec4, b: vec4 }; +@group(0) @binding(0) var inp: In; +struct Out { v: vec2 }; +@group(0) @binding(1) var out: Out; +@compute @workgroup_size(1) +fn main() { + out.v = (inp.a + inp.b).xy; +} +` + code := compileWGSL(t, src) + mustContainMSL(t, code, ").xy") +} + +func TestIntegration_SwizzleOfVariableNoParens(t *testing.T) { + src := ` +struct In { v: vec4 }; +@group(0) @binding(0) var inp: In; +struct Out { v: vec3 }; +@group(0) @binding(1) var out: Out; +@compute @workgroup_size(1) +fn main() { + let tmp = inp.v; + out.v = tmp.xyz; +} +` + code := compileWGSL(t, src) + mustContainMSL(t, code, ".xyz") + mustNotContainMSL(t, code, ").xyz") +} + // ============================================================================= // Test: Modf and frexp math (covers writeModf, writeFrexp) // ============================================================================= diff --git a/snapshot/testdata/golden/msl/7048-multiple-dynamic-2.msl b/snapshot/testdata/golden/msl/7048-multiple-dynamic-2.msl index d0a6e56..1843fc6 100644 --- a/snapshot/testdata/golden/msl/7048-multiple-dynamic-2.msl +++ b/snapshot/testdata/golden/msl/7048-multiple-dynamic-2.msl @@ -25,5 +25,5 @@ fragment fs_mainOutput fs_main( metal::float2 val_0_ = uint(_e9) < 2 ? my_array.inner[_e9] : DefaultConstructible(); int _e11 = index_0_; metal::float2 val_1_ = uint(_e11) < 2 ? my_array.inner[_e11] : DefaultConstructible(); - return fs_mainOutput { val_0_ * val_1_.xxyy }; + return fs_mainOutput { (val_0_ * val_1_).xxyy }; }