diff --git a/apps/example/src/App.tsx b/apps/example/src/App.tsx
index cc85c03b5..821b30d6e 100644
--- a/apps/example/src/App.tsx
+++ b/apps/example/src/App.tsx
@@ -16,6 +16,7 @@ import {
Cubemap,
} from "./Cube";
import { HelloTriangle, HelloTriangleMSAA } from "./Triangle";
+import { ImageBlur } from "./ImageBlur";
import { RenderBundles } from "./RenderBundles";
import { ABuffer } from "./ABuffer";
import { OcclusionQuery } from "./OcclusionQuery";
@@ -73,6 +74,7 @@ function App() {
name="HelloTriangleMSAA"
component={HelloTriangleMSAA}
/>
+
diff --git a/apps/example/src/Home.tsx b/apps/example/src/Home.tsx
index fd72751ad..85ece3978 100644
--- a/apps/example/src/Home.tsx
+++ b/apps/example/src/Home.tsx
@@ -19,6 +19,10 @@ export const examples = [
screen: "HelloTriangleMSAA",
title: "🔺 Hello Triangle MSAA",
},
+ {
+ screen: "ImageBlur",
+ title: "🫧 Image Blur",
+ },
{
screen: "Reanimated",
title: "🐎 Reanimated",
diff --git a/apps/example/src/ImageBlur/ImageBlur.tsx b/apps/example/src/ImageBlur/ImageBlur.tsx
new file mode 100644
index 000000000..2286f3196
--- /dev/null
+++ b/apps/example/src/ImageBlur/ImageBlur.tsx
@@ -0,0 +1,319 @@
+import React, { useEffect, useRef, useState } from "react";
+import { PixelRatio, Pressable, StyleSheet, Text, View } from "react-native";
+import { Canvas, type CanvasRef, useDevice } from "react-native-webgpu";
+
+import { decodeImage } from "../components/useAssets";
+
+import { createBlurPassPlan } from "./blurPlan";
+import { BLUR_SHADER, FULLSCREEN_TEXTURE_SHADER } from "./Shaders";
+
+const FILTER_SIZES = [5, 15, 31] as const;
+const ITERATIONS = [1, 2, 4] as const;
+
+const nextOption = (options: readonly T[], current: T) => {
+ const index = options.indexOf(current);
+ return options[(index + 1) % options.length];
+};
+
+export const ImageBlur = () => {
+ const { device } = useDevice();
+ const canvasRef = useRef(null);
+ const [filterSize, setFilterSize] = useState(15);
+ const [iterations, setIterations] = useState(2);
+ const settingsRef = useRef({ filterSize, iterations });
+ const renderRef = useRef<(() => void) | null>(null);
+ settingsRef.current = { filterSize, iterations };
+
+ useEffect(() => {
+ if (!device) {
+ return;
+ }
+
+ let cancelled = false;
+ let render: (() => void) | null = null;
+ const buffers: GPUBuffer[] = [];
+ const textures: GPUTexture[] = [];
+
+ const trackBuffer = (buffer: GPUBuffer) => {
+ buffers.push(buffer);
+ return buffer;
+ };
+ const trackTexture = (texture: GPUTexture) => {
+ textures.push(texture);
+ return texture;
+ };
+ const createFlipBuffer = (value: number) => {
+ const buffer = trackBuffer(
+ device.createBuffer({
+ size: 4,
+ mappedAtCreation: true,
+ usage: GPUBufferUsage.UNIFORM,
+ }),
+ );
+ new Uint32Array(buffer.getMappedRange())[0] = value;
+ buffer.unmap();
+ return buffer;
+ };
+
+ const start = async () => {
+ const context = canvasRef.current?.getContext("webgpu");
+ if (!context) {
+ return;
+ }
+
+ const image = await decodeImage(require("../assets/Di-3d.png"));
+ if (cancelled) {
+ return;
+ }
+
+ const canvas = context.canvas as HTMLCanvasElement;
+ const pixelRatio = PixelRatio.get();
+ canvas.width = canvas.clientWidth * pixelRatio;
+ canvas.height = canvas.clientHeight * pixelRatio;
+
+ const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
+ context.configure({
+ device,
+ format: presentationFormat,
+ alphaMode: "premultiplied",
+ });
+
+ const blurModule = device.createShaderModule({ code: BLUR_SHADER });
+ const blurPipeline = device.createComputePipeline({
+ layout: "auto",
+ compute: { module: blurModule },
+ });
+ const fullscreenModule = device.createShaderModule({
+ code: FULLSCREEN_TEXTURE_SHADER,
+ });
+ const fullscreenPipeline = device.createRenderPipeline({
+ layout: "auto",
+ vertex: { module: fullscreenModule, entryPoint: "vertexMain" },
+ fragment: {
+ module: fullscreenModule,
+ entryPoint: "fragmentMain",
+ targets: [{ format: presentationFormat }],
+ },
+ primitive: { topology: "triangle-list" },
+ });
+
+ const sampler = device.createSampler({
+ magFilter: "linear",
+ minFilter: "linear",
+ });
+ const sourceTexture = trackTexture(
+ device.createTexture({
+ size: [image.width, image.height, 1],
+ format: "rgba8unorm",
+ usage:
+ GPUTextureUsage.TEXTURE_BINDING |
+ GPUTextureUsage.COPY_DST |
+ GPUTextureUsage.RENDER_ATTACHMENT,
+ }),
+ );
+ device.queue.copyExternalImageToTexture(
+ { source: image },
+ { texture: sourceTexture },
+ [image.width, image.height],
+ );
+
+ const intermediateTextures = [0, 1].map(() =>
+ trackTexture(
+ device.createTexture({
+ size: [image.width, image.height],
+ format: "rgba8unorm",
+ usage:
+ GPUTextureUsage.COPY_DST |
+ GPUTextureUsage.STORAGE_BINDING |
+ GPUTextureUsage.TEXTURE_BINDING,
+ }),
+ ),
+ );
+ const horizontalFlip = createFlipBuffer(0);
+ const verticalFlip = createFlipBuffer(1);
+ const blurParamsBuffer = trackBuffer(
+ device.createBuffer({
+ size: 8,
+ usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.UNIFORM,
+ }),
+ );
+
+ const computeConstants = device.createBindGroup({
+ layout: blurPipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: 0, resource: sampler },
+ { binding: 1, resource: { buffer: blurParamsBuffer } },
+ ],
+ });
+ const createComputeBindGroup = (
+ input: GPUTexture,
+ output: GPUTexture,
+ flip: GPUBuffer,
+ ) =>
+ device.createBindGroup({
+ layout: blurPipeline.getBindGroupLayout(1),
+ entries: [
+ { binding: 1, resource: input.createView() },
+ { binding: 2, resource: output.createView() },
+ { binding: 3, resource: { buffer: flip } },
+ ],
+ });
+ const computeBindGroups = [
+ createComputeBindGroup(
+ sourceTexture,
+ intermediateTextures[0],
+ horizontalFlip,
+ ),
+ createComputeBindGroup(
+ intermediateTextures[0],
+ intermediateTextures[1],
+ verticalFlip,
+ ),
+ createComputeBindGroup(
+ intermediateTextures[1],
+ intermediateTextures[0],
+ horizontalFlip,
+ ),
+ ] as const;
+ const showResultBindGroup = device.createBindGroup({
+ layout: fullscreenPipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: 0, resource: sampler },
+ { binding: 1, resource: intermediateTextures[1].createView() },
+ ],
+ });
+
+ let previousFilterSize = 0;
+ render = () => {
+ if (cancelled) {
+ return;
+ }
+
+ const settings = settingsRef.current;
+ const plan = createBlurPassPlan({
+ width: image.width,
+ height: image.height,
+ filterSize: settings.filterSize,
+ iterations: settings.iterations,
+ });
+ if (settings.filterSize !== previousFilterSize) {
+ device.queue.writeBuffer(
+ blurParamsBuffer,
+ 0,
+ new Uint32Array([plan.filterDim, plan.blockDim]),
+ );
+ previousFilterSize = settings.filterSize;
+ }
+
+ const commandEncoder = device.createCommandEncoder();
+ const computePass = commandEncoder.beginComputePass();
+ computePass.setPipeline(blurPipeline);
+ computePass.setBindGroup(0, computeConstants);
+ for (const pass of plan.passes) {
+ computePass.setBindGroup(1, computeBindGroups[pass.bindGroup]);
+ computePass.dispatchWorkgroups(pass.x, pass.y);
+ }
+ computePass.end();
+
+ const renderPass = commandEncoder.beginRenderPass({
+ colorAttachments: [
+ {
+ view: context.getCurrentTexture().createView(),
+ clearValue: [0, 0, 0, 1],
+ loadOp: "clear",
+ storeOp: "store",
+ },
+ ],
+ });
+ renderPass.setPipeline(fullscreenPipeline);
+ renderPass.setBindGroup(0, showResultBindGroup);
+ renderPass.draw(6);
+ renderPass.end();
+
+ device.queue.submit([commandEncoder.finish()]);
+ context.present();
+ };
+
+ renderRef.current = render;
+ render();
+ };
+
+ void start();
+ return () => {
+ cancelled = true;
+ if (renderRef.current === render) {
+ renderRef.current = null;
+ }
+ for (const buffer of buffers) {
+ buffer.destroy();
+ }
+ for (const texture of textures) {
+ texture.destroy();
+ }
+ };
+ }, [device]);
+
+ useEffect(() => {
+ renderRef.current?.();
+ }, [filterSize, iterations]);
+
+ return (
+
+
+
+
+ setFilterSize((current) => nextOption(FILTER_SIZES, current))
+ }
+ >
+ Kernel
+ {filterSize}
+
+
+ setIterations((current) => nextOption(ITERATIONS, current))
+ }
+ >
+ Passes
+ {iterations * 2}
+
+
+
+ );
+};
+
+const styles = StyleSheet.create({
+ container: {
+ flex: 1,
+ backgroundColor: "black",
+ },
+ controls: {
+ position: "absolute",
+ top: 24,
+ left: 16,
+ flexDirection: "row",
+ gap: 8,
+ },
+ button: {
+ minWidth: 84,
+ borderRadius: 12,
+ paddingHorizontal: 12,
+ paddingVertical: 8,
+ backgroundColor: "rgba(0,0,0,0.65)",
+ },
+ buttonLabel: {
+ color: "rgba(255,255,255,0.65)",
+ fontSize: 11,
+ textTransform: "uppercase",
+ },
+ buttonValue: {
+ color: "white",
+ marginTop: 2,
+ fontSize: 16,
+ fontWeight: "600",
+ },
+});
diff --git a/apps/example/src/ImageBlur/Shaders.ts b/apps/example/src/ImageBlur/Shaders.ts
new file mode 100644
index 000000000..69570a16b
--- /dev/null
+++ b/apps/example/src/ImageBlur/Shaders.ts
@@ -0,0 +1,112 @@
+// Adapted from the official WebGPU image blur sample:
+// https://github.com/webgpu/webgpu-samples/tree/main/sample/imageBlur
+export const BLUR_SHADER = /* wgsl */ `
+struct Params {
+ filterDim: i32,
+ blockDim: u32,
+}
+
+@group(0) @binding(0) var samp: sampler;
+@group(0) @binding(1) var params: Params;
+@group(1) @binding(1) var inputTex: texture_2d;
+@group(1) @binding(2) var outputTex: texture_storage_2d;
+
+struct Flip {
+ value: u32,
+}
+@group(1) @binding(3) var flip: Flip;
+
+var tile: array, 4>;
+
+@compute @workgroup_size(32, 1, 1)
+fn main(
+ @builtin(workgroup_id) workgroupId: vec3u,
+ @builtin(local_invocation_id) localInvocationId: vec3u,
+) {
+ let filterOffset = (params.filterDim - 1) / 2;
+ let dimensions = vec2i(textureDimensions(inputTex, 0));
+ let baseIndex = vec2i(
+ workgroupId.xy * vec2(params.blockDim, 4u) +
+ localInvocationId.xy * vec2(4u, 1u)
+ ) - vec2(filterOffset, 0);
+
+ for (var row = 0; row < 4; row++) {
+ for (var column = 0; column < 4; column++) {
+ var loadIndex = baseIndex + vec2(column, row);
+ if (flip.value != 0u) {
+ loadIndex = loadIndex.yx;
+ }
+ tile[row][4u * localInvocationId.x + u32(column)] = textureSampleLevel(
+ inputTex,
+ samp,
+ (vec2f(loadIndex) + vec2f(0.5)) / vec2f(dimensions),
+ 0.0
+ ).rgb;
+ }
+ }
+
+ workgroupBarrier();
+
+ for (var row = 0; row < 4; row++) {
+ for (var column = 0; column < 4; column++) {
+ var writeIndex = baseIndex + vec2(column, row);
+ if (flip.value != 0u) {
+ writeIndex = writeIndex.yx;
+ }
+ let center = i32(4u * localInvocationId.x) + column;
+ if (
+ center >= filterOffset &&
+ center < 128 - filterOffset &&
+ all(writeIndex < dimensions)
+ ) {
+ var color = vec3f(0.0);
+ for (var index = 0; index < params.filterDim; index++) {
+ let tileIndex = center + index - filterOffset;
+ color += tile[row][tileIndex] / f32(params.filterDim);
+ }
+ textureStore(outputTex, writeIndex, vec4f(color, 1.0));
+ }
+ }
+ }
+}
+`;
+
+export const FULLSCREEN_TEXTURE_SHADER = /* wgsl */ `
+@group(0) @binding(0) var imageSampler: sampler;
+@group(0) @binding(1) var imageTexture: texture_2d;
+
+struct VertexOutput {
+ @builtin(position) position: vec4f,
+ @location(0) uv: vec2f,
+}
+
+@vertex
+fn vertexMain(@builtin(vertex_index) vertexIndex: u32) -> VertexOutput {
+ const positions = array(
+ vec2f( 1.0, 1.0),
+ vec2f( 1.0, -1.0),
+ vec2f(-1.0, -1.0),
+ vec2f( 1.0, 1.0),
+ vec2f(-1.0, -1.0),
+ vec2f(-1.0, 1.0),
+ );
+ const uvs = array(
+ vec2f(1.0, 0.0),
+ vec2f(1.0, 1.0),
+ vec2f(0.0, 1.0),
+ vec2f(1.0, 0.0),
+ vec2f(0.0, 1.0),
+ vec2f(0.0, 0.0),
+ );
+
+ var output: VertexOutput;
+ output.position = vec4f(positions[vertexIndex], 0.0, 1.0);
+ output.uv = uvs[vertexIndex];
+ return output;
+}
+
+@fragment
+fn fragmentMain(@location(0) uv: vec2f) -> @location(0) vec4f {
+ return textureSample(imageTexture, imageSampler, uv);
+}
+`;
diff --git a/apps/example/src/ImageBlur/__tests__/blurPlan.test.ts b/apps/example/src/ImageBlur/__tests__/blurPlan.test.ts
new file mode 100644
index 000000000..7eb6a54b9
--- /dev/null
+++ b/apps/example/src/ImageBlur/__tests__/blurPlan.test.ts
@@ -0,0 +1,45 @@
+///
+
+import { createBlurPassPlan } from "../blurPlan";
+
+describe("createBlurPassPlan", () => {
+ it("alternates horizontal and vertical passes for every iteration", () => {
+ expect(
+ createBlurPassPlan({
+ width: 512,
+ height: 256,
+ filterSize: 15,
+ iterations: 2,
+ }),
+ ).toEqual({
+ blockDim: 114,
+ filterDim: 15,
+ passes: [
+ { bindGroup: 0, x: 5, y: 64 },
+ { bindGroup: 1, x: 3, y: 128 },
+ { bindGroup: 2, x: 5, y: 64 },
+ { bindGroup: 1, x: 3, y: 128 },
+ ],
+ });
+ });
+
+ it("rejects unsupported filter sizes and iteration counts", () => {
+ expect(() =>
+ createBlurPassPlan({
+ width: 512,
+ height: 256,
+ filterSize: 34,
+ iterations: 1,
+ }),
+ ).toThrow("filterSize must be an odd integer between 1 and 33");
+
+ expect(() =>
+ createBlurPassPlan({
+ width: 512,
+ height: 256,
+ filterSize: 15,
+ iterations: 0,
+ }),
+ ).toThrow("iterations must be a positive integer");
+ });
+});
diff --git a/apps/example/src/ImageBlur/blurPlan.ts b/apps/example/src/ImageBlur/blurPlan.ts
new file mode 100644
index 000000000..74415a27f
--- /dev/null
+++ b/apps/example/src/ImageBlur/blurPlan.ts
@@ -0,0 +1,59 @@
+const TILE_DIMENSION = 128;
+const BATCH_HEIGHT = 4;
+
+export type BlurBindGroupIndex = 0 | 1 | 2;
+
+export interface BlurPass {
+ bindGroup: BlurBindGroupIndex;
+ x: number;
+ y: number;
+}
+
+interface BlurPlanOptions {
+ width: number;
+ height: number;
+ filterSize: number;
+ iterations: number;
+}
+
+export const createBlurPassPlan = ({
+ width,
+ height,
+ filterSize,
+ iterations,
+}: BlurPlanOptions) => {
+ if (
+ !Number.isInteger(filterSize) ||
+ filterSize < 1 ||
+ filterSize > 33 ||
+ filterSize % 2 === 0
+ ) {
+ throw new Error("filterSize must be an odd integer between 1 and 33");
+ }
+ if (!Number.isInteger(iterations) || iterations < 1) {
+ throw new Error("iterations must be a positive integer");
+ }
+
+ const blockDim = TILE_DIMENSION - (filterSize - 1);
+ const horizontalPass = {
+ x: Math.ceil(width / blockDim),
+ y: Math.ceil(height / BATCH_HEIGHT),
+ };
+ const verticalPass = {
+ x: Math.ceil(height / blockDim),
+ y: Math.ceil(width / BATCH_HEIGHT),
+ };
+ const passes: BlurPass[] = [
+ { bindGroup: 0, ...horizontalPass },
+ { bindGroup: 1, ...verticalPass },
+ ];
+
+ for (let index = 1; index < iterations; index += 1) {
+ passes.push(
+ { bindGroup: 2, ...horizontalPass },
+ { bindGroup: 1, ...verticalPass },
+ );
+ }
+
+ return { blockDim, filterDim: filterSize, passes };
+};
diff --git a/apps/example/src/ImageBlur/index.ts b/apps/example/src/ImageBlur/index.ts
new file mode 100644
index 000000000..747cf3873
--- /dev/null
+++ b/apps/example/src/ImageBlur/index.ts
@@ -0,0 +1 @@
+export { ImageBlur } from "./ImageBlur";
diff --git a/apps/example/src/Route.ts b/apps/example/src/Route.ts
index a3c4cea20..8022a09ab 100644
--- a/apps/example/src/Route.ts
+++ b/apps/example/src/Route.ts
@@ -2,6 +2,7 @@ export type Routes = {
Home: undefined;
HelloTriangle: undefined;
HelloTriangleMSAA: undefined;
+ ImageBlur: undefined;
Cube: undefined;
ThreeJS: undefined;
Tensorflow: undefined;