WebGPU Compute, WGSL Shaders & GPU Memory Architecture Studio

WebGL 2.0 is a legacy wrapper around OpenGL ES 3.0, designed exclusively for rasterization graphics with no native support for general-purpose compute shaders, storage buffers, or modern GPU hardware queues. To perform parallel math (GPGPU) in WebGL, developers had to render textured quads to offscreen framebuffers (FBOs) and encode numbers as 32-bit RGBA pixel colors. W3C WebGPU is a ground-up redesign modeled directly after modern explicit graphics APIs (Vulkan, Apple Metal, and Microsoft DirectX 12). WebGPU introduces first-class compute pipelines (GPUComputePipeline), writable storage buffers (GPUBufferUsage.STORAGE), and direct access to GPU workgroup shared memory. This enables browser-based large language model (LLM) inference, physical simulations, and real-time computer vision with near-native metal performance.

GPU parallel execution is structured in three nested dimensions: (1) Invocations (Threads): The smallest unit of execution running the WGSL shader function. (2) Workgroups: A cluster of invocations declared in WGSL via @workgroup_size(x, y, z) (e.g. 16, 16, 1 = 256 threads). All invocations within a single workgroup execute concurrently on the same Streaming Multiprocessor (SM) / Compute Unit (CU), share high-speed local memory (var), and can synchronize execution via workgroupBarrier(). (3) Dispatch Grid: The global problem space dispatched via passEncoder.dispatchWorkgroups(gx, gy, gz). The total number of threads dispatched across the GPU is (gx * x) * (gy * y) * (gz * z). At the physical hardware level, GPUs group invocations into lockstep SIMD units called Warps (32 threads on Nvidia) or Wavefronts (32/64 threads on AMD).

Modern GPU VRAM is accessed through high-bandwidth memory controllers in wide cache-line transactions (typically 128 bytes per transaction). When all 32 threads in a warp access consecutive, aligned 4-byte floating-point memory addresses (e.g. thread i reads array[i]), the hardware memory controller coalesces all 32 individual reads into a single 128-byte DRAM burst (100% bus utilization). Conversely, if threads perform strided access (e.g. thread i reads column i in a row-major matrix with stride N), each thread memory request falls into a different 128-byte cache line. The GPU memory controller is forced to issue 32 separate 128-byte memory transactions to retrieve only 128 bytes of useful data (3.1% bus efficiency), causing catastrophic memory pipeline stalls.

A naive matrix multiplication ($C = A imes B$) computes the dot product of row i and column j, requiring $2N^3$ global memory reads from high-latency VRAM (200-400 cycles latency per read). Tiled Matrix Multiplication partitions the input matrices into small square tiles (e.g., 16x16 elements) that fit within high-speed on-chip workgroup shared memory (var, ~2-4 cycles latency). The 256 threads in a workgroup collaboratively load one tile of matrix A and one tile of matrix B from VRAM into shared memory using coalesced reads, execute a workgroupBarrier() to ensure all tile bytes are written, and then compute dot products entirely from shared memory. This reduces global memory traffic by a factor of the tile size (16x reduction), shifting the algorithm from memory-bound to compute-bound.

WGSL enforces strict memory layout rules for structs bound to host memory. For Uniform Buffers (which follow the legacy Std140 layout), vec3 has an alignment requirement of 16 bytes (the size of a vec4), even though its byte size is only 12 bytes. If a developer defines a WGSL struct containing struct Data { a: vec3, b: f32 }, the field b will be placed at offset 12 in C/JavaScript Float32Array, but the GPU expects b at offset 16 due to the 16-byte alignment of vec3. Storage Buffers follow the newer Std430 rules, which allow tighter packing for arrays of scalars, but still enforce strict natural alignment. Passing unaligned JavaScript TypedArrays without accounting for WGSL padding offsets causes silent data misalignment and garbled calculation outputs.

GPU VRAM is physically located on the graphics card PCI Express bus, which is asynchronous from the CPU clock domain. Allowing JavaScript to synchronously read VRAM would force the CPU to stall for hundreds of thousands of clock cycles while waiting for the PCIe bus transfer. Furthermore, modern operating systems enforce strict memory security: VRAM buffers cannot be mapped into user-space CPU memory while actively bound to GPU command queues. To read back compute results, WebGPU requires a two-step staging protocol: (1) Encode a GPU command to copy data from the GPU storage buffer to a staging buffer created with GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST; (2) Submit the command buffer and call stagingBuffer.mapAsync(GPUMapMode.READ). Once the returned Promise resolves, JavaScript can call stagingBuffer.getMappedRange() to inspect the results, followed immediately by stagingBuffer.unmap().