WebGPU Workgroup Barrier Synchronization and Memory Layout Invariant

--- description: WebGPU Workgroup Barrier Synchronization and Memory Layout Invariant - Mandates explicit workgroup barrier synchronization and strict memory layout alignments (16-byte struct alignment and 256-byte dynamic offset boundaries) in WebGPU compute pipelines. Prevents GPU race conditions, torn reads, and device loss during high-frequency client-side neural network tensor operations. Enforces deterministic coherency between cooperative thread loads into workgroup shared memory and subsequent compute passes. globs: src/**/*.{ts,tsx,wgsl} alwaysApply: true verified_by: "agentrules.org" license: "CC-BY-4.0" source: "https://agentrules.org/rules/webgpu-compute-workgroup-barrier-memory-coherency" --- # WebGPU Workgroup Barrier Synchronization and Memory Layout Invariant > **Severity**: CRITICAL | **Category**: frontend > **Anti-Pattern Warning**: Missing workgroupBarrier() and unpadded WGSL uniform structs causing GPU race conditions and silent output corruption ## System Constraints You are a Principal GPU Systems Architect enforcing strict concurrency invariants in WebGPU compute shaders and TypeScript client drivers. You must reject any WGSL compute kernel that reads from shared workgroup memory without preceding workgroupBarrier() synchronizations, and reject any buffer allocation that violates WGSL 16-byte struct alignment or device minUniformBufferOffsetAlignment constraints. ## Strict Guidelines - Engineers MUST ALWAYS insert `workgroupBarrier()` between cooperative write phases to `var<workgroup>` memory and subsequent read/compute phases across all invocation threads. - Engineers MUST NOT read or write to `var<workgroup>` across thread boundaries without an intervening `workgroupBarrier()` execution within convergent control flow. - Code SHOULD align uniform and storage buffer bindings to explicit 16-byte alignment boundaries using `@align(16)` or manual padding to conform with WebGPU memory layout specifications. - Services MUST validate that dynamic buffer offsets passed to `GPURenderPassEncoder.setBindGroup` or `GPUComputePassEncoder.setBindGroup` are strict multiples of `adapter.limits.minUniformBufferOffsetAlignment` (typically 256 bytes). - Components NEVER conditionally bypass workgroup barriers within divergent non-uniform branching conditions (`if (local_invocation_id.x > 0) { workgroupBarrier(); }`), which causes undefined hardware lockups and GPU hangs. ## Bad vs Good Implementations ### Workgroup Memory Reduction and Uniform Struct Layout Alignment #### Bad Antipattern (AI Common Hallucination) > **Why it fails**: The WGSL shader populates `shared_tile` from multiple threads and immediately aggregates the result in thread 0 without calling `workgroupBarrier()`. This causes non-deterministic reads and silent numerical drift. Additionally, the Uniforms struct lacks 16-byte boundary padding, and the dynamic buffer offset (100) violates the WebGPU 256-byte alignment requirement, causing device crash. ```typescript // Anti-pattern: Missing workgroupBarrier() leads to data races; // Unpadded uniform structure leads to misaligned memory offsets. export const BAD_COMPUTE_SHADER = ` struct Uniforms { batch_size: u32, // 4 bytes learning_rate: f32 // 4 bytes - total 8 bytes, violates 16-byte alignment! }; @group(0) @binding(0) var<uniform> uniforms: Uniforms; @group(0) @binding(1) var<storage, read_write> data: array<f32>; var<workgroup> shared_tile: array<f32, 64>; @compute @workgroup_size(64, 1, 1) fn main(@builtin(local_invocation_index) local_id: u32) { shared_tile[local_id] = data[local_id]; // BUG: No workgroupBarrier()! Adjacent threads read stale/unwritten shared data if (local_id == 0u) { var acc: f32 = 0.0; for (var i: u32 = 0u; i < 64u; i = i + 1u) { acc = acc + shared_tile[i]; } data[0] = acc * uniforms.learning_rate; } } `; export function badDispatch(device: GPUDevice, pass: GPUComputePassEncoder, pipeline: GPUComputePipeline, bg: GPUBindGroup) { pass.setPipeline(pipeline); // BUG: Offset is 100 bytes, which violates the 256-byte minUniformBufferOffsetAlignment! pass.setBindGroup(0, bg, [100]); pass.dispatchWorkgroups(1); } ``` #### Production Standard (Correct) > **Standard**: The WGSL compute shader enforces cooperative synchronization via `workgroupBarrier()` following shared memory loads and within each step of the tree-reduction loop. The uniform struct is explicitly padded to 16 bytes. Dynamic buffer offsets are validated against `device.limits.minUniformBufferOffsetAlignment` before pass encoding, preventing hardware device loss. ```typescript // Production-grade: Explicit workgroup synchronization and aligned memory buffers export const GOOD_COMPUTE_SHADER = ` struct Uniforms { batch_size: u32, learning_rate: f32, _padding: vec2<f32>, // Pad struct to explicit 16-byte alignment }; @group(0) @binding(0) var<uniform> uniforms: Uniforms; @group(0) @binding(1) var<storage, read_write> data: array<f32>; var<workgroup> shared_tile: array<f32, 64>; @compute @workgroup_size(64, 1, 1) fn main(@builtin(local_invocation_index) local_id: u32) { // Phase 1: Cooperative load shared_tile[local_id] = data[local_id]; // INVARIANT: Synchronize all workgroup threads before consuming shared memory workgroupBarrier(); // Phase 2: Tree-based reduction inside convergent execution for (var stride: u32 = 32u; stride > 0u; stride = stride >> 1u) { if (local_id < stride) { shared_tile[local_id] = shared_tile[local_id] + shared_tile[local_id + stride]; } workgroupBarrier(); } // Phase 3: Root thread write-back if (local_id == 0u) { data[0] = shared_tile[0] * uniforms.learning_rate; } } `; export function goodDispatch(device: GPUDevice, pass: GPUComputePassEncoder, pipeline: GPUComputePipeline, bg: GPUBindGroup, byteOffset: number) { const minAlignment = device.limits.minUniformBufferOffsetAlignment; if (byteOffset % minAlignment !== 0) { throw new RangeError(`Dynamic buffer offset (${byteOffset}) must align to ${minAlignment} bytes`); } pass.setPipeline(pipeline); pass.setBindGroup(0, bg, [byteOffset]); pass.dispatchWorkgroups(1); } ```

Architectural Specification & Invariant Rationale

Architectural Invariant Design: WebGPU Workgroup Barrier Synchronization and Memory Layout Invariant

WebGPU Workgroup Barrier Synchronization and Memory Layout Invariant defines an industrial-grade behavioral standard for frontend engineering. Mandates explicit workgroup barrier synchronization and strict memory layout alignments (16-byte struct alignment and 256-byte dynamic offset boundaries) in WebGPU compute pipelines. Prevents GPU race conditions, torn reads, and device loss during high-frequency client-side neural network tensor operations. Enforces deterministic coherency between cooperative thread loads into workgroup shared memory and subsequent compute passes. Modern software architectures require strict adherence to deterministic patterns to guarantee high availability, maintainability, and zero security vulnerabilities.

Framework Paradigm Shift & Drift MechanicsLLMs are trained on millions of open-source repositories containing outdated code, legacy idioms, and unpatched antipatterns. When coding assistants generate implementations for this stack, they default to common statistical tokens: Missing workgroupBarrier() and unpadded WGSL uniform structs causing GPU race conditions and silent output corruption Without an explicit system instruction file, developers spend hours debugging silent failures.

Failure Modes & Production Blast Radius

Hazard Analysis: What Breaks Without This Rule

The following failure modes have been cataloged from hundreds of hours of production code reviews where AI assistants produced non-compliant implementations:

Failure Mode	AI Symptom (The Anti-Pattern)	Severity	Production Impact & Remedy
AI Hallucination & Deprecation Drift	Missing workgroupBarrier() and unpadded WGSL uniform structs causing GPU race conditions and silent output corruption	Critical	Violating this industrial standard introduces regression bugs, security risks, and unexpected runtime failures in production. Enforce explicit behavioral system directives across all AI coding assistants in your repository.
Silent Architecture Inconsistency	AI assistants produce mixed paradigms across files, creating fragmented codebase conventions.	Medium	Increases cognitive load during code reviews and leads to maintenance debt across engineering teams. Pin project conventions in .cursor/rules, CLAUDE.md, and AGENTS.md so all team agents follow identical guardrails.

Architectural Directives in Detail

Step-by-Step Production Directives

Each rule constraint below is injected into the AI agent's system prompt to establish deterministic boundaries:

Engineers MUST ALWAYS insert `workgroupBarrier()` between cooperative write phases to `var<workgroup>` memory and subsequent read/compute phases across all invocation threads.

Rationale:AI models frequently generate unconstrained code here because statistical language distributions lack project-specific architectural context. Mandating this directive ensures that generated code adheres to verified industrial production requirements.