release: v0.17.16

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
 

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 |

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(")")
 	}

msl/internal/codegen/int64_atomics_test.go

@@ -227,6 +227,16 @@ var<workgroup> value: atomic<u64>;
 @compute @workgroup_size(1) fn main() {
     atomicMax(&value, 1lu);
 }
+`,
+		},
+		{
+			name:      "signed add",
+			operation: "add",
+			source: `
+@group(0) @binding(0) var<storage, read_write> value: atomic<i64>;
+@compute @workgroup_size(1) fn main() {
+    atomicAdd(&value, 1li);
+}
 `,
 		},
 	}

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<f32> };
+@group(0) @binding(0) var<uniform> u: Uniforms;
+struct In { pos: vec4<f32> };
+@group(0) @binding(1) var<storage, read> inp: In;
+struct Out { v: vec3<f32> };
+@group(0) @binding(2) var<storage, read_write> 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<f32>, b: vec4<f32> };
+@group(0) @binding(0) var<storage, read> inp: In;
+struct Out { v: vec2<f32> };
+@group(0) @binding(1) var<storage, read_write> 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<f32> };
+@group(0) @binding(0) var<storage, read> inp: In;
+struct Out { v: vec3<f32> };
+@group(0) @binding(1) var<storage, read_write> 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)
 // =============================================================================

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 };
 }