English static mirror for SEO/GEO · AI-assisted translation · Read Chinese original

Born Appendix B: WGSL Operator Catalog for the WebGPU Backend

Forum topic · QianXun · 2026-06-13

Summary

This appendix from the serial technical book 'Born' catalogs the WebGPU backend's embedded WGSL compute shaders: 53 shaders across 9 operator categories. Element-wise binary operations (add, sub, mul, div) use workgroup_size 256, while unary ops include neg, exp, log, sqrt, relu, sigmoid, and silu. Matrix operations cover standard matmul (16,16 workgroup) and batched matmul (8,8,1). Convolution and pooling are handled by conv2d and maxPool2d shaders with (16,16) workgroups. Reduction ops include global sum, cross-workgroup sum, per-dimension sum, and argmax. Comparison and logic shaders implement >, <, ==, !=, and boolean operations. A tiled Flash Attention shader (8,8,1 workgroup) rounds out the list. The document serves as a reference index for a from-scratch GPU inference library written in WGSL.

Born's WebGPU backend includes 53 embedded WGSL compute shaders covering 9 operator categories.

Element-wise Binary Operations (workgroup_size = 256)

| Shader | Expression | |--------|------------| | addShader | result[i] = a[i] + b[i] | | subShader | result[i] = a[i] - b[i] | | mulShader | result[i] = a[i] * b[i] | | divShader | result[i] = a[i] / b[i] |

Element-wise Unary Operations

| Shader | Expression | |--------|------------| | negShader | -a[i] | | expShader | exp(a[i]) | | logShader | log(a[i]) | | sqrtShader | sqrt(a[i]) | | reluShader | max(0.0, a[i]) | | sigmoidShader | 1.0 / (1.0 + exp(-a[i])) | | siluShader | a[i] / (1.0 + exp(-a[i])) |

Matrix Operations

| Shader | Workgroup | Description | |--------|-----------|-------------| | matmulShader | (16,16) | Standard matrix multiplication | | batchMatMulShader | (8,8,1) | Batched matrix multiplication |

Convolution and Pooling

| Shader | Workgroup | Description | |--------|-----------|-------------| | conv2dShader | (16,16) | 2D convolution | | maxPool2dShader | (16,16) | Max pooling |

Reduction Operations

| Shader | Description | |--------|-------------| | sumShader | Global sum | | globalSumShader | Cross-workgroup sum | | sumDimShader | Sum along an axis | | argmaxShader | Argmax along an axis |

Comparison and Logic

| Shader | Description | |--------|-------------| | greaterShader, lowerShader | >, < | | equalShader, notEqualShader | ==, != | | andShader, orShader, notShader | Logical operations |

Flash Attention

| Shader | Workgroup | Description | |--------|-----------|-------------| | flashAttentionShader | (8,8,1) | Tiled Flash Attention |

---

Total: 53 shaders

📘 *Born* is a serialized technical book. This is Appendix B, part 4.

Tags

#webgpu#wgsl#compute-shaders#flash-attention#deep-learning#gpu-compute#operators#born

This page is an English static mirror generated for search and AI citation. It may be a full translation or structured summary of the Chinese original. Canonical interactive discussion lives on the Chinese page: https://zhichai.net/topic/177981256