hello example shows pink screen on macOS

[blagodus]

hello example shows pink screen on macOS

Platform: macOS
Build command: CGO_ENABLED=0 go build .

Steps to reproduce

1. Clone the repo
2. cd examples/hello
3. CGO_ENABLED=0 go build . && ./hello

Result

Pink screen. See attached screenshot.

Console output

2026/03/13 19:17:21 render: ensure pipelines: SDF pipeline: compile sdf_render shader: metal: failed to compile MSL: program_source:87:26: error: cannot reserve 'buffer' resource location at index 0
    constant ClipParams& clip [[buffer(0)]]) {
                         ^
program_source:129:26: error: cannot reserve 'buffer' resource location at index 0
    constant ClipParams& clip [[buffer(0)]]) {
                         ^
 | Error Domain=MTLLibraryErrorDomain Code=3 "program_source:87:26: error: cannot reserve 'buffer' resource location at index 0
    constant ClipParams& clip [[buffer(0)]]) {
                         ^
program_source:129:26: error: cannot reserve 'buffer' resource location at index 0
    constant ClipParams& clip [[buffer(0)]]) {
                         ^
" UserInfo={NSLocalizedDescription=program_source:87:26: error: cannot reserve 'buffer' resource location at index 0
    constant ClipParams& clip [[buffer(0)]]) {
                         ^
program_source:129:26: error: cannot reserve 'buffer' resource location at index 0
    constant ClipParams& clip [[buffer(0)]]) {
                         ^
} | {
    NSLocalizedDescription = "program_source:87:26: error: cannot reserve 'buffer' resource location at index 0\n    constant ClipParams& clip [[buffer(0)]]) {\n                         ^\nprogram_source:129:26: error: cannot reserve 'buffer' resource location at index 0\n    constant ClipParams& clip [[buffer(0)]]) {\n                         ^\n";
}
MSL:
#include <metal_stdlib>
#include <simd/simd.h>

using metal::uint;

struct Uniforms {
    metal::float2 viewport;
    metal::float2 _pad;
};

struct VertexInput {
    metal::float2 position_;
    metal::float2 local;
    float shape_kind;
    float param1;
    float param2;
    float param3;
    float half_stroke;
    float is_stroked;
    metal::float4 color_;
};

struct VertexOutput {
    metal::float4 clip_position;
    metal::float2 local;
    float shape_kind;
    float param1;
    float param2;
    float param3;
    float half_stroke;
    float is_stroked;
    metal::float4 color_;
};

struct ClipParams {
    metal::float4 clip_rect;
    float clip_radius;
    float clip_enabled;
    metal::float2 _pad;
};

// Safe division helper (handles zero divisor)
template <typename T, typename D>
T _naga_div(T lhs, D rhs) {
    D nz = D(rhs != D(0));
    return lhs / (nz * rhs + D(!nz));
}

// Safe modulo helper (handles zero divisor)
template <typename T, typename D>
T _naga_mod(T lhs, D rhs) {
    D nz = D(rhs != D(0));
    return lhs % (nz * rhs + D(!nz));
}

float rrect_clip_coverage(metal::float2 frag_pos,
    constant ClipParams& clip) {

[kolkov]

Thanks for the report!

This is a shader compilation bug in the gg rendering engine (v0.36.4). The new ClipParams uniform for GPU rounded-rectangle clipping binds to buffer(0) on Metal, which conflicts with the existing Uniforms struct. This causes all SDF shaders to fail, resulting in the pink screen.

Upstream issue filed: gogpu/gg#209

The fix needs to happen in gg's shader binding layout (naga Metal backend mapping). Once gg patches this, we'll update the dependency and the hello example will work on macOS again.

Workaround: Unfortunately there's no easy workaround at the ui level — the shader conflict happens in the GPU pipeline before any rendering occurs.

[kolkov]

Fix released: naga v0.14.7

The root cause was in the naga shader compiler (MSL backend). It used only the WGSL @binding(N) index and ignored @group(N) when generating Metal buffer indices, so Uniforms and ClipParams both mapped to [[buffer(0)]].

Fixed in gogpu/naga v0.14.7 — Metal buffer indices are now sequential across bind groups.

To verify the fix before we release updated gg/ui, you can add this to your go.mod:

require github.com/gogpu/naga v0.14.7

Then rebuild and the pink screen should be gone. Full gg + ui releases with the updated dependency coming soon.

[blagodus]

Thanks for the super fast fix! The pink screen is indeed gone after updating to naga v0.14.7!

However, I'm now seeing a new issue — the window renders completely black with a small white/bright spot in the top-left corner (screenshot attached).

Happy to provide any additional debug info if helpful.

[kolkov]

The Metal shader fix (naga v0.14.7) resolved the pink screen — great!

The black screen with a bright spot in the corner is a different issue. Could you enable diagnostic logging to help us investigate? Add these lines at the start of your main():

import (
    "log/slog"
    "os"

    "github.com/gogpu/gg"
)

func main() {
    gg.SetLogger(slog.New(slog.NewTextHandler(os.Stderr, &slog.HandlerOptions{
        Level: slog.LevelDebug,
    })))
    // ... rest of main
}

Then run and share the first 20-30 lines of stderr output. This will show us viewport dimensions, GPU pipeline state, and surface configuration.

Also — which macOS version, which Mac model, and which display (Retina or external)?

There is also a known but separate macOS rendering offset issue being investigated in gogpu/gg#171, though your symptom may be unrelated.

[blagodus]

Thanks for the follow-up. Here are the diagnostic details you requested.

Machine: MacBook Pro 16-inch (Apple M1 Pro)
Display: Built-in Retina (3456x2160 native)
macOS: Tahoe 16.2

---

stderr output

time=2026-03-15T00:29:24.394+01:00 level=INFO msg="metal: framework initialized"
time=2026-03-15T00:29:24.395+01:00 level=INFO msg="metal: instance created"
time=2026-03-15T00:29:24.395+01:00 level=DEBUG msg="metal: enumerating adapters" count=1
time=2026-03-15T00:29:24.395+01:00 level=INFO msg="metal: adapter found" name="Apple M1 Pro" type=DiscreteGPU lowPower=false removable=false headless=false maxBuffer=9534832640 depth24Stencil8=false
time=2026-03-15T00:29:24.395+01:00 level=INFO msg="metal: device created" name="Apple M1 Pro"
time=2026-03-15T00:29:24.395+01:00 level=DEBUG msg="metal: adapter opened" maxFramesInFlight=2 blockSupport=true
time=2026-03-15T00:29:24.395+01:00 level=INFO msg="metal: surface configured" width=1600 height=1800 format=BGRA8Unorm presentMode=Fifo vsync=true
time=2026-03-15T00:29:24.464+01:00 level=DEBUG msg="metal: surface texture acquired" drawable=105553149381472 texture=5629916336 width=1600 height=1800
time=2026-03-15T00:29:24.465+01:00 level=INFO msg=ggcanvas.NewWithScale logical_w=800 logical_h=900 scale=2 physical_w=1600 physical_h=1800
time=2026-03-15T00:29:24.465+01:00 level=DEBUG msg="metal: WGSL→MSL compilation" elapsed=408.333µs mslBytes=3043
time=2026-03-15T00:29:24.543+01:00 level=DEBUG msg="VelloAccelerator init: compute pipeline unavailable" error="vello compute: create shader module for pathtag_reduce: metal: failed to compile MSL: program_source:70:12: error: unknown type name 'type_6'; did you mean 'type_8'?\n    device type_6& scene [[buffer(1)]],\n           ^~~~~~\n           type_8\nprogram_source:32:8: note: 'type_8' declared here\nstruct type_8 {\n       ^\nprogram_source:71:12: error: unknown type name 'type_7'; did you mean 'type_8'?\n    device type_7& reduced [[buffer(2)]]) {\n           ^~~~~~\n           type_8\nprogram_source:32:8: note: 'type_8' declared here\nstruct type_8 {\n       ^\nprogram_source:76:49: error: type 'device type_8' does not provide a subscript operator\n        PathMonoid _fc16 = new_path_monoid(scene[(config.pathtag_base + global_id[0])]);\n                                           ~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~\nprogram_source:84:5: error: use of undeclared identifier 'sh_scratch'\n    sh_scratch.inner[local_id[0]] = m;\n    ^\nprogram_source:85:25: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\n    threadgroup_barrier(mem_flags::mem_threadgroup);\n                        ^~~~~~~~~\n                        metal::mem_flags\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\nenum class mem_flags\n           ^\nprogram_source:91:40: error: use of undeclared identifier 'sh_scratch'\n            PathMonoid _fc34 = combine(sh_scratch.inner[(local_id[0] - (1u << i))], m);\n                                       ^\nprogram_source:94:29: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\n        threadgroup_barrier(mem_flags::mem_threadgroup);\n                            ^~~~~~~~~\n                            metal::mem_flags\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\nenum class mem_flags\n           ^\nprogram_source:95:9: error: use of undeclared identifier 'sh_scratch'\n        sh_scratch.inner[local_id[0]] = m;\n        ^\nprogram_source:96:29: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\n        threadgroup_barrier(mem_flags::mem_threadgroup);\n                            ^~~~~~~~~\n                            metal::mem_flags\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\nenum class mem_flags\n           ^\nprogram_source:99:22: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]\n    if ((local_id[0] == 0u)) {\n         ~~~~~~~~~~~~^~~~~\nprogram_source:99:22: note: remove extraneous parentheses around the comparison to silence this warning\n    if ((local_id[0] == 0u)) {\n        ~            ^    ~\nprogram_source:99:22: note: use '=' to turn this equality comparison into an assignment\n    if ((local_id[0] == 0u)) {\n                     ^~\n                     =\nprogram_source:100:16: error: type 'device type_8' does not provide a subscript operator\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\n        ~~~~~~~^~~~~~~~~\nprogram_source:100:29: error: use of undeclared identifier 'sh_scratch'\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\n                            ^\n | {\n    NSLocalizedDescription = \"program_source:70:12: error: unknown type name 'type_6'; did you mean 'type_8'?\\n    device type_6& scene [[buffer(1)]],\\n           ^~~~~~\\n           type_8\\nprogram_source:32:8: note: 'type_8' declared here\\nstruct type_8 {\\n       ^\\nprogram_source:71:12: error: unknown type name 'type_7'; did you mean 'type_8'?\\n    device type_7& reduced [[buffer(2)]]) {\\n           ^~~~~~\\n           type_8\\nprogram_source:32:8: note: 'type_8' declared here\\nstruct type_8 {\\n       ^\\nprogram_source:76:49: error: type 'device type_8' does not provide a subscript operator\\n        PathMonoid _fc16 = new_path_monoid(scene[(config.pathtag_base + global_id[0])]);\\n                                           ~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~\\nprogram_source:84:5: error: use of undeclared identifier 'sh_scratch'\\n    sh_scratch.inner[local_id[0]] = m;\\n    ^\\nprogram_source:85:25: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\\n    threadgroup_barrier(mem_flags::mem_threadgroup);\\n                        ^~~~~~~~~\\n                        metal::mem_flags\\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\\nenum class mem_flags\\n           ^\\nprogram_source:91:40: error: use of undeclared identifier 'sh_scratch'\\n            PathMonoid _fc34 = combine(sh_scratch.inner[(local_id[0] - (1u << i))], m);\\n                                       ^\\nprogram_source:94:29: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\\n        threadgroup_barrier(mem_flags::mem_threadgroup);\\n                            ^~~~~~~~~\\n                            metal::mem_flags\\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\\nenum class mem_flags\\n           ^\\nprogram_source:95:9: error: use of undeclared identifier 'sh_scratch'\\n        sh_scratch.inner[local_id[0]] = m;\\n        ^\\nprogram_source:96:29: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\\n        threadgroup_barrier(mem_flags::mem_threadgroup);\\n                            ^~~~~~~~~\\n                            metal::mem_flags\\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\\nenum class mem_flags\\n           ^\\nprogram_source:99:22: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]\\n    if ((local_id[0] == 0u)) {\\n         ~~~~~~~~~~~~^~~~~\\nprogram_source:99:22: note: remove extraneous parentheses around the comparison to silence this warning\\n    if ((local_id[0] == 0u)) {\\n        ~            ^    ~\\nprogram_source:99:22: note: use '=' to turn this equality comparison into an assignment\\n    if ((local_id[0] == 0u)) {\\n                     ^~\\n                     =\\nprogram_source:100:16: error: type 'device type_8' does not provide a subscript operator\\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\\n        ~~~~~~~^~~~~~~~~\\nprogram_source:100:29: error: use of undeclared identifier 'sh_scratch'\\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\\n                            ^\\n\";\n}\nMSL:\n#include <metal_stdlib>\n#include <simd/simd.h>\n\nusing metal::uint;\n\nstruct PathMonoid {\n    uint trans_ix;\n    uint path_seg_ix;\n    uint path_seg_offset;\n    uint style_ix;\n    uint path_ix;\n};\n\nstruct Config {\n    uint width_in_tiles;\n    uint height_in_tiles;\n    uint target_width;\n    uint target_height;\n    uint n_drawobj;\n    uint n_path;\n    uint n_clip;\n    uint pathtag_base;\n    uint pathdata_base;\n    uint drawtag_base;\n    uint drawdata_base;\n    uint transform_base;\n    uint style_base;\n    uint n_lines;\n    uint bg_color;\n};\n\nstruct type_8 {\n    PathMonoid inner[256];\n};\n\nconstant uint WG_SIZE = 256u;\n\n// Safe division helper (handles zero divisor)\ntemplate <typename T, typename D>\nT _naga_div(T lhs, D rhs) {\n    D nz = D(rhs != D(0));\n    return lhs / (nz * rhs + D(!nz));\n}\n\n// Safe modulo helper (handles zero divisor)\ntemplate <typename T, typename D>\nT _naga_mod(T lhs, D rhs) {\n    D nz = D(rhs != D(0));\n    return lhs % (nz * rhs + D(!nz));\n}\n\nPathMonoid new_path_monoid(uint tag_word) {\n    uint a = (((tag_word & 50529027u) + ((tag_word >> 2u) & 16843009u)) + (((tag_word & 50529027u) + ((tag_word >> 2u) & 16843009u)) & (((tag_word >> 3u) & 16843009u) * 15u)));\n    \n    a = (a + (a >> 8u));\n    a = (a + (a >> 16u));\n    return PathMonoid{metal::popcount((tag_word & 538976288u)), metal::popcount((((tag_word & 50529027u) * 7u) & 67372036u)), (a & 255u), metal::popcount((tag_word & 1077952576u)), metal::popcount((tag_word & 269488144u))};\n}\n\nPathMonoid combine(PathMonoid a, PathMonoid b) {\n    return PathMonoid{(a.trans_ix + b.trans_ix), (a.path_seg_ix + b.path_seg_ix), (a.path_seg_offset + b.path_seg_offset), (a.style_ix + b.style_ix), (a.path_ix + b.path_ix)};\n}\n\nPathMonoid identity() {\n    return PathMonoid{0u, 0u, 0u, 0u, 0u};\n}\n\nkernel void main_(metal::uint3 global_id [[thread_position_in_grid]], metal::uint3 local_id [[thread_position_in_threadgroup]], metal::uint3 wg_id [[threadgroup_position_in_grid]],\n    constant Config& config [[buffer(0)]],\n    device type_6& scene [[buffer(1)]],\n    device type_7& reduced [[buffer(2)]]) {\n    PathMonoid m = PathMonoid();\n    uint i = 0u;\n    \n    if ((global_id[0] < (config.pathdata_base - config.pathtag_base))) {\n        PathMonoid _fc16 = new_path_monoid(scene[(config.pathtag_base + global_id[0])]);\n        m = _fc16;\n    } else {\n        {\n            PathMonoid _fc17 = identity();\n            m = _fc17;\n        }\n    }\n    sh_scratch.inner[local_id[0]] = m;\n    threadgroup_barrier(mem_flags::mem_threadgroup);\n    while (true) {\n        if (!((i < 8u))) {\n            break;\n        }\n        if ((local_id[0] >= (1u << i))) {\n            PathMonoid _fc34 = combine(sh_scratch.inner[(local_id[0] - (1u << i))], m);\n            m = _fc34;\n        }\n        threadgroup_barrier(mem_flags::mem_threadgroup);\n        sh_scratch.inner[local_id[0]] = m;\n        threadgroup_barrier(mem_flags::mem_threadgroup);\n        i = (i + 1u);\n    }\n    if ((local_id[0] == 0u)) {\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\n    }\n}\n\n"
time=2026-03-15T00:29:24.543+01:00 level=INFO msg="switched to shared GPU device" adapter=*metal.Device
time=2026-03-15T00:29:24.543+01:00 level=INFO msg="NewContext HiDPI" logical_w=800 logical_h=900 scale=2 physical_w=1600 physical_h=1800
time=2026-03-15T00:29:24.544+01:00 level=DEBUG msg="GPURenderSession.SetSurfaceTarget changed" surface=true width=1600 height=1800 modeChanged=true sizeChanged=true
time=2026-03-15T00:29:24.544+01:00 level=DEBUG msg="SetSurfaceTarget configured" width=1600 height=1800
time=2026-03-15T00:29:24.632+01:00 level=INFO msg="LayoutText result" text="gogpu/ui — Widget Demo" glyphs=22 quads=19 outlineSkip=3 atlasSkip=0 boundsSkip=0 viewport=1600x1800
time=2026-03-15T00:29:24.681+01:00 level=INFO msg="LayoutText result" text=Checkboxes glyphs=10 quads=10 outlineSkip=0 atlasSkip=0 boundsSkip=0 viewport=1600x1800
time=2026-03-15T00:29:24.764+01:00 level=INFO msg="LayoutText result" text="Enable notifications" glyphs=20 quads=19 outlineSkip=1 atlasSkip=0 boundsSkip=0 viewport=1600x1800
time=2026-03-15T00:29:24.794+01:00 level=INFO msg="LayoutText result" text="Dark mode" glyphs=9 quads=8 outlineSkip=1 atlasSkip=0 boundsSkip=0 viewport=1600x1800
time=2026-03-15T00:29:24.802+01:00 level=INFO msg="LayoutText result" text="Disabled checkbox" glyphs=17 quads=16 outlineSkip=1 atlasSkip=0 boundsSkip=0 viewport=1600x1800

---

Let me know if you need anything else to investigate further.

[kolkov]

@blagodus Thank you for the diagnostic logs — they're very informative!

Two findings from your output:

1. MSL compute shader errors (type_6/type_7, sh_scratch undeclared) — This is a naga MSL codegen issue with workgroup-shared memory in compute shaders. The VelloAccelerator correctly falls back to SDF rendering when compute shaders fail, so this doesn't affect visual output. We're tracking this separately.

2. Content offset / black screen — This is the same root cause as gg#171. We've now fixed it.

What happened: We've finally resolved a significant technical debt — for historical reasons, both gogpu and gg were using the wgpu HAL layer directly, bypassing the public API and its validation/lifecycle management. This meant surface configuration, viewport setup, and DPI handling had unclear ownership across layers. We've completed a major migration to the proper three-layer wgpu architecture (public API → core → HAL) with clear separation of responsibilities. This was a substantial refactoring we'd been postponing, but the macOS issues made it clear we couldn't delay any longer.

The specific fixes for your issue:

• Explicit SetViewport in all GPU render passes (was relying on Metal's default)
• PrepareFrame per-frame contentsScale refresh (was only set at surface creation)

Please update to:

github.com/gogpu/gogpu v0.24.0
github.com/gogpu/gg v0.37.0

With the architectural debt resolved, we can now focus on fixing remaining bugs properly. Please let us know if the content offset is resolved!

[blagodus]

@kolkov Tested again after the latest fixes, but still getting the same black screen.

---

log output

./hello         
time=2026-03-15T14:54:06.721+01:00 level=INFO msg="metal: framework initialized"
time=2026-03-15T14:54:06.722+01:00 level=INFO msg="metal: instance created"
time=2026-03-15T14:54:06.722+01:00 level=DEBUG msg="metal: enumerating adapters" count=1
time=2026-03-15T14:54:06.722+01:00 level=INFO msg="metal: adapter found" name="Apple M1 Pro" type=DiscreteGPU lowPower=false removable=false headless=false maxBuffer=9534832640 depth24Stencil8=false
2026/03/15 14:54:06 INFO adapter selected name="Apple M1 Pro" type=DiscreteGPU
time=2026-03-15T14:54:06.722+01:00 level=INFO msg="metal: device created" name="Apple M1 Pro"
time=2026-03-15T14:54:06.722+01:00 level=DEBUG msg="metal: adapter opened" maxFramesInFlight=2 blockSupport=true
time=2026-03-15T14:54:06.722+01:00 level=INFO msg="metal: surface configured" width=1600 height=1800 format=BGRA8Unorm presentMode=Fifo vsync=true
time=2026-03-15T14:54:06.754+01:00 level=DEBUG msg="metal: surface texture acquired" drawable=105553119072640 texture=4504710672 width=1600 height=1800
time=2026-03-15T14:54:06.754+01:00 level=INFO msg=ggcanvas.NewWithScale logical_w=800 logical_h=900 scale=2 physical_w=1600 physical_h=1800
time=2026-03-15T14:54:06.754+01:00 level=DEBUG msg="metal: WGSL→MSL compilation" elapsed=425.625µs mslBytes=3043
time=2026-03-15T14:54:06.894+01:00 level=DEBUG msg="VelloAccelerator init: compute pipeline unavailable" error="vello compute: create shader module for pathtag_reduce: wgpu: failed to create shader module: metal: failed to compile MSL: program_source:70:12: error: unknown type name 'type_6'; did you mean 'type_8'?\n    device type_6& scene [[buffer(1)]],\n           ^~~~~~\n           type_8\nprogram_source:32:8: note: 'type_8' declared here\nstruct type_8 {\n       ^\nprogram_source:71:12: error: unknown type name 'type_7'; did you mean 'type_8'?\n    device type_7& reduced [[buffer(2)]]) {\n           ^~~~~~\n           type_8\nprogram_source:32:8: note: 'type_8' declared here\nstruct type_8 {\n       ^\nprogram_source:76:49: error: type 'device type_8' does not provide a subscript operator\n        PathMonoid _fc16 = new_path_monoid(scene[(config.pathtag_base + global_id[0])]);\n                                           ~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~\nprogram_source:84:5: error: use of undeclared identifier 'sh_scratch'\n    sh_scratch.inner[local_id[0]] = m;\n    ^\nprogram_source:85:25: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\n    threadgroup_barrier(mem_flags::mem_threadgroup);\n                        ^~~~~~~~~\n                        metal::mem_flags\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\nenum class mem_flags\n           ^\nprogram_source:91:40: error: use of undeclared identifier 'sh_scratch'\n            PathMonoid _fc34 = combine(sh_scratch.inner[(local_id[0] - (1u << i))], m);\n                                       ^\nprogram_source:94:29: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\n        threadgroup_barrier(mem_flags::mem_threadgroup);\n                            ^~~~~~~~~\n                            metal::mem_flags\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\nenum class mem_flags\n           ^\nprogram_source:95:9: error: use of undeclared identifier 'sh_scratch'\n        sh_scratch.inner[local_id[0]] = m;\n        ^\nprogram_source:96:29: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\n        threadgroup_barrier(mem_flags::mem_threadgroup);\n                            ^~~~~~~~~\n                            metal::mem_flags\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\nenum class mem_flags\n           ^\nprogram_source:99:22: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]\n    if ((local_id[0] == 0u)) {\n         ~~~~~~~~~~~~^~~~~\nprogram_source:99:22: note: remove extraneous parentheses around the comparison to silence this warning\n    if ((local_id[0] == 0u)) {\n        ~            ^    ~\nprogram_source:99:22: note: use '=' to turn this equality comparison into an assignment\n    if ((local_id[0] == 0u)) {\n                     ^~\n                     =\nprogram_source:100:16: error: type 'device type_8' does not provide a subscript operator\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\n        ~~~~~~~^~~~~~~~~\nprogram_source:100:29: error: use of undeclared identifier 'sh_scratch'\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\n                            ^\n | {\n    NSLocalizedDescription = \"program_source:70:12: error: unknown type name 'type_6'; did you mean 'type_8'?\\n    device type_6& scene [[buffer(1)]],\\n           ^~~~~~\\n           type_8\\nprogram_source:32:8: note: 'type_8' declared here\\nstruct type_8 {\\n       ^\\nprogram_source:71:12: error: unknown type name 'type_7'; did you mean 'type_8'?\\n    device type_7& reduced [[buffer(2)]]) {\\n           ^~~~~~\\n           type_8\\nprogram_source:32:8: note: 'type_8' declared here\\nstruct type_8 {\\n       ^\\nprogram_source:76:49: error: type 'device type_8' does not provide a subscript operator\\n        PathMonoid _fc16 = new_path_monoid(scene[(config.pathtag_base + global_id[0])]);\\n                                           ~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~\\nprogram_source:84:5: error: use of undeclared identifier 'sh_scratch'\\n    sh_scratch.inner[local_id[0]] = m;\\n    ^\\nprogram_source:85:25: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\\n    threadgroup_barrier(mem_flags::mem_threadgroup);\\n                        ^~~~~~~~~\\n                        metal::mem_flags\\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\\nenum class mem_flags\\n           ^\\nprogram_source:91:40: error: use of undeclared identifier 'sh_scratch'\\n            PathMonoid _fc34 = combine(sh_scratch.inner[(local_id[0] - (1u << i))], m);\\n                                       ^\\nprogram_source:94:29: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\\n        threadgroup_barrier(mem_flags::mem_threadgroup);\\n                            ^~~~~~~~~\\n                            metal::mem_flags\\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\\nenum class mem_flags\\n           ^\\nprogram_source:95:9: error: use of undeclared identifier 'sh_scratch'\\n        sh_scratch.inner[local_id[0]] = m;\\n        ^\\nprogram_source:96:29: error: use of undeclared identifier 'mem_flags'; did you mean 'metal::mem_flags'?\\n        threadgroup_barrier(mem_flags::mem_threadgroup);\\n                            ^~~~~~~~~\\n                            metal::mem_flags\\n/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.850/include/metal/metal_compute:13:12: note: 'metal::mem_flags' declared here\\nenum class mem_flags\\n           ^\\nprogram_source:99:22: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]\\n    if ((local_id[0] == 0u)) {\\n         ~~~~~~~~~~~~^~~~~\\nprogram_source:99:22: note: remove extraneous parentheses around the comparison to silence this warning\\n    if ((local_id[0] == 0u)) {\\n        ~            ^    ~\\nprogram_source:99:22: note: use '=' to turn this equality comparison into an assignment\\n    if ((local_id[0] == 0u)) {\\n                     ^~\\n                     =\\nprogram_source:100:16: error: type 'device type_8' does not provide a subscript operator\\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\\n        ~~~~~~~^~~~~~~~~\\nprogram_source:100:29: error: use of undeclared identifier 'sh_scratch'\\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\\n                            ^\\n\";\n}\nMSL:\n#include <metal_stdlib>\n#include <simd/simd.h>\n\nusing metal::uint;\n\nstruct PathMonoid {\n    uint trans_ix;\n    uint path_seg_ix;\n    uint path_seg_offset;\n    uint style_ix;\n    uint path_ix;\n};\n\nstruct Config {\n    uint width_in_tiles;\n    uint height_in_tiles;\n    uint target_width;\n    uint target_height;\n    uint n_drawobj;\n    uint n_path;\n    uint n_clip;\n    uint pathtag_base;\n    uint pathdata_base;\n    uint drawtag_base;\n    uint drawdata_base;\n    uint transform_base;\n    uint style_base;\n    uint n_lines;\n    uint bg_color;\n};\n\nstruct type_8 {\n    PathMonoid inner[256];\n};\n\nconstant uint WG_SIZE = 256u;\n\n// Safe division helper (handles zero divisor)\ntemplate <typename T, typename D>\nT _naga_div(T lhs, D rhs) {\n    D nz = D(rhs != D(0));\n    return lhs / (nz * rhs + D(!nz));\n}\n\n// Safe modulo helper (handles zero divisor)\ntemplate <typename T, typename D>\nT _naga_mod(T lhs, D rhs) {\n    D nz = D(rhs != D(0));\n    return lhs % (nz * rhs + D(!nz));\n}\n\nPathMonoid new_path_monoid(uint tag_word) {\n    uint a = (((tag_word & 50529027u) + ((tag_word >> 2u) & 16843009u)) + (((tag_word & 50529027u) + ((tag_word >> 2u) & 16843009u)) & (((tag_word >> 3u) & 16843009u) * 15u)));\n    \n    a = (a + (a >> 8u));\n    a = (a + (a >> 16u));\n    return PathMonoid{metal::popcount((tag_word & 538976288u)), metal::popcount((((tag_word & 50529027u) * 7u) & 67372036u)), (a & 255u), metal::popcount((tag_word & 1077952576u)), metal::popcount((tag_word & 269488144u))};\n}\n\nPathMonoid combine(PathMonoid a, PathMonoid b) {\n    return PathMonoid{(a.trans_ix + b.trans_ix), (a.path_seg_ix + b.path_seg_ix), (a.path_seg_offset + b.path_seg_offset), (a.style_ix + b.style_ix), (a.path_ix + b.path_ix)};\n}\n\nPathMonoid identity() {\n    return PathMonoid{0u, 0u, 0u, 0u, 0u};\n}\n\nkernel void main_(metal::uint3 global_id [[thread_position_in_grid]], metal::uint3 local_id [[thread_position_in_threadgroup]], metal::uint3 wg_id [[threadgroup_position_in_grid]],\n    constant Config& config [[buffer(0)]],\n    device type_6& scene [[buffer(1)]],\n    device type_7& reduced [[buffer(2)]]) {\n    PathMonoid m = PathMonoid();\n    uint i = 0u;\n    \n    if ((global_id[0] < (config.pathdata_base - config.pathtag_base))) {\n        PathMonoid _fc16 = new_path_monoid(scene[(config.pathtag_base + global_id[0])]);\n        m = _fc16;\n    } else {\n        {\n            PathMonoid _fc17 = identity();\n            m = _fc17;\n        }\n    }\n    sh_scratch.inner[local_id[0]] = m;\n    threadgroup_barrier(mem_flags::mem_threadgroup);\n    while (true) {\n        if (!((i < 8u))) {\n            break;\n        }\n        if ((local_id[0] >= (1u << i))) {\n            PathMonoid _fc34 = combine(sh_scratch.inner[(local_id[0] - (1u << i))], m);\n            m = _fc34;\n        }\n        threadgroup_barrier(mem_flags::mem_threadgroup);\n        sh_scratch.inner[local_id[0]] = m;\n        threadgroup_barrier(mem_flags::mem_threadgroup);\n        i = (i + 1u);\n    }\n    if ((local_id[0] == 0u)) {\n        reduced[wg_id[0]] = sh_scratch.inner[(WG_SIZE - 1u)];\n    }\n}\n\n"
time=2026-03-15T14:54:06.894+01:00 level=INFO msg="switched to shared GPU device" adapter=*wgpu.Device
time=2026-03-15T14:54:06.894+01:00 level=INFO msg="NewContext HiDPI" logical_w=800 logical_h=900 scale=2 physical_w=1600 physical_h=1800
time=2026-03-15T14:54:06.894+01:00 level=DEBUG msg="GPURenderSession.SetSurfaceTarget changed" surface=true width=1600 height=1800 modeChanged=true sizeChanged=true
time=2026-03-15T14:54:06.894+01:00 level=DEBUG msg="SetSurfaceTarget configured" width=1600 height=1800
time=2026-03-15T14:54:06.981+01:00 level=INFO msg="LayoutText result" text="gogpu/ui — Widget Demo" glyphs=22 quads=19 outlineSkip=3 atlasSkip=0 boundsSkip=0 viewport=1600x1800
time=2026-03-15T14:54:07.029+01:00 level=INFO msg="LayoutText result" text=Checkboxes glyphs=10 quads=10 outlineSkip=0 atlasSkip=0 boundsSkip=0 viewport=1600x1800
time=2026-03-15T14:54:07.108+01:00 level=INFO msg="LayoutText result" text="Enable notifications" glyphs=20 quads=19 outlineSkip=1 atlasSkip=0 boundsSkip=0 viewport=1600x1800
time=2026-03-15T14:54:07.135+01:00 level=INFO msg="LayoutText result" text="Dark mode" glyphs=9 quads=8 outlineSkip=1 atlasSkip=0 boundsSkip=0 viewport=1600x1800

[kolkov]

@blagodus Thank you for testing again. The persistent black screen confirms this is the same macOS Metal issue being tracked in gg#171. Our fixes work on Windows/Linux but don't resolve the root cause on Metal.

To narrow down the problem, could you try building with our Rust backend? It wraps the well-tested wgpu-native library with a completely different Metal implementation:

# 1. Download wgpu-native for macOS arm64:
# https://github.com/gfx-rs/wgpu-native/releases
# Place libwgpu_native.dylib in /usr/local/lib/ or next to the binary

# 2. Build with Rust backend:
go build -tags rust .

This isolates whether the problem is in our Pure Go Metal HAL or in the rendering pipeline itself.

The MSL compute shader errors (sh_scratch, type_6) are a separate naga codegen issue — they don't cause the black screen (SDF fallback handles rendering). We'll track that separately.