From Theoretical Deductions to 72 FPS in Production: Implementing Dawn's Metal Bindless Backend and a Zero-Compromise WebGPU Backend for NoGraphicsAPI

Abstract: In my previous post, I shared a technical discussion with former Unity senior rendering research lead and legendary graphics architect Sebastian Aaltonen on the feasibility of bringing NoGraphicsAPI to WebGPU. While we deduced a theoretical path using experimental Google Dawn extensions, we faced strict architectural red lines set by Aaltonen: zero ugly memory-access macros in shaders (Zero Macros), alongside overcoming the absence of 64-bit hardware pointers (Buffer Device Address), persistent CPU mapping, and a mature cross-platform bindless backend in WebGPU.

> Theory is one thing; shipping code is another. Over the past few days, I turned those architectural deductions into a full-stack reality across both native desktop and our cross-process WebAssembly runtime:

1. Completing Google Dawn's Metal Bindless Backend: Built ResourceTableMTL from the ground up using Metal 3 Argument Buffers Tier 2 (direct 64-bit gpuResourceID writes) and TextureView-granular residency tracking, passing 100% of Dawn's 44 built-in unit tests and 47 end-to-end tests under full hardware GPU validation;

2. Refactoring emdawnwebgpu's Cross-Process Wire Codegen: Automated all 249 WebGPU C API pass-through functions and binary-search dispatch tables from dawn.json, eliminating nearly 600 lines of hand-maintained macros and unlocking cross-process serialization for WGPUResourceTable;

3. Building a Zero-Compromise WebGPU Backend for NoGraphicsAPI: By combining a Slang generic smart pointer (GpuPtr<T>) with build-time shadow header rewriting, ByteAddressBuffer-based cross-word-size C layout alignment, and transparent Mesh Shader-to-Vertex Pulling lowering, we got all three official examples—triangle, cube, and the 262,144-object deferred_renderer—running at 135 FPS on macOS Native and a locked 72 FPS inside our GFX_Runtime WebAssembly container, with zero lines of changes to any business .slang shader or *_shared.h header.

> This post documents the core architectural designs and the most memorable debugging war stories from these three milestones.

---

Ported NoGraphicsAPI to WebGPU bindless experimental
https://github.com/SetoKaiba/NoGraphicsAPI/tree/webgpu_bindless

Added Metal backend to Google Dawn bindless
https://github.com/SetoKaiba/dawn/tree/bindless_metal

WebGPU bindless on iPhone 13 Pro Max with low battery health charging, frequency degraded

WebGPU bindless on Mac Mini M4

1. Completing the First Puzzle Piece: A Full Metal Bindless Backend for Google Dawn

While surveying Google Dawn's (chromium/8103 branch) source code in the previous article, we discovered the ChromiumExperimentalSamplingResourceTable extension spearheaded by Dawn Tech Lead Corentin Wallez (@Kangz). However, as soon as we tried to enable it on macOS and iOS, we hit an immediate roadblock:

While Dawn already had frontend validation, Tint WGSL-to-MSL compiler transforms (resource_table_helper), and backend implementations for D3D12 (ResourceTableD3D12) and Vulkan (ResourceTableVk), the Metal backend (ResourceTableMTL) was completely unimplemented!

Without a Metal backend in Dawn, WebGPU Bindless on Apple Silicon was nothing more than a mirage.

1. Architectural Design: Bypassing MTLArgumentEncoder with Direct 64-Bit gpuResourceID Writes

In legacy Metal 2 Argument Buffers (Tier 1), binding textures and samplers required encoding through MTLArgumentEncoder, which was stateful and constrained by strict texture type matching. On modern Apple Silicon (Metal 3/4) hardware supporting Argument Buffers Tier 2, every texture view and sampler possesses a 64-bit hardware resource handle: MTLResourceID (backed by a raw uint64_t _impl).

Inspecting Dawn's Tint compiler source (src/tint/lang/msl/writer/raise/resource_table.cc) reveals that when Tint lowers WGSL's getResource<T>(index) to MSL, it synthesizes two implicit buffer bindings:

• Argument Buffer (Slot 28): Effectively a tightly packed array of 8-byte uint64_t handles (ulong* tint_resource_table), which the shader casts directly into Metal's texture2d<float> or sampler handles;
• Metadata Buffer (Slot 29): Starts with APISize (the user-requested table capacity), followed by an array of ResourceType enum IDs tracking the exact type bound to each slot. Before sampling, the generated shader queries the Metadata Buffer for bounds and type-safety validation: if an index is out of bounds or its bound type mismatches the requested type, execution safely redirects to one of the 35 Default Fallback Resources appended at the tail of the table ([APISize .. APISize + 34]). This guarantees Web sandbox safety (never faulting the GPU on invalid indices) while preserving dynamic indexing.

flowchart LR
    subgraph WGSL ["WGSL Shader Invocation"]
        Call["getResource Dynamic Texture Fetch (idx)"]
    end

    subgraph TintMSL ["Tint-Generated MSL Validation & Lookup"]
        Meta["Metadata Buffer (Slot 29)\n[APISize, TypeID_0, TypeID_1, ...]"]
        Table["Argument Buffer (Slot 28)\n[gpuResourceID_0, ..., Fallback_0..34]"]
        Check{"idx less than APISize AND\nTypeID[idx] matches requested type?"}
        Valid["Read Table[idx]\nReinterpret cast to metal::texture2d"]
        Fallback["Read Table[APISize + DefaultSlot]\nSafe fallback to 1x1 zero texture"]
    end

    Call --> Check
    Meta --> Check
    Check -->|Validation Passed| Valid
    Check -->|Out of Bounds or Type Mismatch| Fallback
    Table --> Valid
    Table --> Fallback

With Tint's ABI contract clear, we implemented ResourceTableMTL.{h,mm} in Dawn's Metal backend and wired up the full pipeline:

1. Direct 64-Bit Handle Writes: ResourceTableMTL allocates an MTLResourceStorageModeShared buffer of size (APISize + 35) * sizeof(uint64_t), completely bypassing MTLArgumentEncoder and writing [mtlTexture gpuResourceID]._impl or [mtlSampler gpuResourceID]._impl directly into the 8-byte slots (while enabling supportArgumentBuffers = YES in SamplerMTL.mm);
2. Copy-on-Write (CoW) In-Flight Protection: When the CPU mutates slots via wgpuResourceTableUpdate, ResourceTableMTL checks mLastUsageSerial > queue->GetCompletedCommandSerial(). If an in-flight GPU pass is still reading the table buffer, it clones a fresh MTLBuffer for subsequent CPU writes—eliminating CPU-GPU race conditions with zero pipeline stalls.

2. Debugging War Story: The Depth24Plus Swizzle View Residency Trap

Once compiled, all 44 dawn_unittests and basic Metal end-to-end tests passed immediately. However, when we enabled full Metal backend validation (--enable-backend-validation=full, activating the Metal Debug Layer and GPU Shader Validation) and ran ResourceTableTests.UseDefaultSamplers, the GPU validation layer threw a fatal error:

Invalid Metal usage: Invalid texture type MTLTextureTypeInvalid bound to shader, expected MTLTextureType2D

The slot held a valid gpuResourceID and the Metadata Buffer check took the valid branch—so why did the Metal GPU claim the bound texture was MTLTextureTypeInvalid?

Digging deeper uncovered two critical Metal GPU Residency invariants:

1. TextureView—Not the Parent Texture—Is the Atomic Unit of Residency: In Metal, when resources are accessed indirectly via 64-bit IDs inside an Argument Buffer, the driver cannot statically deduce which textures a shader will sample. The command encoder must explicitly call [encoder useResources:count:usage:stages:] to make every referenced resource resident in the GPU page table. Our initial implementation tracked base Texture* pointers and passed their id<MTLTexture> handles to useResources. However, when depthTexture.CreateView() is called on a Depth24Plus texture, Dawn invokes [mtlTexture newTextureViewWithPixelFormat:... swizzle:kR001Swizzle] to satisfy WebGPU's (r, 0, 0, 1) depth sampling semantics—creating a distinct child id<MTLTexture> object! The parent texture had gpuResourceID: 2, while the swizzled TextureView had gpuResourceID: 3! Passing only the parent texture to useResources left handle 3 unresident in the GPU hardware page table, causing an immediate GPU validation fault upon dereference.
2. All 35 Default Fallback Textures Must Always Be Resident: Even when a shader accesses valid slots, speculative execution or fallback testing can touch the default fallback resources at [APISize .. APISize + 34]. Omitting those 35 built-in 1×1 fallback textures from useResources likewise triggers GPU Shader Validation faults.

The Fix:
We refactored ResourceTableMTL to track Ref<TextureViewBase> directly (registering each view's exact texView->GetMTLTexture()) and cached all 35 default fallback textures in mDefaultResidentResources during initialization. With zero modifications to Dawn's test suite:

• dawn_unittests --gtest_filter="*ResourceTable*": 44 / 44 PASSED (100%);
• dawn_end2end_tests --backend=metal --enable-backend-validation=full --gtest_filter="*ResourceTable*": 47 / 47 PASSED (100%, 0 Metal API/GPU validation errors, 0 leaked bytes under /usr/bin/leaks)!

---

2. Unblocking the Cross-Process Artery: Auto-Generating webgpuwire from dawn.json

With the host-side Metal Bindless backend ready, the next step was exposing wgpuDeviceCreateResourceTable and its sibling APIs to WebAssembly client code running inside our hybrid container.

In our hybrid runtime architecture (detailed in Post #1), WASM applications statically link libemdawnwebgpu_wire.a to serialize WebGPU C calls into a binary command stream sent over IPC to the native host. However, after rebasing onto Dawn's chromium/8103 branch, inspecting third_party/emdawnwebgpu/pkg/webgpu/src/webgpuwire.cpp revealed a glaring maintenance bottleneck:

webgpuwire.cpp relied on a nearly 600-line hand-maintained WEBGPU_FOREACH_PROC macro table, and all 11 newly added WGPUResourceTable functions—along with SetImmediates and other new extensions—were completely missing!

Why Not Let Dawn's Code Generator Do the Work?

Examining src/dawn/wire/client/ResourceTable.cpp showed that slot lifecycle tracking (determining when a removed slot can be safely recycled via Queue::GetCompletedSubmitIndex) and wire command serialization are already 100% self-contained inside dawn::wire::client. They require no custom JavaScript bridging in webgpuwire.cpp—just a pure 1:1 pass-through:

WGPU_EXPORT WGPUResourceTable wgpuDeviceCreateResourceTable(
    WGPUDevice device, WGPUResourceTableDescriptor const * descriptor) {
    return wgpuDawnWireClientDeviceCreateResourceTable(device, descriptor);
}

In fact, across all 249 WebGPU C APIs, only 14 functions require custom logic in webgpuwire.cpp (such as flushing the IPC wire on wgpuQueueSubmit, binding HTML canvas surfaces, or driving future polling). The remaining 235 functions are pure mechanical forwards to wgpuDawnWireClient*!

To eliminate this maintenance burden permanently, we overhauled emdawnwebgpu's build pipeline:

1. Added a New Jinja2 Codegen Template (webgpu_wire_procs.h): Hooked into generator/dawn_json_generator.py using the exact same tag filter as dawn_wire (enabled_tags=['dawn', 'deprecated'], disabled_tags=['native']);
2. Default Pass-Throughs with Macro Opt-Out: Generated #ifndef WEBGPU_WIRE_CUSTOM_<WGPUName> guarded pass-through implementations for all 249 functions, plus a compile-time lexicographically sorted binary-search table (sEmwgpuWireProcMap) that 100% auto-generates wgpuGetProcAddress;
3. Streamlined webgpuwire.cpp: Deleted nearly 600 lines of hand-written macro tables, retaining only the 14 custom functions tagged with WEBGPU_WIRE_CUSTOM_*. All 11 WGPUResourceTable bindless APIs were immediately exported, and future Dawn upgrades will automatically stay in sync.

---

3. The Ultimate Test: Bringing NoGraphicsAPI to WebGPU with Zero Macro Pollution

With the underlying foundation complete, we arrived at the main event: building a WebGPU backend (NoGraphicsAPIWebGPU.cpp) for Sebastian Aaltonen's NoGraphicsAPI, and getting all three official examples (triangle, cube, and deferred_renderer simulating 262,144 dynamic lit cubes) running on both macOS Dawn Native and GFX_Runtime (Emscripten WASM .a Wire mode).

To stay true to NoGraphicsAPI's minimalist philosophy, I enforced three strict rules:

1. Zero Business Code Intrusion: Not a single line of .slang shader code in the examples may be altered, and no platform-specific macros or manual _pad fields may be added to user *_shared.h headers;
2. Strict Compliance with AGENTS.md: Pure C webgpu.h implementation, zero C++ standard library headers (except <initializer_list> for Span), -fno-exceptions -fno-rtti everywhere, no lambdas, and no dynamic containers;
3. Single-Source Dual-Target: The exact same C++ and Slang codebase must compile natively on macOS and to WebAssembly via Emscripten.

Here is how we conquered 64-bit GPU pointers (BDA), cross-word-size memory layout alignment, unified Bindless heaps, and Mesh Shader lowering.

flowchart TD
    subgraph Source ["100% Unmodified Example Source Code"]
        SlangSrc["Business Shaders (*.slang)\nroot.vertices[i] / gpu_texture(idx)"]
        SharedH["Shared Headers (*_shared.h)\nstruct Root { Vertex* vertices; ... }"]
    end

    subgraph BuildPipeline ["CMake Build-Time Transparent Pipeline (NoGraphicsAPI_slang.cmake)"]
        Rewrite["1. rewrite_bda.py (Shadow Header Generation)\nAutomatically rewrites raw T* to generic GpuPtr"]
        SlangC["2. slangc -target wgsl -fvk-use-c-layout\nLowers GpuPtr subscripting to __mega_buffer.Load"]
        FixWGSL["3. fix_wgsl.py (WGSL AST Post-Processing)\nHandle var-to-let fix + automatic vertex Y-flip"]
    end

    subgraph Runtime ["NoGraphicsAPIWebGPU Runtime Foundation"]
        MegaBuf["@binding(0) Global Mega-Buffer (ByteAddressBuffer)\nBacks all GpuHeap virtual addresses & CPU shadow sync"]
        RootBuf["@binding(1) Dynamic Root Ring Buffer\n256-byte aligned, C-layout raw byte loads"]
        ResTable["WGPUResourceTable (65,536 Slots)\n0..32767 Texture Heap | 32768..65535 Sampler Heap"]
    end

    SharedH --> Rewrite
    SlangSrc --> SlangC
    Rewrite --> SlangC
    SlangC --> FixWGSL
    FixWGSL --> MegaBuf
    FixWGSL --> RootBuf
    FixWGSL --> ResTable

1. Conquering 64-Bit Pointers: Slang's Generic GpuPtr<T> + Build-Time Shadow Headers

In the native Vulkan and Metal backends, shared headers like cube_shared.h declare raw pointers directly inside root structs:

struct CubeRootArguments
{
    gpu_float3* vertices;
    uint32 texture_index;
    uint32 sampler_index;
    float4x4 mvp;
};

As Sebastian Aaltonen emphasized, forcing users to rewrite their shaders with ugly macros like LOAD_FLOAT3(mega_buffer, root.vertices_offset + vertex_id * 12) would ruin the developer experience.

Our zero-macro solution works in two steps:

Step 1: Define a generic smart-pointer struct GpuPtr<T> in include/NoGraphicsAPI/shader.slang using Slang's __subscript property.
As a modern shading language, Slang supports overloading the subscript operator (__subscript with get and set accessors) as well as pointer arithmetic (operator+) on user-defined generic types:

#elif defined(NOGRAPHICSAPI_WEBGPU)

#if defined(NOGRAPHICSAPI_STAGE_COMPUTE)
[[vk::binding(0, 0)]] RWByteAddressBuffer __mega_buffer;
#else
[[vk::binding(0, 0)]] ByteAddressBuffer __mega_buffer;
#endif

struct GpuPtr<T>
{
#if defined(__EMSCRIPTEN__)
    uint offset;     // 4-byte pointer on wasm32 (high 4 bits: heap_id, low 28 bits: byte offset)
#else
    uint64_t offset; // 8-byte pointer on 64-bit native platforms
#endif

    __subscript(uint index) -> T
    {
        get
        {
            uint byte_offset = uint(offset & 0x0FFFFFFFu) + index * sizeof(T);
            return __mega_buffer.Load<T>(byte_offset);
        }
#if defined(NOGRAPHICSAPI_STAGE_COMPUTE)
        [nonmutating]
        set
        {
            uint byte_offset = uint(offset & 0x0FFFFFFFu) + index * sizeof(T);
            __mega_buffer.Store<T>(byte_offset, newValue);
        }
#endif
    }
};

GpuPtr<T> operator+<T>(GpuPtr<T> ptr, uint count)
{
    GpuPtr<T> result;
    result.offset = ptr.offset + count * sizeof(T);
    return result;
}

Step 2: Transparent Build-Time Shadow Header Rewriting.
When CMake compiles Slang shaders (NoGraphicsAPI_slang.cmake), it copies dependent *_shared.h headers into a temporary .shadow build directory and runs a 6-line Python regex script (rewrite_bda.py) that transforms struct member declarations from T* field; into GpuPtr<T> field; before invoking slangc -target wgsl.

As a result:

• The C++ compiler still sees genuine T* raw pointers (storing virtual GPU byte offsets on the CPU);
• The Slang compiler sees GpuPtr<T>, automatically lowering root.vertices[i] and compute shader writes like root.objects[id] = object into typed Load<T> and Store<T> operations against the global @binding(0) Mega-Buffer;
• Business shaders and shared headers remain 100% macro-free and untouched!

---

2. Banishing Manual _pad Fields: Solving Cross-Platform Struct Alignment with ByteAddressBuffer + -fvk-use-c-layout

While bringing up cube and deferred_renderer on WebAssembly, we hit a subtle cross-architecture struct alignment trap.

Look closely at CubeRootArguments again:

struct CubeRootArguments
{
    gpu_float3* vertices; // 8 bytes on 64-bit (offset 0..7); 4 bytes on wasm32 (offset 0..3)
    uint32 texture_index; // 4 bytes (offset 8..11 on 64-bit; offset 4..7 on wasm32)
    uint32 sampler_index; // 4 bytes (offset 12..15 on 64-bit; offset 8..11 on wasm32)
    float4x4 mvp;         // float[16] in C++, requiring only 4-byte alignment!
};

• On 64-bit macOS Native: vertices (8B) + texture_index (4B) + sampler_index (4B) = 16 bytes, placing mvp at byte offset 16;
• On 32-bit Emscripten (wasm32): A C++ pointer gpu_float3* is only 4 bytes! The first three fields total 12 bytes, so C++ places mvp immediately at byte offset 12!
• Worse still: if GPU_ROOT is declared in WGSL as a StructuredBuffer<CubeRootArguments> (var<storage> root: CubeRootArguments), WGSL's std430-like layout rules mandate 16-byte alignment for mat4x4<f32>. The WGSL shader will stubbornly read mvp from offset 16, shifting the matrix by 4 bytes on WASM and corrupting the entire transform!

The Initial Detour & Architectural Reflection:
Initially, an agent worked around this by inserting explicit uint32 _pad[2] / _pad[3] fields into cube_shared.h and gbuffer_shared.h.
During review, I immediately rejected this approach: alignment padding must be handled transparently when consumed by the backend—never leaked into user headers! If users have to manually calculate 32-bit vs. 64-bit pointer sizes and WGSL alignment rules just to declare a root struct, the ergonomics of NoGraphicsAPI are defeated.

The Clean Solution:

1. In include/NoGraphicsAPI/shader.slang, we changed GPU_ROOT on WebGPU from StructuredBuffer<type> to ByteAddressBuffer + .Load<type>(0):

      #define GPU_ROOT_DEFAULT(type, name) \
          [[vk::binding(1, 0)]] ByteAddressBuffer name##_storage; \
          static type name = name##_storage.Load<type>(0)
      

   Because ByteAddressBuffer lowers to a raw array<u32> in WGSL, it completely bypasses WGSL's 16-byte struct member alignment restriction on mat4x4<f32>!
2. Paired with slangc's -fvk-use-c-layout flag and GpuPtr<T> switching its internal offset field between uint (4B on __EMSCRIPTEN__) and uint64_t (8B on 64-bit native), Slang automatically emits u32 word loads at the exact C struct byte offsets for each target platform (offset 16 on 64-bit, offset 12 on wasm32).

With this in place, every single _pad field was removed from the example headers, restoring a clean zero-diff against main across all examples/*_shared.h files.

---

3. Unified Bindless Resource Table & Transparent Mesh Shader Lowering for 262k Objects

With pointers and root structs solved, two core pieces remained:

(1) Single-Table Dual-Region Bindless Mapping & WGSL's var Handle Restriction

[ForceInline]
T gpu_texture<T : IOpaqueDescriptor>(uint index)
{
    __requirePrelude(R"(enable chromium_experimental_resource_table;)");
    __intrinsic_asm "getResource<$TR>($0)";
}

[ForceInline]
T gpu_sampler<T : IOpaqueDescriptor = SamplerState>(uint index)
{
    __requirePrelude(R"(enable chromium_experimental_resource_table;)");
    __intrinsic_asm "getResource<$TR>($0 + 32768u)";
}

(2) Transparent Lowering from Mesh Shaders to Programmable Vertex Pulling

To run deferred_renderer without touching deferred_renderer.cpp, we capitalized on the fact that NoGraphicsAPI already stores all geometry and instance data in the global GpuHeap (Mega-Buffer):

• At build time, when NOGRAPHICSAPI_USE_WEBGPU is active and the shader stage is mesh, CMake compiles the entry point as a programmable vertex-pulling vertex stage;
• At runtime, draw_meshlets(cmd, group_x, group_y, group_z) translates to wgpuRenderPassEncoderDraw(meshlet_vertex_count, group_x * group_y * group_z, 0, 0), fetching per-meshlet vertices and instance transforms directly from the Mega-Buffer via @builtin(vertex_index) and @builtin(instance_index).

---

4. Three "Devil-in-the-Details" Pitfalls During Live Bring-Up

Once all three examples compiled and rendered their first frames, the most educational debugging challenges emerged when comparing visual output and lifecycle behavior across native Metal 4, macOS Dawn Native, and GFX_Runtime WASM Wire.

Pitfall 1: Why Was WebGPU Upside-Down—and Why Did Flipping ddx/ddy in cube.slang Turn Native Metal Pitch-Black?

When example_cube first rendered on WebGPU, two puzzling artifacts appeared:

1. Compared to the native Metal 4 build, the WebGPU triangle and cube were vertically flipped;
2. To fix the cube's unlit, pitch-black faces on WebGPU, an agent initially swapped the screen-space derivative cross product in cube.slang from cross(ddx(pos), ddy(pos)) to cross(ddy(pos), ddx(pos)). That made WebGPU light up properly—but when we recompiled the native Metal 4 backend, the Metal cube turned pitch-black instead!

Root Cause Analysis:
Both symptoms stemmed from a single underlying cause!

• NoGraphicsAPI's math conventions follow Vulkan window-space coordinates (where NDC $y\; =\; -1$ is at the top of the screen and $y\; =\; +1$ is at the bottom);
• In the native Metal backend (NoGraphicsAPIMetal.mm), the library flips the Y-axis during rasterization by passing a negative viewport height: (x, y + height, width, -height);
• WebGPU, however, explicitly forbids negative viewport heights! Without that viewport flip, not only was the image upside-down, but the screen-space partial derivative along $y\_w$ (ddy(input.world_pos)) also inverted its sign! By the right-hand rule, cross(ddx, ddy) produced inward-facing surface normals, clamping the Lambertian dot product to zero.

The Zero-Intrusion Fix:
Modifying cube.slang was off the table. Instead, in fix_wgsl.py, we automatically inject a Y-axis negation right before every @vertex entry point returns its @builtin(position):

var _nga_ret = ...;
_nga_ret.position.y = -_nga_ret.position.y;
return _nga_ret;

---

Pitfall 2: Dawn Wire's "Poison Pill" Mechanism — Why WGPUDawnTogglesDescriptor Cannot Be Passed into Emscripten

Because ChromiumExperimentalSamplingResourceTable is marked FeatureState::Experimental in Dawn, PhysicalDeviceBase::ValidateFeatureSupportedWithToggles enforces a dual gate:

1. The client must request the feature in WGPUDeviceDescriptor::requiredFeatures;
2. The underlying instance/device must enable the "allow_unsafe_apis" toggle via WGPUDawnTogglesDescriptor, or device creation fails with: Feature chromium-experimental-sampling-resource-table is guarded by toggle allow_unsafe_apis.

Early on, to bypass this gate quickly, an agent had mutated FeatureInfo in Dawn's Features.cpp to Stable and injected features inside GFX_Runtime. I asked to revert both: Dawn's feature stability metadata should remain untouched, and GFX_Runtime should only unlock the sandbox toggle (allow_unsafe_apis) on the host Instance—leaving feature requests strictly up to the client application (NoGraphicsAPIWebGPU.cpp).

That raised an interesting question:

"In client code (NoGraphicsAPIWebGPU.cpp), can we pass WGPUDawnTogglesDescriptor (allow_unsafe_apis) unconditionally on both Native and Emscripten without an #if !defined(__EMSCRIPTEN__) guard? Under Emscripten Wire mode, would it simply be a harmless no-op, or would it crash?"

Tracing Dawn Wire's serialization source revealed that Dawn Wire is far from passive—instead of silently ignoring unknown native structs, it actively injects a "poison pill" to force server-side validation failure!

flowchart TD
    subgraph CaseA ["Case A: Chained on WGPUInstanceDescriptor.nextInChain"]
        A1["WASM Client calls wgpuCreateInstance"] --> A2["Wire Client inspects nextInChain"]
        A2 --> A3["Unrecognized sType hits default branch\nReturns WireResult::FatalError (nullptr)"]
    end

    subgraph CaseB ["Case B: Chained on WGPUDeviceDescriptor.nextInChain"]
        B1["WASM Client calls wgpuAdapterRequestDevice"] --> B2["Wire Client serializes DeviceDescriptor"]
        B2 --> B3["toggles carries 'native' tag (omitted from Wire switch)\nActively injects poison pill: DawnInjectedInvalidSType"]
        B3 --> B4["Host Wire Server unpacks DawnInjectedInvalidSType\nImmediately throws DAWN_VALIDATION_ERROR & rejects Device"]
    end

1. If chained on wgpuCreateInstance: dawn::wire::client::Instance::Initialize only whitelists two WGSL control structs in its switch (chain->sType). Encountering WGPUDawnTogglesDescriptor hits default: and returns WireResult::FatalError, causing wgpuCreateInstance to return nullptr immediately;
2. If chained on wgpuAdapterRequestDevice: Because "dawn toggles descriptor" is tagged "native" in dawn.json, dawn_json_generator.py excludes it from WireCmd_autogen.cpp. When the wire client serializer encounters an unrecognized sType in nextInChain, it deliberately serializes a wgpu::SType::DawnInjectedInvalidSType transfer record into the wire buffer! When the host server unpacks the chain and sees DawnInjectedInvalidSType, it immediately raises DAWN_VALIDATION_ERROR("Unexpected chained struct of type...") and aborts device creation!

This design reflects Chromium and Dawn's strict security boundary: an untrusted Web/WASM client across the wire must never be allowed to override host-level native driver toggles via nextInChain!
Thus, the clean architectural separation is:

• Host Runtime (GFX_Runtime): Enables the allow_unsafe_apis toggle when constructing the native dawn::native::Instance, unlocking the gate for experimental features without force-enabling any specific feature;
• Client Backend (NoGraphicsAPIWebGPU.cpp): Chains WGPUDawnTogglesDescriptor only under #if !defined(__EMSCRIPTEN__) (for standalone native execution), while unconditionally requesting WGPUFeatureName_ChromiumExperimentalSamplingResourceTable in WGPUDeviceDescriptor::requiredFeatures across both Native and WASM.

---

Pitfall 3: The Mystery exit code: 6 (SIGABRT) Caused by a Partial Static Library Rebuild

After reverting Features.cpp in the Dawn repository, a quick incremental build (ninja dawn_native) was run and only libdawn_native.a was copied into GFX_Runtime. Launching GFX_Runtime.app resulted in an instant crash right after the first log line:

[Native] DawnBridgeCore initializing, cleared old log file.
Program ended with exit code: 6

Why did the app abort (SIGABRT, exit code 6) at startup when no code was broken?

Inspecting the generated dawn_proc.cpp revealed the culprit: Dawn 8103 enforces a strict Git Commit SHA1 version assertion inside dawnProcSetProcs:

void dawnProcSetProcs(const DawnProcTable* procs_) {
    if (procs_) {
        DAWN_CHECK(std::ranges::equal(procs_->version, dawn::kDawnVersion));
        procs = *procs_;
    }
}

Running the full build_dawn_mac.sh and copy_dawn_mac.sh pipeline to synchronize all static libraries resolved the crash immediately.

---

5. Final Results & Performance Benchmarks

With Google Dawn's Metal Bindless backend, emdawnwebgpu's automated wire pass-throughs, and NoGraphicsAPI's WebGPU backend all in place, the entire end-to-end stack is now complete.

We executed a comprehensive 159-test, 4-tier end-to-end verification suite (ranging from Tier 1 feature unit tests to Tier 4 live window framebuffer pixel verification) across Apple Silicon macOS Native and GFX_Runtime, achieving a 159 / 159 (100% pass) rate:

Looking back at this journey, the most rewarding outcome isn't just watching 262,144 dynamic lit cubes fly at a locked 72 FPS inside a WebAssembly container—it's proving that with thoughtful compiler lowering (Slang + Tint) and targeted runtime engineering (Dawn ResourceTable + Wire), we can honor every Web security sandbox invariant while bringing NoGraphicsAPI's zero-ceremony, pointer-driven programming model cleanly into the WebGPU era.