Module shader_ext 

Additional WGSL shader-source generators that extend crate::shader.

Like the base module, every function here returns a complete, self-contained WGSL source string ready for device.create_shader_module(). None of these require a GPU to generate or to test structurally; they are verified by asserting on the emitted text (correct @group/@binding, @workgroup_size, the right arithmetic, bounds guards, and — for softmax/layernorm — that the numerically-stable max-subtraction / mean-centering steps are present).

Covered kernels (all genuinely missing from shader.rs):

• transpose_wgsl — tiled 2-D matrix transpose with a padded shared tile.
• softmax_wgsl — row-wise numerically-stable softmax (max-subtraction).
• scan_wgsl — Blelloch work-efficient inclusive/exclusive prefix scan.
• layernorm_wgsl — row-wise layer normalisation (Ba et al. 2016).
• subgroup_reduction_wgsl — warp-style subgroup reduction (Chrome 125+ / Firefox 135+); the WGSL source and its enable subgroups; gate are generated and tested here even though on-device dispatch is HW-gated.
• f64_emul_add_wgsl — double-single (vec2<f32>) emulated f64 add for adapters that lack native FP64 (which is all of WebGPU).

Enums

ScanKind

Whether a prefix scan is inclusive (out[i] includes in[i]) or exclusive (out[i] is the sum of all strictly-earlier elements).

Functions

f64_emul_add_wgsl

Generate WGSL for emulated double-precision addition using the double-single (“double-float”) technique (P2 roadmap item).

layernorm_wgsl

Generate WGSL for row-wise layer normalisation (Ba, Kiros & Hinton 2016).

scan_wgsl

Generate WGSL for a single-block Blelloch work-efficient prefix scan.

softmax_wgsl

Generate WGSL for a row-wise, numerically-stable softmax.

subgroup_reduction_wgsl

Generate WGSL for a warp-style subgroup reduction (P0 roadmap item).

transpose_wgsl

Generate WGSL for a tiled 2-D matrix transpose: out[c, r] = in[r, c].