fix(msl): use references (&) instead of pointers (*) for buffer parameters

Buffer parameters now generate 'constant T& name' instead of 'constant T* name'.
Pointer syntax required '->' for member access while expression writer generates
dot access, causing Metal compilation errors on Apple Silicon (ui#23).

Author: kolkov

CHANGELOG.md

@@ -5,6 +5,13 @@ All notable changes to this project will be documented in this file.
 The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.1.0/),
 and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html).
 
+## [0.14.5] - 2026-03-04
+
+### Fixed
+
+#### MSL Backend
+- **Buffer parameters use references (`&`) instead of pointers (`*`)** — buffer parameters now generate `constant Uniforms& u [[buffer(0)]]` (reference) instead of `constant Uniforms* u [[buffer(0)]]` (pointer); pointer syntax required `->` or `(*u).` for member access while the expression writer generates `.` access, causing Metal compilation errors on Apple Silicon ([gogpu/ui#23](https://github.com/gogpu/ui/issues/23))
+
 ## [0.14.4] - 2026-03-01
 
 ### Fixed

msl/expressions.go

@@ -1146,12 +1146,9 @@ func (w *Writer) isIntegerExpression(handle ir.ExpressionHandle) bool {
 }
 
 func (w *Writer) pointerNeedsDeref(pt ir.PointerType) bool {
-	switch pt.Space {
-	case ir.SpaceUniform, ir.SpaceStorage, ir.SpacePushConstant:
-		return true
-	default:
-		return false
-	}
+	// Buffer parameters use references (&) in MSL, not pointers (*).
+	// References don't need explicit dereference.
+	return false
 }
 
 func (w *Writer) shouldDerefPointer(handle ir.ExpressionHandle) bool {

msl/functions.go

@@ -528,7 +528,7 @@ func (w *Writer) writeGlobalResourceParam(handle uint32, global *ir.GlobalVariab
 		typeName := w.writeTypeName(global.Type, StorageAccess(0))
 
 		if space == spaceConstant || space == spaceDevice {
-			w.write("%s %s* %s [[buffer(%d)]]", space, typeName, name, binding)
+			w.write("%s %s& %s [[buffer(%d)]]", space, typeName, name, binding)
 		} else {
 			w.write("%s %s [[buffer(%d)]]", typeName, name, binding)
 		}

snapshot/testdata/golden/msl/collatz.msl

@@ -42,7 +42,7 @@ uint collatz_iterations(uint n_base) {
 }
 
 kernel void main_(metal::uint3 gid [[thread_position_in_grid]],
-    device Data* data [[buffer(0)]]) {
+    device Data& data [[buffer(0)]]) {
     uint _fc9 = collatz_iterations(data.values[gid[0]]);
     data.values[gid[0]] = _fc9;
 }

snapshot/testdata/golden/msl/globals.msl

@@ -27,7 +27,7 @@ T _naga_mod(T lhs, D rhs) {
 }
 
 kernel void main_(uint local_idx [[thread_index_in_threadgroup]],
-    constant Uniforms* uniforms [[buffer(0)]]) {
+    constant Uniforms& uniforms [[buffer(0)]]) {
     shared_data.inner[local_idx] = (uniforms.scale * float(local_idx));
     threadgroup_barrier(mem_flags::mem_threadgroup);
 }

snapshot/testdata/golden/msl/uniforms_mvp.msl

@@ -38,7 +38,7 @@ struct vs_main_Output {
 };
 
 vertex vs_main_Output vs_main(vs_main_Input _input [[stage_in]],
-    constant Uniforms* uniforms [[buffer(0)]]) {
+    constant Uniforms& uniforms [[buffer(0)]]) {
     auto position_ = _input.position_;
     VertexOutput out_ = VertexOutput();
 
@@ -57,7 +57,7 @@ struct fs_main_Input {
 };
 
 fragment metal::float4 fs_main(fs_main_Input _input [[stage_in]],
-    constant Uniforms* uniforms [[buffer(0)]]) {
+    constant Uniforms& uniforms [[buffer(0)]]) {
     auto world_pos = _input.world_pos;
 
     return metal::float4(((metal::normalize(world_pos) * 0.5) + 0.5), 1.0);

snapshot/testdata/golden/msl/workgroup_memory.msl

@@ -26,8 +26,8 @@ T _naga_mod(T lhs, D rhs) {
 }
 
 kernel void main_(uint lid [[thread_index_in_threadgroup]], metal::uint3 gid [[thread_position_in_grid]],
-    constant Params* params [[buffer(0)]],
-    device type_6* output [[buffer(1)]]) {
+    constant Params& params [[buffer(0)]],
+    device type_6& output [[buffer(1)]]) {
     shared_data.inner[lid] = (float(gid[0]) * 0.5);
     threadgroup_barrier(mem_flags::mem_threadgroup);
     if ((lid < 128u)) {