diff --git a/README.md b/README.md
index b3ec7eb8e..8eeb1cba1 100644
--- a/README.md
+++ b/README.md
@@ -226,19 +226,21 @@ device.queue.copyExternalImageToTexture(
React Native WebGPU exposes Dawn's `SharedTextureMemory` so you can import a native pixel surface (an `IOSurface`-backed `CVPixelBuffer` on iOS, an `AHardwareBuffer` on Android) as a sampleable `GPUTexture` without copying pixels through the CPU. This is the path you want for camera frames, video frames, or anything coming out of a hardware producer.
-We expose a single umbrella feature name, `"rnwebgpu/shared-texture-memory"`. Request it at device creation.
+Like `importExternalTexture` on the web, this is **enabled by default**, there is nothing to request at device creation. The only thing to check is that the device supports it before importing. It always does on iOS/macOS; it can be missing on some Android drivers and emulators.
```tsx
-import type { VideoFrame } from "react-native-wgpu";
-
-const FEATURE = "rnwebgpu/shared-texture-memory" as GPUFeatureName;
+import type { NativeVideoFrame } from "react-native-wgpu";
const adapter = await navigator.gpu.requestAdapter();
-const requiredFeatures = adapter!.features.has(FEATURE) ? [FEATURE] : [];
-const device = await adapter!.requestDevice({ requiredFeatures });
+const device = await adapter!.requestDevice();
+
+// On by default when supported; this is the only check you need.
+if (!device.features.has("rnwebgpu/native-texture" as GPUFeatureName)) {
+ return; // rare: some Android drivers/emulators can't import native surfaces
+}
-// `frame` here is a VideoFrame whose .handle is the native surface
-// (IOSurfaceRef / AHardwareBuffer*). VideoFrames are produced by helpers
+// `frame` here is a NativeVideoFrame whose .handle is the native surface
+// (IOSurfaceRef / AHardwareBuffer*). NativeVideoFrames are produced by helpers
// like RNWebGPU.createVideoPlayer or RNWebGPU.createTestVideoFrame, or by
// any third-party module that hands you a compatible native pointer.
const memory = device.importSharedTextureMemory({
@@ -257,6 +259,53 @@ frame.release();
`beginAccess`/`endAccess` bracket the GPU's read window on the shared surface. Pass `initialized: true` when the producer has already written meaningful pixels (the typical video/camera case) and `false` when the next pass will fully overwrite the texture.
+### Importing External Textures
+
+`GPUDevice.importExternalTexture` is the higher-level path for sampling a native surface. You hand it a `NativeVideoFrame` and get back a `GPUExternalTexture` that you bind as a `texture_external` and read with `textureSampleBaseClampToEdge`. It does two things for you on top of `SharedTextureMemory`:
+
+- **Color conversion.** Camera and video surfaces are usually biplanar YUV (NV12), not RGB. An external texture carries the YUV→RGB matrix and the source/destination color-space transfer functions, so on the supported paths the sampler returns ready-to-use RGB in hardware. With raw `SharedTextureMemory` you would sample the luma/chroma planes and do that conversion by hand in WGSL. This is the main reason to prefer it for camera and video frames.
+- **Lifecycle.** It owns the `SharedTextureMemory` + `createTexture` + `beginAccess`/`endAccess` sequence internally, so you just import the frame and `destroy()` the result.
+
+It builds on the same default-on capability as Shared Texture Memory above, so feature-detect the device the same way before importing.
+
+> **Android note:** the hardware YUV→RGB conversion is fully automatic on iOS (NV12 `IOSurface`). On Android, camera frames arrive as an _opaque_ YCbCr `AHardwareBuffer`, and Dawn's Vulkan path forces an identity (`RGB_IDENTITY`) sampler conversion, so the external sample comes back as raw `[Y, Cb, Cr]`. You still get the zero-copy import and the rotation/mirror transform, but you need to apply the YUV→RGB matrix yourself in the shader. See the `CAMERA_PRELUDE` in the [VisionCamera example](/apps/example/src/VisionCamera/shaders.ts) for a ready-made BT.709 decode.
+
+```tsx
+const adapter = await navigator.gpu.requestAdapter();
+const device = await adapter!.requestDevice();
+// Feature-detect as shown above before importing on unsupported hardware.
+
+const render = () => {
+ // A GPUExternalTexture expires once the queue work that used it is submitted,
+ // so re-import one every frame.
+ const externalTexture = device.importExternalTexture({
+ source: frame, // a NativeVideoFrame
+ label: "video-frame",
+ });
+
+ const bindGroup = device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: 0, resource: externalTexture },
+ { binding: 1, resource: sampler },
+ ],
+ });
+
+ // ... encode a pass that samples `externalTexture`, then:
+ device.queue.submit([encoder.finish()]);
+
+ // Release the surface's access window right after the submit that sampled it.
+ externalTexture.destroy();
+ context.present();
+};
+```
+
+Camera frames arrive in the sensor's native orientation, so `importExternalTexture` also accepts non-spec `rotation` (`0` | `90` | `180` | `270`, in degrees) and `mirrored` (horizontal flip) options. Dawn bakes them into the sampling transform, so the shader sees an upright frame. They map directly onto VisionCamera's `frame.orientation` / `frame.isMirrored`.
+
+#### Calling `destroy()`
+
+A `GPUExternalTexture` keeps an open access window on the underlying native surface until the wrapper is destroyed. On the Web `importExternalTexture` is core and the lifetime is handled for you; here the window is tied to the JavaScript object's lifetime. Call `externalTexture.destroy()` right after the `queue.submit()` that sampled it (never before) to release the surface back to its producer immediately. `destroy()` is idempotent, and the surface is also released when the object is garbage-collected, but relying on GC can starve a producer's buffer pool (e.g. an `AVPlayer`'s recycled `IOSurface`s) and pile up GPU resources, so prefer the explicit call in a render loop.
+
### Reanimated Integration
React Native WebGPU supports running WebGPU rendering on the UI thread using [React Native Reanimated](https://docs.swmansion.com/react-native-reanimated/) and [React Native Worklets](https://github.com/margelo/react-native-worklets).
diff --git a/apps/example/src/App.tsx b/apps/example/src/App.tsx
index 354a1c572..f355f2bf7 100644
--- a/apps/example/src/App.tsx
+++ b/apps/example/src/App.tsx
@@ -37,7 +37,8 @@ import { AsyncStarvation } from "./Diagnostics/AsyncStarvation";
import { DeviceLostHang } from "./Diagnostics/DeviceLostHang";
import { StorageBufferVertices } from "./StorageBufferVertices";
import { SharedTextureMemory } from "./SharedTextureMemory";
-import { Camera } from "./Camera";
+import { ImportExternalTexture } from "./ImportExternalTexture";
+import { VisionCamera } from "./VisionCamera";
// The two lines below are needed by three.js
import "fast-text-encoding";
@@ -103,7 +104,11 @@ function App() {
name="SharedTextureMemory"
component={SharedTextureMemory}
/>
-
+
+
diff --git a/apps/example/src/Camera/Camera.tsx b/apps/example/src/Camera/Camera.tsx
deleted file mode 100644
index b0e9529a2..000000000
--- a/apps/example/src/Camera/Camera.tsx
+++ /dev/null
@@ -1,60 +0,0 @@
-import React, { useEffect } from "react";
-import { StyleSheet, Text, View } from "react-native";
-import {
- Camera as VisionCamera,
- useCameraDevice,
- useCameraPermission,
-} from "react-native-vision-camera";
-
-const styles = StyleSheet.create({
- container: {
- flex: 1,
- backgroundColor: "black",
- },
- camera: {
- flex: 1,
- },
- centered: {
- flex: 1,
- alignItems: "center",
- justifyContent: "center",
- padding: 24,
- },
- message: {
- color: "white",
- textAlign: "center",
- },
-});
-
-export const Camera = () => {
- const { hasPermission, requestPermission } = useCameraPermission();
- const device = useCameraDevice("back");
-
- useEffect(() => {
- if (!hasPermission) {
- requestPermission();
- }
- }, [hasPermission, requestPermission]);
-
- if (!hasPermission) {
- return (
-
- Requesting camera permission...
-
- );
- }
-
- if (device == null) {
- return (
-
- No camera device available.
-
- );
- }
-
- return (
-
-
-
- );
-};
diff --git a/apps/example/src/Camera/index.ts b/apps/example/src/Camera/index.ts
deleted file mode 100644
index 17290593b..000000000
--- a/apps/example/src/Camera/index.ts
+++ /dev/null
@@ -1 +0,0 @@
-export { Camera } from "./Camera";
diff --git a/apps/example/src/Home.tsx b/apps/example/src/Home.tsx
index 74c153c0f..a3a63731e 100644
--- a/apps/example/src/Home.tsx
+++ b/apps/example/src/Home.tsx
@@ -132,8 +132,12 @@ export const examples = [
title: "🎞️ Shared Texture Memory",
},
{
- screen: "Camera",
- title: "📷 Camera",
+ screen: "ImportExternalTexture",
+ title: "🎬 Import External Texture",
+ },
+ {
+ screen: "VisionCamera",
+ title: "📷 VisionCamera integration",
},
];
diff --git a/apps/example/src/ImportExternalTexture/ImportExternalTexture.tsx b/apps/example/src/ImportExternalTexture/ImportExternalTexture.tsx
new file mode 100644
index 000000000..f8399ee8a
--- /dev/null
+++ b/apps/example/src/ImportExternalTexture/ImportExternalTexture.tsx
@@ -0,0 +1,289 @@
+import React, { useEffect, useRef, useState } from "react";
+import { PixelRatio, Platform, StyleSheet, Text, View } from "react-native";
+import {
+ Canvas,
+ useCanvasRef,
+ useDevice,
+ type NativeCanvas,
+ type NativeVideoFrame,
+ type VideoPlayer,
+} from "react-native-wgpu";
+
+// This is the SharedTextureMemory demo, rewritten to use
+// GPUDevice.importExternalTexture instead of the manual
+// importSharedTextureMemory + createTexture + beginAccess/endAccess dance.
+//
+// The visible result is identical (the same video frame stretched 'cover' over
+// the canvas), but the sampling path differs: the source is bound as a
+// `texture_external` and read with textureSampleBaseClampToEdge, and Dawn owns
+// the shared-memory begin/end-access lifecycle internally. A GPUExternalTexture
+// expires once the queue work that used it is submitted, so we re-import one
+// every frame.
+const SHADER = /* wgsl */ `
+struct VsOut {
+ @builtin(position) position: vec4f,
+ @location(0) uv: vec2f,
+};
+
+struct Uniforms {
+ // Per-axis scale applied to UVs *around the center* so that the canvas
+ // samples a sub-rectangle of the texture matching the canvas aspect ratio.
+ // 'cover' fit: one axis is 1.0, the other is canvasAR / textureAR (or its
+ // reciprocal), whichever is < 1 — i.e. we crop on the longer axis.
+ uvScale: vec2f,
+};
+
+@group(0) @binding(0) var srcTex: texture_external;
+@group(0) @binding(1) var srcSampler: sampler;
+@group(0) @binding(2) var u: Uniforms;
+
+@vertex
+fn vs_main(@builtin(vertex_index) vid: u32) -> VsOut {
+ // Full-screen triangle.
+ var positions = array(
+ vec2f(-1.0, -3.0),
+ vec2f(-1.0, 1.0),
+ vec2f( 3.0, 1.0),
+ );
+ var uvs = array(
+ vec2f(0.0, 2.0),
+ vec2f(0.0, 0.0),
+ vec2f(2.0, 0.0),
+ );
+ var out: VsOut;
+ out.position = vec4f(positions[vid], 0.0, 1.0);
+ out.uv = uvs[vid];
+ return out;
+}
+
+@fragment
+fn fs_main(in: VsOut) -> @location(0) vec4f {
+ let uv = vec2f(0.5) + (in.uv - vec2f(0.5)) * u.uvScale;
+ return textureSampleBaseClampToEdge(srcTex, srcSampler, uv);
+}
+`;
+
+// importExternalTexture is backed by the "rnwebgpu/native-texture" umbrella
+// feature (it imports the frame's IOSurface / AHardwareBuffer as shared texture
+// memory, then wraps that as an external texture). Like importExternalTexture on
+// the web, that capability is now enabled by default: requestDevice / useDevice
+// turns it on automatically whenever the adapter supports it, so we don't pass
+// anything in requiredFeatures. We keep the name only to feature-detect and
+// degrade gracefully on hardware that doesn't support it.
+const FEATURE = "rnwebgpu/native-texture" as GPUFeatureName;
+
+export const ImportExternalTexture = () => {
+ const ref = useCanvasRef();
+ const [error, setError] = useState(null);
+ const rafRef = useRef(null);
+
+ const { device, adapter } = useDevice();
+
+ useEffect(() => {
+ if (!device) {
+ return;
+ }
+ // native-texture is auto-enabled when supported; if it is still missing,
+ // this hardware/driver can't back importExternalTexture.
+ if (!device.features.has(FEATURE)) {
+ setError(
+ `This device doesn't support ${FEATURE} (importExternalTexture). Adapter supports: ${
+ adapter
+ ? [...adapter.features]
+ .filter((f) => f.toString().startsWith("shared-"))
+ .join(", ") || "none"
+ : "n/a"
+ }`,
+ );
+ return;
+ }
+
+ const context = ref.current?.getContext("webgpu");
+ if (!context) {
+ return;
+ }
+ const canvas = context.canvas as unknown as NativeCanvas;
+ canvas.width = canvas.clientWidth * PixelRatio.get();
+ canvas.height = canvas.clientHeight * PixelRatio.get();
+ const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
+ context.configure({
+ device,
+ format: presentationFormat,
+ alphaMode: "premultiplied",
+ });
+
+ // Pick a frame source per platform. On iOS we stream a real video via
+ // AVPlayer; elsewhere we don't have a video pipeline yet, so we use a
+ // single synthetic IOSurface/AHardwareBuffer frame. A VideoPlayer exposes
+ // copyLatestFrame() (a fresh frame each tick) while a NativeVideoFrame does
+ // not — the render loop tells them apart with that property.
+ let source: VideoPlayer | NativeVideoFrame;
+ if (Platform.OS === "ios") {
+ const VIDEO_URL =
+ "https://test-videos.co.uk/vids/bigbuckbunny/mp4/h264/1080/Big_Buck_Bunny_1080_10s_5MB.mp4";
+ const player = RNWebGPU.createVideoPlayer(VIDEO_URL);
+ player.play();
+ source = player;
+ } else {
+ source = RNWebGPU.createTestVideoFrame(1024, 1024);
+ }
+
+ const module = device.createShaderModule({ code: SHADER });
+ const pipeline = device.createRenderPipeline({
+ layout: "auto",
+ vertex: { module, entryPoint: "vs_main" },
+ fragment: {
+ module,
+ entryPoint: "fs_main",
+ targets: [{ format: presentationFormat }],
+ },
+ primitive: { topology: "triangle-list" },
+ });
+ const sampler = device.createSampler({
+ magFilter: "linear",
+ minFilter: "linear",
+ });
+ // One persistent uniform buffer; rewritten whenever the frame dimensions
+ // change.
+ const uniformBuffer = device.createBuffer({
+ size: 16, // vec2 padded to 16-byte uniform alignment
+ usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
+ });
+
+ // 'cover' fit: scale UVs around their center so the longer axis of the
+ // texture is cropped to match the canvas aspect ratio.
+ const computeUvScale = (texW: number, texH: number): [number, number] => {
+ const canvasAR = canvas.width / canvas.height;
+ const texAR = texW / texH;
+ if (texAR > canvasAR) {
+ // Texture is wider than the canvas: crop horizontally.
+ return [canvasAR / texAR, 1];
+ } else {
+ // Texture is taller than (or equal to) the canvas: crop vertically.
+ return [1, texAR / canvasAR];
+ }
+ };
+
+ // Hold the current frame across rAF ticks. For a VideoPlayer we pull the
+ // latest frame each tick (keeping the previous one when the decoder hasn't
+ // produced a new one yet, to avoid a black flash); for a one-shot
+ // NativeVideoFrame we just keep re-importing the same frame.
+ let currentFrame: NativeVideoFrame | null =
+ "copyLatestFrame" in source ? null : source;
+ let lastDims: [number, number] | null = null;
+
+ const render = () => {
+ if ("copyLatestFrame" in source) {
+ const newFrame = source.copyLatestFrame();
+ if (newFrame) {
+ if (currentFrame) {
+ currentFrame.release();
+ }
+ currentFrame = newFrame;
+ }
+ }
+
+ const encoder = device.createCommandEncoder();
+ const pass = encoder.beginRenderPass({
+ colorAttachments: [
+ {
+ view: context.getCurrentTexture().createView(),
+ clearValue: { r: 0, g: 0, b: 0, a: 1 },
+ loadOp: "clear",
+ storeOp: "store",
+ },
+ ],
+ });
+
+ // A GPUExternalTexture expires after each submit, so re-import one every
+ // tick, even when sampling the same frame as last tick. Unlike on the
+ // web, the shared-memory begin/end-access window is tied to this
+ // wrapper's lifetime, so we destroy() it right after submit (below) to
+ // release the surface promptly instead of waiting for GC.
+ let externalTex: GPUExternalTexture | null = null;
+ if (currentFrame) {
+ try {
+ externalTex = device.importExternalTexture({
+ source: currentFrame,
+ label: "video-frame",
+ });
+ } catch (e) {
+ console.warn("[ImportExternalTexture] import failed:", e);
+ }
+
+ if (externalTex) {
+ if (
+ !lastDims ||
+ lastDims[0] !== currentFrame.width ||
+ lastDims[1] !== currentFrame.height
+ ) {
+ const [sx, sy] = computeUvScale(
+ currentFrame.width,
+ currentFrame.height,
+ );
+ device.queue.writeBuffer(
+ uniformBuffer,
+ 0,
+ new Float32Array([sx, sy]),
+ );
+ lastDims = [currentFrame.width, currentFrame.height];
+ }
+ const bindGroup = device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: 0, resource: externalTex },
+ { binding: 1, resource: sampler },
+ { binding: 2, resource: { buffer: uniformBuffer } },
+ ],
+ });
+ pass.setPipeline(pipeline);
+ pass.setBindGroup(0, bindGroup);
+ pass.draw(3);
+ }
+ }
+
+ pass.end();
+ device.queue.submit([encoder.finish()]);
+ // Now that the work sampling it has been submitted, end the external
+ // texture's access window so the frame's surface is released promptly.
+ externalTex?.destroy();
+ context.present();
+ rafRef.current = requestAnimationFrame(render);
+ };
+ rafRef.current = requestAnimationFrame(render);
+
+ return () => {
+ if (rafRef.current !== null) {
+ cancelAnimationFrame(rafRef.current);
+ }
+ if (currentFrame) {
+ currentFrame.release();
+ currentFrame = null;
+ }
+ uniformBuffer.destroy();
+ // For the player, release it; the one-shot frame was released above as
+ // currentFrame (same object), so don't double-release it here.
+ if ("copyLatestFrame" in source) {
+ source.release();
+ }
+ };
+ }, [device, adapter, ref]);
+
+ if (error) {
+ return (
+
+ {error}
+
+ );
+ }
+ return (
+
+
+
+ );
+};
+
+const styles = StyleSheet.create({
+ errorContainer: { flex: 1, padding: 16, justifyContent: "center" },
+ errorText: { color: "red", fontSize: 14 },
+});
diff --git a/apps/example/src/ImportExternalTexture/index.ts b/apps/example/src/ImportExternalTexture/index.ts
new file mode 100644
index 000000000..d08addd8a
--- /dev/null
+++ b/apps/example/src/ImportExternalTexture/index.ts
@@ -0,0 +1 @@
+export * from "./ImportExternalTexture";
diff --git a/apps/example/src/Route.ts b/apps/example/src/Route.ts
index 5c43b2a36..5b381edee 100644
--- a/apps/example/src/Route.ts
+++ b/apps/example/src/Route.ts
@@ -30,5 +30,6 @@ export type Routes = {
DeviceLostHang: undefined;
StorageBufferVertices: undefined;
SharedTextureMemory: undefined;
- Camera: undefined;
+ ImportExternalTexture: undefined;
+ VisionCamera: undefined;
};
diff --git a/apps/example/src/SharedTextureMemory/SharedTextureMemory.tsx b/apps/example/src/SharedTextureMemory/SharedTextureMemory.tsx
index 7805cb7ff..b5627cc43 100644
--- a/apps/example/src/SharedTextureMemory/SharedTextureMemory.tsx
+++ b/apps/example/src/SharedTextureMemory/SharedTextureMemory.tsx
@@ -6,7 +6,7 @@ import {
useDevice,
type GPUSharedTextureMemory,
type NativeCanvas,
- type VideoFrame,
+ type NativeVideoFrame,
} from "react-native-wgpu";
const SHADER = /* wgsl */ `
@@ -53,7 +53,12 @@ fn fs_main(in: VsOut) -> @location(0) vec4f {
}
`;
-const REQUIRED_FEATURES = ["rnwebgpu/shared-texture-memory" as GPUFeatureName];
+// This screen keeps the explicit feature request for documentation. As of the
+// default-on change, "rnwebgpu/native-texture" is enabled automatically by
+// requestDevice / useDevice whenever the adapter supports it (like
+// importExternalTexture on the web), so passing it in requiredFeatures is
+// optional. We still list it here to show how to gate on the capability.
+const REQUIRED_FEATURES = ["rnwebgpu/native-texture" as GPUFeatureName];
export const SharedTextureMemory = () => {
const ref = useCanvasRef();
@@ -103,7 +108,7 @@ export const SharedTextureMemory = () => {
// from copyLatestFrame() as "keep showing the previous frame", which means
// a one-shot source renders correctly without any other change.
interface FrameSource {
- copyLatestFrame(): VideoFrame | null;
+ copyLatestFrame(): NativeVideoFrame | null;
release(): void;
}
let source: FrameSource;
@@ -117,7 +122,7 @@ export const SharedTextureMemory = () => {
release: () => player.release(),
};
} else {
- let pending: VideoFrame | null = RNWebGPU.createTestVideoFrame(
+ let pending: NativeVideoFrame | null = RNWebGPU.createTestVideoFrame(
1024,
1024,
);
@@ -168,7 +173,7 @@ export const SharedTextureMemory = () => {
// demo we rely on AVPlayer recycling its IOSurface pool, which is safe as
// long as we end-access before letting the player reclaim the buffer.
type Bound = {
- frame: VideoFrame;
+ frame: NativeVideoFrame;
memory: GPUSharedTextureMemory;
texture: GPUTexture;
bindGroup: GPUBindGroup;
@@ -190,7 +195,7 @@ export const SharedTextureMemory = () => {
}
};
- const bindFrame = (frame: VideoFrame): Bound | null => {
+ const bindFrame = (frame: NativeVideoFrame): Bound | null => {
try {
const memory = device.importSharedTextureMemory({
handle: frame.handle,
diff --git a/apps/example/src/Tests.tsx b/apps/example/src/Tests.tsx
index c6ae7ae1b..fc9ebf635 100644
--- a/apps/example/src/Tests.tsx
+++ b/apps/example/src/Tests.tsx
@@ -14,12 +14,6 @@ export const CI = process.env.CI === "true";
const { width } = Dimensions.get("window");
const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
-// Umbrella feature owned by react-native-wgpu: expands to the platform's
-// Dawn feature pair natively and appears on adapter.features when supported.
-const OPTIONAL_SHARED_TEXTURE_MEMORY_FEATURES = [
- "rnwebgpu/shared-texture-memory" as GPUFeatureName,
-];
-
export const Tests = ({ assets: { di3D, saturn, moon } }: AssetProps) => {
const [texture, setTexture] = useState(null);
const [adapter, setAdapter] = useState(null);
@@ -33,10 +27,11 @@ export const Tests = ({ assets: { di3D, saturn, moon } }: AssetProps) => {
if (!a) {
throw new Error("No appropriate GPUAdapter found.");
}
- const requiredFeatures = OPTIONAL_SHARED_TEXTURE_MEMORY_FEATURES.filter(
- (f) => a.features.has(f),
- );
- const d = await a.requestDevice({ requiredFeatures });
+ // "rnwebgpu/native-texture" is enabled by default whenever the adapter
+ // supports it, so the shared-texture / importExternalTexture specs get
+ // the capability without requesting it here. Specs that need it still
+ // gate on device.features.has(...) and skip where it is unavailable.
+ const d = await a.requestDevice();
if (!d) {
throw new Error("No appropriate GPUDevice found.");
}
diff --git a/apps/example/src/VisionCamera/EffectToolbar.tsx b/apps/example/src/VisionCamera/EffectToolbar.tsx
new file mode 100644
index 000000000..15ca54665
--- /dev/null
+++ b/apps/example/src/VisionCamera/EffectToolbar.tsx
@@ -0,0 +1,73 @@
+import React from "react";
+import { Pressable, ScrollView, StyleSheet, Text, View } from "react-native";
+
+import { FEATURES, type Modes } from "./features";
+
+type Props = {
+ modes: Modes;
+ onCycle: (key: keyof Modes, optionsCount: number) => void;
+};
+
+export const EffectToolbar = ({ modes, onCycle }: Props) => {
+ return (
+
+
+ {FEATURES.map((f) => (
+ onCycle(f.key, f.labels.length)}
+ style={({ pressed }) => [
+ styles.button,
+ pressed && styles.buttonPressed,
+ ]}
+ >
+ {f.title}
+ {f.labels[modes[f.key]]}
+
+ ))}
+
+
+ );
+};
+
+const styles = StyleSheet.create({
+ toolbar: {
+ position: "absolute",
+ left: 0,
+ right: 0,
+ top: 60,
+ },
+ toolbarContent: {
+ paddingHorizontal: 12,
+ gap: 8,
+ },
+ button: {
+ backgroundColor: "rgba(0,0,0,0.55)",
+ borderColor: "rgba(255,255,255,0.18)",
+ borderWidth: 1,
+ borderRadius: 14,
+ paddingHorizontal: 12,
+ paddingVertical: 8,
+ minWidth: 84,
+ },
+ buttonPressed: {
+ backgroundColor: "rgba(255,255,255,0.18)",
+ },
+ buttonTitle: {
+ color: "rgba(255,255,255,0.65)",
+ fontSize: 11,
+ fontWeight: "500",
+ letterSpacing: 0.4,
+ textTransform: "uppercase",
+ },
+ buttonValue: {
+ color: "white",
+ fontSize: 15,
+ fontWeight: "600",
+ marginTop: 2,
+ },
+});
diff --git a/apps/example/src/VisionCamera/VisionCamera.tsx b/apps/example/src/VisionCamera/VisionCamera.tsx
new file mode 100644
index 000000000..c4adcfaa0
--- /dev/null
+++ b/apps/example/src/VisionCamera/VisionCamera.tsx
@@ -0,0 +1,700 @@
+import React, { useCallback, useEffect, useMemo, useState } from "react";
+import {
+ Linking,
+ PixelRatio,
+ Platform,
+ StyleSheet,
+ Text,
+ TouchableOpacity,
+ View,
+} from "react-native";
+import {
+ Canvas,
+ useCanvasRef,
+ type NativeCanvas,
+ type RNCanvasContext,
+} from "react-native-wgpu";
+import {
+ useCamera,
+ useCameraDevices,
+ useCameraPermission,
+ useFrameOutput,
+} from "react-native-vision-camera";
+
+import { BLUR_SHADER, PREPASS_SHADER } from "./blurShaders";
+import { EffectToolbar } from "./EffectToolbar";
+import {
+ ABERRATION_STRENGTHS,
+ BLUR_ITERATIONS,
+ INITIAL_MODES,
+ PIXELATE_BLOCKS,
+ type Modes,
+} from "./features";
+import { CAMERA_PRELUDE, SHADER } from "./shaders";
+
+// Camera frame → SharedTextureMemory (NV12 biplanar) → GPUExternalTexture →
+// textureSampleBaseClampToEdge with hardware YUV/sRGB conversion → chromatic
+// aberration in WGSL.
+//
+// Everything past frame arrival runs on Vision Camera's worklet runtime.
+// react-native-wgpu's `registerWebGPUForReanimated` (loaded from main on
+// startup) registers a Worklets custom serializer for GPUDevice / canvas /
+// pipeline / sampler / buffer, so the closure references below auto-box on
+// the way into the worklet and auto-unbox on the way out.
+//
+// The WGSL (SHADER + CAMERA_PRELUDE) lives in ./shaders.
+
+const REQUIRED_FEATURES: GPUFeatureName[] = [
+ "rnwebgpu/native-texture" as GPUFeatureName,
+ "dawn-multi-planar-formats" as GPUFeatureName,
+];
+
+// Android-only feature, gates Dawn's "wrap a YCbCr AHB as a GPUExternalTexture
+// with implicit SamplerYcbcrConversion" path. Without it our native
+// `importExternalTexture` flow on Android can't produce a usable external
+// texture from a camera frame. We probe the adapter for it and surface a
+// clear error if the device's Vulkan driver doesn't advertise it (e.g. some
+// Android-Desktop / Chromebook configurations).
+const OPAQUE_YCBCR_EXT =
+ "opaque-ycbcr-android-for-external-texture" as GPUFeatureName;
+
+// Blur infrastructure: prepass writes the
+// cover-fit camera image into a 1/4-res rgba8unorm, the separable box-blur
+// compute pings between two storage textures, and the main pass linearly
+// upsamples the final result so the effective sigma is ~4x the per-iteration
+// kernel.
+const BLUR_SCALE = 4;
+const BLUR_FILTER_SIZE = 31;
+const BLUR_TILE_DIM = 128;
+const BLUR_BATCH = 4;
+const BLUR_BLOCK_DIM = BLUR_TILE_DIM - BLUR_FILTER_SIZE;
+
+export const VisionCamera = () => {
+ const { hasPermission, requestPermission } = useCameraPermission();
+ useEffect(() => {
+ if (!hasPermission) {
+ requestPermission();
+ }
+ }, [hasPermission, requestPermission]);
+
+ if (!hasPermission) {
+ return (
+
+
+ Camera access is required. Grant it in Settings or tap below.
+
+ Linking.openSettings()}
+ style={styles.permissionButton}
+ >
+ Open Settings
+
+
+ );
+ }
+ return ;
+};
+
+const CameraView = () => {
+ const ref = useCanvasRef();
+ const [gpu, setGpu] = useState<{
+ adapter: GPUAdapter;
+ device: GPUDevice;
+ } | null>(null);
+ const [deviceError, setDeviceError] = useState(null);
+ useEffect(() => {
+ let cancelled = false;
+ (async () => {
+ try {
+ const adapter = await navigator.gpu.requestAdapter();
+ if (!adapter) {
+ throw new Error("requestAdapter returned null");
+ }
+ const adapterFeatures = [...adapter.features].sort();
+ console.log(
+ "[VisionCamera] adapter features (" +
+ adapterFeatures.length +
+ "): " +
+ adapterFeatures.join(", "),
+ );
+ const hasOpaqueYCbCrExt =
+ Platform.OS !== "android" || adapter.features.has(OPAQUE_YCBCR_EXT);
+ if (Platform.OS === "android" && !hasOpaqueYCbCrExt) {
+ throw new Error(
+ "This Android device's Vulkan driver doesn't advertise " +
+ "opaque-ycbcr-android-for-external-texture. Camera-frame import " +
+ "as a GPUExternalTexture isn't supported here. (This is a " +
+ "device/driver limitation, not a code issue.)",
+ );
+ }
+ const featuresToRequest: GPUFeatureName[] = [
+ ...REQUIRED_FEATURES,
+ ...(Platform.OS === "android" ? [OPAQUE_YCBCR_EXT] : []),
+ ];
+ console.log(
+ "[VisionCamera] requesting device with features: " +
+ featuresToRequest.join(", "),
+ );
+ const device = await adapter.requestDevice({
+ requiredFeatures: featuresToRequest,
+ });
+ if (cancelled) {
+ return;
+ }
+ console.log(
+ "[VisionCamera] device created, features: " +
+ [...device.features].sort().join(", "),
+ );
+ setGpu({ adapter, device });
+ } catch (e) {
+ if (cancelled) {
+ return;
+ }
+ console.warn("[VisionCamera] device creation failed: " + String(e));
+ setDeviceError(String(e));
+ }
+ })();
+ return () => {
+ cancelled = true;
+ };
+ }, []);
+ const device = gpu?.device ?? null;
+ const adapter = gpu?.adapter ?? null;
+ // Capture the RNWebGPU singleton into a local so the frame-processor worklet
+ // closes over it. RNWebGPU is a registered, boxable WebGPU NativeObject, so
+ // the Worklets custom serializer ships it across the worklet boundary the same
+ // way it does `device` (it is NOT installed as a global on worklet runtimes).
+ // This is what lets us call the interop factory off RNWebGPU, where the native
+ // platform context already lives, instead of off `device`.
+ const rnwgpu = RNWebGPU;
+ const devices = useCameraDevices();
+ // Pick back camera if available, otherwise front, otherwise anything. The
+ // iOS simulator returns an empty list since there are no cameras, in which
+ // case we surface a clear error rather than letting useCamera throw.
+ const cameraDevice = useMemo(
+ () =>
+ devices.find((d) => d.position === "back") ??
+ devices.find((d) => d.position === "front") ??
+ devices[0],
+ [devices],
+ );
+
+ const [pipelineState, setPipelineState] = useState<{
+ pipeline: GPURenderPipeline;
+ sampler: GPUSampler;
+ uniformBuffer: GPUBuffer;
+ context: RNCanvasContext;
+ canvasWidth: number;
+ canvasHeight: number;
+ prepassPipeline: GPURenderPipeline;
+ prepassUniformBuffer: GPUBuffer;
+ blurPipeline: GPUComputePipeline;
+ blurConstants: GPUBindGroup;
+ blurBindGroup0: GPUBindGroup;
+ blurBindGroup1: GPUBindGroup;
+ blurBindGroup2: GPUBindGroup;
+ blurSrcTexture: GPUTexture;
+ blurredView: GPUTextureView;
+ blurWidth: number;
+ blurHeight: number;
+ } | null>(null);
+ const [error, setError] = useState(null);
+ const [modes, setModes] = useState(INITIAL_MODES);
+ const cycle = useCallback((key: keyof Modes, optionsCount: number) => {
+ setModes((prev) => ({ ...prev, [key]: (prev[key] + 1) % optionsCount }));
+ }, []);
+
+ // Initialize pipeline once device + canvas are both ready.
+ useEffect(() => {
+ if (!device || pipelineState) {
+ return;
+ }
+ const missing = REQUIRED_FEATURES.filter((f) => !device.features.has(f));
+ if (missing.length > 0) {
+ setError(
+ `Device missing features [${missing.join(", ")}]. Adapter: ${
+ adapter
+ ? [...adapter.features]
+ .filter((f) => f.toString().startsWith("shared-"))
+ .join(", ") || "none"
+ : "n/a"
+ }`,
+ );
+ return;
+ }
+ const context = ref.current?.getContext("webgpu");
+ if (!context) {
+ return;
+ }
+ const canvas = context.canvas as unknown as NativeCanvas;
+ canvas.width = canvas.clientWidth * PixelRatio.get();
+ canvas.height = canvas.clientHeight * PixelRatio.get();
+ const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
+ context.configure({
+ device,
+ format: presentationFormat,
+ alphaMode: "premultiplied",
+ });
+
+ const module = device.createShaderModule({
+ code: CAMERA_PRELUDE + SHADER,
+ });
+ const pipeline = device.createRenderPipeline({
+ layout: "auto",
+ vertex: { module, entryPoint: "vs_main" },
+ fragment: {
+ module,
+ entryPoint: "fs_main",
+ targets: [{ format: presentationFormat }],
+ },
+ primitive: { topology: "triangle-list" },
+ });
+ const sampler = device.createSampler({
+ magFilter: "linear",
+ minFilter: "linear",
+ });
+ const uniformBuffer = device.createBuffer({
+ size: 32,
+ usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
+ });
+
+ // ----- Blur infrastructure -----
+ const blurWidth = Math.max(
+ BLUR_TILE_DIM,
+ Math.ceil(canvas.width / BLUR_SCALE),
+ );
+ const blurHeight = Math.max(
+ BLUR_TILE_DIM,
+ Math.ceil(canvas.height / BLUR_SCALE),
+ );
+
+ const prepassModule = device.createShaderModule({
+ code: CAMERA_PRELUDE + PREPASS_SHADER,
+ });
+ const prepassPipeline = device.createRenderPipeline({
+ layout: "auto",
+ vertex: { module: prepassModule, entryPoint: "vs_main" },
+ fragment: {
+ module: prepassModule,
+ entryPoint: "fs_main",
+ targets: [{ format: "rgba8unorm" }],
+ },
+ primitive: { topology: "triangle-list" },
+ });
+ const prepassUniformBuffer = device.createBuffer({
+ size: 16,
+ usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
+ });
+
+ const blurSrcTexture = device.createTexture({
+ size: [blurWidth, blurHeight],
+ format: "rgba8unorm",
+ usage:
+ GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.TEXTURE_BINDING,
+ });
+ const blurPing = [0, 1].map(() =>
+ device.createTexture({
+ size: [blurWidth, blurHeight],
+ format: "rgba8unorm",
+ usage:
+ GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.TEXTURE_BINDING,
+ }),
+ );
+
+ const blurPipeline = device.createComputePipeline({
+ layout: "auto",
+ compute: {
+ module: device.createShaderModule({ code: BLUR_SHADER }),
+ entryPoint: "main",
+ },
+ });
+
+ const flip0Buffer = device.createBuffer({
+ size: 4,
+ mappedAtCreation: true,
+ usage: GPUBufferUsage.UNIFORM,
+ });
+ new Uint32Array(flip0Buffer.getMappedRange())[0] = 0;
+ flip0Buffer.unmap();
+ const flip1Buffer = device.createBuffer({
+ size: 4,
+ mappedAtCreation: true,
+ usage: GPUBufferUsage.UNIFORM,
+ });
+ new Uint32Array(flip1Buffer.getMappedRange())[0] = 1;
+ flip1Buffer.unmap();
+
+ const blurParamsBuffer = device.createBuffer({
+ size: 8,
+ usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
+ });
+ device.queue.writeBuffer(
+ blurParamsBuffer,
+ 0,
+ new Uint32Array([BLUR_FILTER_SIZE + 1, BLUR_BLOCK_DIM]),
+ );
+
+ const blurConstants = device.createBindGroup({
+ layout: blurPipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: 0, resource: sampler },
+ { binding: 1, resource: { buffer: blurParamsBuffer } },
+ ],
+ });
+ // H: blurSrcTexture -> blurPing[0]
+ const blurBindGroup0 = device.createBindGroup({
+ layout: blurPipeline.getBindGroupLayout(1),
+ entries: [
+ { binding: 1, resource: blurSrcTexture.createView() },
+ { binding: 2, resource: blurPing[0].createView() },
+ { binding: 3, resource: { buffer: flip0Buffer } },
+ ],
+ });
+ // V: blurPing[0] -> blurPing[1]
+ const blurBindGroup1 = device.createBindGroup({
+ layout: blurPipeline.getBindGroupLayout(1),
+ entries: [
+ { binding: 1, resource: blurPing[0].createView() },
+ { binding: 2, resource: blurPing[1].createView() },
+ { binding: 3, resource: { buffer: flip1Buffer } },
+ ],
+ });
+ // H (iteration N>=2): blurPing[1] -> blurPing[0]
+ const blurBindGroup2 = device.createBindGroup({
+ layout: blurPipeline.getBindGroupLayout(1),
+ entries: [
+ { binding: 1, resource: blurPing[1].createView() },
+ { binding: 2, resource: blurPing[0].createView() },
+ { binding: 3, resource: { buffer: flip0Buffer } },
+ ],
+ });
+ // Final iteration's V pass always lands in blurPing[1].
+ const blurredView = blurPing[1].createView();
+
+ setPipelineState({
+ pipeline,
+ sampler,
+ uniformBuffer,
+ context,
+ canvasWidth: canvas.width,
+ canvasHeight: canvas.height,
+ prepassPipeline,
+ prepassUniformBuffer,
+ blurPipeline,
+ blurConstants,
+ blurBindGroup0,
+ blurBindGroup1,
+ blurBindGroup2,
+ blurSrcTexture,
+ blurredView,
+ blurWidth,
+ blurHeight,
+ });
+ }, [device, adapter, ref, pipelineState]);
+
+ // Build the frame processor worklet. Captured WebGPU objects flow into the
+ // worklet runtime via the registerWebGPUForReanimated custom serializer.
+ // Diagnostic: log once on the *first call only* by capturing a plain bool
+ // box. Worklets serializes the box once on closure creation, so flipping
+ // .seen mutates the worklet-side copy; we use it strictly to avoid spamming
+ // metro logs and don't rely on main-thread visibility.
+ const logBox = useMemo(() => ({ seen: false }), []);
+ const frameOutput = useFrameOutput({
+ pixelFormat: "native", // zero-copy, gives us NV12 IOSurfaces on iOS
+ onFrame: (frame) => {
+ "worklet";
+ if (!logBox.seen) {
+ logBox.seen = true;
+
+ console.log(
+ "[VisionCamera] worklet first frame, hasPipeline=" +
+ String(pipelineState != null) +
+ " hasDevice=" +
+ String(device != null) +
+ " frame=" +
+ String(frame.width) +
+ "x" +
+ String(frame.height) +
+ " orientation=" +
+ String(frame.orientation) +
+ " isMirrored=" +
+ String(frame.isMirrored),
+ );
+ }
+ if (!pipelineState || !device) {
+ frame.dispose();
+ return;
+ }
+ const {
+ pipeline,
+ sampler,
+ uniformBuffer,
+ context,
+ canvasWidth,
+ canvasHeight,
+ prepassPipeline,
+ prepassUniformBuffer,
+ blurPipeline,
+ blurConstants,
+ blurBindGroup0,
+ blurBindGroup1,
+ blurBindGroup2,
+ blurSrcTexture,
+ blurredView,
+ blurWidth,
+ blurHeight,
+ } = pipelineState;
+ const blurIterations = BLUR_ITERATIONS[modes.blur] ?? 0;
+ const blurMode = blurIterations > 0 ? modes.blur : 0;
+ const nativeBuffer = frame.getNativeBuffer();
+ try {
+ let videoFrame;
+ try {
+ // Call createVideoFrameFromNativeBuffer on the captured RNWebGPU
+ // singleton (see `rnwgpu` above). It rides into this worklet via the
+ // WebGPU custom serializer, same as `device`, so the factory can live
+ // on RNWebGPU where the native platform context already is — no
+ // GPUDevice-level workaround and no PlatformContext global singleton.
+ videoFrame = rnwgpu.createVideoFrameFromNativeBuffer(
+ nativeBuffer.pointer,
+ );
+ } catch (e) {
+ console.warn(
+ "[VisionCamera] createVideoFrameFromNativeBuffer threw: " +
+ String(e),
+ );
+ throw e;
+ }
+ try {
+ // Orientation. The sensor buffer is landscape; frame.orientation is
+ // the rotation needed to bring it upright. We hand that to Dawn via
+ // importExternalTexture's `rotation`, which de-rotates the frame into
+ // the portrait canvas. The residual vertical flip (Android buffer
+ // Y-origin) and YUV->RGB are corrected in the shader (CAMERA_PRELUDE).
+ let rotationDeg: 0 | 90 | 180 | 270 = 0;
+ if (frame.orientation === "right") {
+ rotationDeg = 90;
+ } else if (frame.orientation === "down") {
+ rotationDeg = 180;
+ } else if (frame.orientation === "left") {
+ rotationDeg = 270;
+ }
+ // A 90/270 rotation swaps the displayed width & height, so cover-fit
+ // uses the post-rotation dimensions.
+ const rotated = rotationDeg === 90 || rotationDeg === 270;
+ const dispW = rotated ? videoFrame.height : videoFrame.width;
+ const dispH = rotated ? videoFrame.width : videoFrame.height;
+ // Compute cover-fit uvScale based on frame & canvas aspect ratios.
+ // On most phones the back camera is landscape (e.g. 1920x1080) and
+ // the canvas is portrait, so the y-axis gets cropped.
+ const canvasAR = canvasWidth / canvasHeight;
+ const frameAR = dispW / dispH;
+ let sx = 1;
+ let sy = 1;
+ if (frameAR > canvasAR) {
+ sx = canvasAR / frameAR;
+ } else {
+ sy = frameAR / canvasAR;
+ }
+ // 32-byte uniform: vec4f params + vec4u modes. Built on a single
+ // ArrayBuffer so the f32/u32 halves go up in one writeBuffer call.
+ const uniformData = new ArrayBuffer(32);
+ const uniformF32 = new Float32Array(uniformData);
+ const uniformU32 = new Uint32Array(uniformData);
+ uniformF32[0] = sx;
+ uniformF32[1] = sy;
+ uniformF32[2] = ABERRATION_STRENGTHS[modes.aberration] ?? 0;
+ uniformF32[3] = PIXELATE_BLOCKS[modes.pixelate] ?? 0;
+ uniformU32[4] = modes.effect;
+ uniformU32[5] = modes.tint;
+ uniformU32[6] = modes.vignette;
+ uniformU32[7] = blurMode;
+ device.queue.writeBuffer(uniformBuffer, 0, uniformData);
+
+ let externalTex;
+ try {
+ externalTex = device.importExternalTexture({
+ source: videoFrame,
+ label: "camera-frame",
+ rotation: rotationDeg,
+ mirrored: frame.isMirrored,
+ });
+ } catch (e) {
+ console.warn(
+ "[VisionCamera] importExternalTexture threw: " + String(e),
+ );
+ throw e;
+ }
+ const bindGroup = device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: 0, resource: externalTex },
+ { binding: 1, resource: sampler },
+ { binding: 2, resource: { buffer: uniformBuffer } },
+ { binding: 3, resource: blurredView },
+ ],
+ });
+
+ const encoder = device.createCommandEncoder();
+
+ if (blurIterations > 0) {
+ // Prepass: cover-fit external (YUV) -> rgba8unorm at 1/4 res.
+ device.queue.writeBuffer(
+ prepassUniformBuffer,
+ 0,
+ new Float32Array([dispW, dispH, canvasWidth, canvasHeight]),
+ );
+ const prepassBindGroup = device.createBindGroup({
+ layout: prepassPipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: 0, resource: externalTex },
+ { binding: 1, resource: sampler },
+ { binding: 2, resource: { buffer: prepassUniformBuffer } },
+ ],
+ });
+ const prepass = encoder.beginRenderPass({
+ colorAttachments: [
+ {
+ view: blurSrcTexture.createView(),
+ clearValue: { r: 0, g: 0, b: 0, a: 1 },
+ loadOp: "clear",
+ storeOp: "store",
+ },
+ ],
+ });
+ prepass.setPipeline(prepassPipeline);
+ prepass.setBindGroup(0, prepassBindGroup);
+ prepass.draw(3);
+ prepass.end();
+
+ // Separable box-blur compute. Iteration 1 reads from
+ // blurSrcTexture; subsequent iterations ping-pong between
+ // blurPing[1] -> blurPing[0] -> blurPing[1].
+ const compute = encoder.beginComputePass();
+ compute.setPipeline(blurPipeline);
+ compute.setBindGroup(0, blurConstants);
+ compute.setBindGroup(1, blurBindGroup0);
+ compute.dispatchWorkgroups(
+ Math.ceil(blurWidth / BLUR_BLOCK_DIM),
+ Math.ceil(blurHeight / BLUR_BATCH),
+ );
+ compute.setBindGroup(1, blurBindGroup1);
+ compute.dispatchWorkgroups(
+ Math.ceil(blurHeight / BLUR_BLOCK_DIM),
+ Math.ceil(blurWidth / BLUR_BATCH),
+ );
+ for (let i = 0; i < blurIterations - 1; i++) {
+ compute.setBindGroup(1, blurBindGroup2);
+ compute.dispatchWorkgroups(
+ Math.ceil(blurWidth / BLUR_BLOCK_DIM),
+ Math.ceil(blurHeight / BLUR_BATCH),
+ );
+ compute.setBindGroup(1, blurBindGroup1);
+ compute.dispatchWorkgroups(
+ Math.ceil(blurHeight / BLUR_BLOCK_DIM),
+ Math.ceil(blurWidth / BLUR_BATCH),
+ );
+ }
+ compute.end();
+ }
+
+ const pass = encoder.beginRenderPass({
+ colorAttachments: [
+ {
+ view: context.getCurrentTexture().createView(),
+ clearValue: { r: 0, g: 0, b: 0, a: 1 },
+ loadOp: "clear",
+ storeOp: "store",
+ },
+ ],
+ });
+ pass.setPipeline(pipeline);
+ pass.setBindGroup(0, bindGroup);
+ pass.draw(3);
+ pass.end();
+ device.queue.submit([encoder.finish()]);
+ // The work sampling it is submitted, so end the external texture's
+ // access window now to release the camera frame's surface promptly
+ // (don't wait for GC, which would starve the frame buffer pool).
+ externalTex.destroy();
+ context.present();
+ } finally {
+ videoFrame.release();
+ }
+ } finally {
+ nativeBuffer.release();
+ frame.dispose();
+ }
+ },
+ });
+
+ // We have to call useCamera unconditionally (hook order). Pass a stub
+ // device when none exists so the hook doesn't throw, but keep isActive
+ // false so it never tries to start the session.
+ useCamera({
+ isActive: pipelineState != null && cameraDevice != null,
+ device: cameraDevice as NonNullable,
+ outputs: [frameOutput],
+ });
+
+ if (deviceError) {
+ return (
+
+
+ Device creation failed: {deviceError}
+
+
+ );
+ }
+ if (error) {
+ return (
+
+ {error}
+
+ );
+ }
+ if (!device) {
+ return (
+
+ Waiting for GPU device...
+
+ );
+ }
+ if (cameraDevice == null) {
+ return (
+
+
+ No camera available. This screen needs a physical device with a camera
+ (the iOS Simulator does not have one).
+
+
+ );
+ }
+ return (
+
+
+
+
+ );
+};
+
+const styles = StyleSheet.create({
+ root: { flex: 1, backgroundColor: "black" },
+ canvas: { flex: 1 },
+ errorContainer: { flex: 1, padding: 16, justifyContent: "center" },
+ errorText: { color: "red", fontSize: 14 },
+ permissionContainer: {
+ flex: 1,
+ padding: 24,
+ justifyContent: "center",
+ alignItems: "center",
+ },
+ permissionText: { fontSize: 16, textAlign: "center", marginBottom: 16 },
+ permissionButton: {
+ backgroundColor: "#007AFF",
+ paddingHorizontal: 24,
+ paddingVertical: 12,
+ borderRadius: 8,
+ },
+ permissionButtonText: { color: "white", fontSize: 16, fontWeight: "600" },
+});
diff --git a/apps/example/src/VisionCamera/blurShaders.ts b/apps/example/src/VisionCamera/blurShaders.ts
new file mode 100644
index 000000000..eddbfb8d7
--- /dev/null
+++ b/apps/example/src/VisionCamera/blurShaders.ts
@@ -0,0 +1,144 @@
+// Two shaders that together feed the Blur effect with a true large-kernel
+// gaussian, computed once on the GPU per frame and resampled by the main pass.
+//
+// 1. PREPASS_SHADER. Renders the external (YUV) camera texture into a regular
+// rgba8unorm texture, cover-projected at 1/4 canvas resolution. This solves
+// two problems at once:
+// a. texture_external can only be sampled with textureSampleBaseClampToEdge
+// and cannot be bound to a compute pipeline that expects texture_2d,
+// so we need a fixup pass anyway.
+// b. The blur runs at 1/4 res, and the main fragment's linear sampler
+// upsamples for free, which makes the effective kernel ~4x wider than
+// what we actually compute. Big blur, cheap.
+//
+// 2. BLUR_SHADER. The tile-based separable box-blur compute shader from the
+// WebGPU samples repo, used unmodified. One workgroup loads a 128 x 4 tile
+// into shared memory, then 32 threads write filterDim-averaged outputs.
+// A `flip` uniform swaps x/y so the same shader does both axes. Iterating
+// H-V passes (the variance adds) approximates a gaussian.
+
+export const PREPASS_SHADER = /* wgsl */ `
+struct PrepassUniforms {
+ texSize: vec2f,
+ canvasSize: vec2f,
+};
+
+@group(0) @binding(0) var srcTex: texture_external;
+@group(0) @binding(1) var srcSampler: sampler;
+@group(0) @binding(2) var u: PrepassUniforms;
+
+struct VsOut {
+ @builtin(position) position: vec4f,
+ @location(0) uv: vec2f,
+};
+
+@vertex
+fn vs_main(@builtin(vertex_index) vid: u32) -> VsOut {
+ var positions = array(
+ vec2f(-1.0, -3.0),
+ vec2f(-1.0, 1.0),
+ vec2f( 3.0, 1.0),
+ );
+ var uvs = array(
+ vec2f(0.0, 2.0),
+ vec2f(0.0, 0.0),
+ vec2f(2.0, 0.0),
+ );
+ var out: VsOut;
+ out.position = vec4f(positions[vid], 0.0, 1.0);
+ out.uv = uvs[vid];
+ return out;
+}
+
+@fragment
+fn fs_main(in: VsOut) -> @location(0) vec4f {
+ // Cover-fit, same math the worklet sends via uvScale to the main pass so
+ // the blurred image lines up with the unblurred path.
+ let canvasAR = u.canvasSize.x / u.canvasSize.y;
+ let texAR = u.texSize.x / u.texSize.y;
+ var scale: vec2f;
+ if (texAR > canvasAR) {
+ scale = vec2f(canvasAR / texAR, 1.0);
+ } else {
+ scale = vec2f(1.0, texAR / canvasAR);
+ }
+ let uv = vec2f(0.5) + (in.uv - vec2f(0.5)) * scale;
+ // cameraCoord (vertical flip) + cameraDecode (YUV->RGB) come from
+ // CAMERA_PRELUDE, prepended when this module is compiled. They are no-ops on
+ // iOS and handle the Android opaque-YUV case.
+ let c = cameraDecode(textureSampleBaseClampToEdge(
+ srcTex,
+ srcSampler,
+ cameraCoord(clamp(uv, vec2f(0.0), vec2f(1.0))),
+ ));
+ return vec4f(c.rgb, 1.0);
+}
+`;
+
+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 dims = vec2i(textureDimensions(inputTex, 0));
+ let baseIndex = vec2i(WorkGroupID.xy * vec2(params.blockDim, 4u) +
+ LocalInvocationID.xy * vec2(4u, 1u))
+ - vec2(filterOffset, 0);
+
+ for (var r = 0; r < 4; r++) {
+ for (var c = 0; c < 4; c++) {
+ var loadIndex = baseIndex + vec2(c, r);
+ if (flip.value != 0u) {
+ loadIndex = loadIndex.yx;
+ }
+ tile[r][4u * LocalInvocationID.x + u32(c)] = textureSampleLevel(
+ inputTex,
+ samp,
+ (vec2f(loadIndex) + vec2f(0.5)) / vec2f(dims),
+ 0.0,
+ ).rgb;
+ }
+ }
+
+ workgroupBarrier();
+
+ for (var r = 0; r < 4; r++) {
+ for (var c = 0; c < 4; c++) {
+ var writeIndex = baseIndex + vec2(c, r);
+ if (flip.value != 0u) {
+ writeIndex = writeIndex.yx;
+ }
+ let center = i32(4u * LocalInvocationID.x) + c;
+ if (center >= filterOffset &&
+ center < 128 - filterOffset &&
+ all(writeIndex < dims)) {
+ var acc = vec3f(0.0);
+ for (var f = 0; f < params.filterDim; f++) {
+ let i = center + f - filterOffset;
+ acc = acc + (1.0 / f32(params.filterDim)) * tile[r][i];
+ }
+ textureStore(outputTex, writeIndex, vec4f(acc, 1.0));
+ }
+ }
+ }
+}
+`;
diff --git a/apps/example/src/VisionCamera/features.ts b/apps/example/src/VisionCamera/features.ts
new file mode 100644
index 000000000..eb46bc224
--- /dev/null
+++ b/apps/example/src/VisionCamera/features.ts
@@ -0,0 +1,48 @@
+const EFFECT_LABELS = ["Off", "Gray", "Sepia", "Invert", "Vibrant"] as const;
+const TINT_LABELS = ["Off", "Warm", "Cool"] as const;
+const ABERRATION_LABELS = ["Off", "Soft", "Strong"] as const;
+const BLUR_LABELS = ["Off", "Strong", "Overlay"] as const;
+const TOGGLE_LABELS = ["Off", "On"] as const;
+
+export type Modes = {
+ effect: number;
+ tint: number;
+ aberration: number;
+ blur: number;
+ vignette: number;
+ pixelate: number;
+};
+
+export const INITIAL_MODES: Modes = {
+ effect: 0,
+ tint: 0,
+ aberration: 0,
+ blur: 0,
+ vignette: 0,
+ pixelate: 0,
+};
+
+// Aberration strength in UV units per level. Soft matches the original demo.
+export const ABERRATION_STRENGTHS = [0.0, 0.006, 0.018] as const;
+// Block size in UV units per pixelate level. Larger value, chunkier pixels.
+export const PIXELATE_BLOCKS = [0.0, 0.02] as const;
+// H-V iteration count per blur level. Each iteration is a separable box pass
+// at 1/4 canvas res; variance adds, so the effective sigma grows as sqrt(N).
+// Strong and Overlay share the same iteration count, the difference is in the
+// fragment shader (full-canvas blurred vs. card-shaped sharp inset).
+export const BLUR_ITERATIONS = [0, 3, 3] as const;
+
+export type Feature = {
+ title: string;
+ key: keyof Modes;
+ labels: readonly string[];
+};
+
+export const FEATURES: Feature[] = [
+ { title: "Effect", key: "effect", labels: EFFECT_LABELS },
+ { title: "Tint", key: "tint", labels: TINT_LABELS },
+ { title: "Aberration", key: "aberration", labels: ABERRATION_LABELS },
+ { title: "Blur", key: "blur", labels: BLUR_LABELS },
+ { title: "Vignette", key: "vignette", labels: TOGGLE_LABELS },
+ { title: "Pixelate", key: "pixelate", labels: TOGGLE_LABELS },
+];
diff --git a/apps/example/src/VisionCamera/index.ts b/apps/example/src/VisionCamera/index.ts
new file mode 100644
index 000000000..fc62066eb
--- /dev/null
+++ b/apps/example/src/VisionCamera/index.ts
@@ -0,0 +1 @@
+export * from "./VisionCamera";
diff --git a/apps/example/src/VisionCamera/shaders.ts b/apps/example/src/VisionCamera/shaders.ts
new file mode 100644
index 000000000..e93520107
--- /dev/null
+++ b/apps/example/src/VisionCamera/shaders.ts
@@ -0,0 +1,219 @@
+import { Platform } from "react-native";
+
+// Main camera-effects shader: samples the imported external texture (camera
+// frame) with cover-fit + optional chromatic aberration / pixelate, optionally
+// mixes in the pre-blurred backdrop, then applies effect / tint / vignette.
+// `cameraCoord` / `cameraDecode` come from CAMERA_PRELUDE, which is prepended
+// at shader-module creation time.
+export const SHADER = /* wgsl */ `
+struct VsOut {
+ @builtin(position) position: vec4f,
+ @location(0) uv: vec2f,
+};
+
+struct Uniforms {
+ // x, y: 'cover'-fit UV scale around (0.5, 0.5).
+ // z: chromatic aberration offset in UV units (0 disables).
+ // w: pixelate block size in UV units (0 disables).
+ params: vec4f,
+ // x: effect (0 off, 1 gray, 2 sepia, 3 invert, 4 vibrant)
+ // y: tint (0 off, 1 warm, 2 cool)
+ // z: vignette (0 off, 1 on)
+ // w: blurMode (0 off, 1 strong - blurred everywhere (prepass bakes
+ // cover-fit), 2 overlay - blurred backdrop + sharp card)
+ modes: vec4u,
+};
+
+@group(0) @binding(0) var srcTex: texture_external;
+@group(0) @binding(1) var srcSampler: sampler;
+@group(0) @binding(2) var u: Uniforms;
+@group(0) @binding(3) var blurredTex: texture_2d;
+
+@vertex
+fn vs_main(@builtin(vertex_index) vid: u32) -> VsOut {
+ var positions = array(
+ vec2f(-1.0, -3.0),
+ vec2f(-1.0, 1.0),
+ vec2f( 3.0, 1.0),
+ );
+ var uvs = array(
+ vec2f(0.0, 2.0),
+ vec2f(0.0, 0.0),
+ vec2f(2.0, 0.0),
+ );
+ var out: VsOut;
+ out.position = vec4f(positions[vid], 0.0, 1.0);
+ out.uv = uvs[vid];
+ return out;
+}
+
+fn snap(uv: vec2f, block: f32) -> vec2f {
+ if (block <= 0.0) {
+ return uv;
+ }
+ return (floor(uv / block) + vec2f(0.5)) * block;
+}
+
+fn sampleExternal(uv: vec2f, block: f32) -> vec4f {
+ return cameraDecode(
+ textureSampleBaseClampToEdge(srcTex, srcSampler, cameraCoord(snap(uv, block))),
+ );
+}
+
+fn sampleBlurred(uv: vec2f, block: f32) -> vec4f {
+ return textureSampleLevel(
+ blurredTex,
+ srcSampler,
+ clamp(snap(uv, block), vec2f(0.0), vec2f(1.0)),
+ 0.0,
+ );
+}
+
+// RGB-split sample of the external camera with cover-fit + optional pixelate.
+fn rgbSplitExternal(uv: vec2f, aberration: f32, block: f32) -> vec3f {
+ let r = sampleExternal(uv + vec2f( aberration, 0.0), block).r;
+ let g = sampleExternal(uv, block).g;
+ let b = sampleExternal(uv + vec2f(-aberration, 0.0), block).b;
+ return vec3f(r, g, b);
+}
+
+// RGB-split sample of the pre-blurred 2D texture (cover-fit baked in).
+fn rgbSplitBlurred(uv: vec2f, aberration: f32, block: f32) -> vec3f {
+ let r = sampleBlurred(uv + vec2f( aberration, 0.0), block).r;
+ let g = sampleBlurred(uv, block).g;
+ let b = sampleBlurred(uv + vec2f(-aberration, 0.0), block).b;
+ return vec3f(r, g, b);
+}
+
+fn applyEffect(rgb: vec3f, mode: u32) -> vec3f {
+ if (mode == 1u) {
+ let l = dot(rgb, vec3f(0.2126, 0.7152, 0.0722));
+ return vec3f(l);
+ }
+ if (mode == 2u) {
+ return vec3f(
+ dot(rgb, vec3f(0.393, 0.769, 0.189)),
+ dot(rgb, vec3f(0.349, 0.686, 0.168)),
+ dot(rgb, vec3f(0.272, 0.534, 0.131))
+ );
+ }
+ if (mode == 3u) {
+ return vec3f(1.0) - rgb;
+ }
+ if (mode == 4u) {
+ let l = dot(rgb, vec3f(0.2126, 0.7152, 0.0722));
+ let sat = mix(vec3f(l), rgb, 1.55);
+ return clamp((sat - 0.5) * 1.18 + 0.5, vec3f(0.0), vec3f(1.0));
+ }
+ return rgb;
+}
+
+fn applyTint(rgb: vec3f, mode: u32) -> vec3f {
+ if (mode == 1u) {
+ return clamp(rgb * vec3f(1.10, 1.02, 0.86), vec3f(0.0), vec3f(1.0));
+ }
+ if (mode == 2u) {
+ return clamp(rgb * vec3f(0.86, 0.98, 1.16), vec3f(0.0), vec3f(1.0));
+ }
+ return rgb;
+}
+
+@fragment
+fn fs_main(in: VsOut) -> @location(0) vec4f {
+ let uvScale = u.params.xy;
+ let aberration = u.params.z;
+ let pixelate = u.params.w;
+ let effect = u.modes.x;
+ let tint = u.modes.y;
+ let vignette = u.modes.z;
+ let blurMode = u.modes.w;
+
+ // Overlay card geometry. NDC-space rect, the camera image is cover-fit
+ // inside the card (same uvScale as the off path), the blurred backdrop
+ // fills outside.
+ let overlayPadding = 0.08;
+ let edgeAA = 0.004;
+
+ var color: vec3f;
+ if (blurMode == 1u) {
+ // Strong: prepass already baked cover-fit, so feed in.uv straight in.
+ color = rgbSplitBlurred(in.uv, aberration, pixelate);
+ } else if (blurMode == 2u) {
+ // Overlay: per-fragment edge factor 0 strictly inside the card,
+ // 1 strictly outside, smoothstep band for AA. Uniform within the branch
+ // (blurMode came from the uniform buffer), so calling
+ // textureSampleBaseClampToEdge on the external texture is allowed.
+ let cardHalf = vec2f(0.5 - overlayPadding);
+ let p = abs(in.uv - vec2f(0.5));
+ let edgeDist = max(p.x - cardHalf.x, p.y - cardHalf.y);
+ let outside = smoothstep(-edgeAA, edgeAA, edgeDist);
+
+ let cardUv = (in.uv - vec2f(overlayPadding)) /
+ (1.0 - 2.0 * overlayPadding);
+ let sharpUv = vec2f(0.5) + (cardUv - vec2f(0.5)) * uvScale;
+ let sharp = rgbSplitExternal(sharpUv, aberration, pixelate);
+ let backdrop = rgbSplitBlurred(in.uv, aberration, pixelate);
+ color = mix(sharp, backdrop, outside);
+ } else {
+ let uv = vec2f(0.5) + (in.uv - vec2f(0.5)) * uvScale;
+ color = rgbSplitExternal(uv, aberration, pixelate);
+ }
+
+ color = applyEffect(color, effect);
+ color = applyTint(color, tint);
+
+ if (vignette == 1u) {
+ let d = distance(in.uv, vec2f(0.5));
+ let v = 1.0 - smoothstep(0.35, 0.85, d);
+ color = color * v;
+ }
+
+ return vec4f(color, 1.0);
+}
+`;
+
+// Android delivers camera frames as an external-format (opaque) YUV buffer.
+// Dawn's OpaqueYCbCrAndroidForExternalTexture path samples it through a Vulkan
+// YCbCr conversion that is hard-coded to RGB_IDENTITY (Dawn's
+// SamplerVk.cpp::GetYCbCrForTextureView; see crbug.com/497675620), so the
+// external sample comes back as raw [Y, Cb, Cr] on *every* device — this is by
+// design, not a driver quirk — and we do the BT.709 YUV->RGB ourselves below.
+// (Dawn's own SharedTextureMemoryOpaqueYCbCrAndroidForExternalTexture
+// .NoopSampleY8Cb8Cr8AHB test asserts the same raw passthrough.) The frame also
+// comes out mirrored on both axes relative to the canvas (Android buffer
+// origin), so we flip X and Y. iOS goes through the native two-plane path,
+// which already converts and orients, so this correction is Android-only. The
+// prelude is prepended to every shader module that samples the camera (main
+// pass + blur prepass).
+export const CAMERA_PRELUDE = /* wgsl */ `
+const CAMERA_IS_YUV: bool = ${Platform.OS === "android"};
+const CAMERA_FLIP_X: bool = ${Platform.OS === "android"};
+const CAMERA_FLIP_Y: bool = ${Platform.OS === "android"};
+
+fn cameraCoord(uv: vec2f) -> vec2f {
+ var c = uv;
+ if (CAMERA_FLIP_X) {
+ c.x = 1.0 - c.x;
+ }
+ if (CAMERA_FLIP_Y) {
+ c.y = 1.0 - c.y;
+ }
+ return c;
+}
+
+// BT.709 limited-range YUV -> RGB. On the Android opaque path the sampled
+// channels are always raw [Y, Cb, Cr] (Dawn forces an RGB_IDENTITY Vulkan
+// conversion); a no-op passthrough on every other platform.
+fn cameraDecode(c: vec4f) -> vec4f {
+ if (!CAMERA_IS_YUV) {
+ return c;
+ }
+ let y = c.r - 0.0627451;
+ let cb = c.g - 0.5;
+ let cr = c.b - 0.5;
+ let r = 1.164384 * y + 1.792741 * cr;
+ let g = 1.164384 * y - 0.213249 * cb - 0.532909 * cr;
+ let b = 1.164384 * y + 2.112402 * cb;
+ return vec4f(clamp(vec3f(r, g, b), vec3f(0.0), vec3f(1.0)), 1.0);
+}
+`;
diff --git a/packages/webgpu/README.md b/packages/webgpu/README.md
index b3ec7eb8e..8eeb1cba1 100644
--- a/packages/webgpu/README.md
+++ b/packages/webgpu/README.md
@@ -226,19 +226,21 @@ device.queue.copyExternalImageToTexture(
React Native WebGPU exposes Dawn's `SharedTextureMemory` so you can import a native pixel surface (an `IOSurface`-backed `CVPixelBuffer` on iOS, an `AHardwareBuffer` on Android) as a sampleable `GPUTexture` without copying pixels through the CPU. This is the path you want for camera frames, video frames, or anything coming out of a hardware producer.
-We expose a single umbrella feature name, `"rnwebgpu/shared-texture-memory"`. Request it at device creation.
+Like `importExternalTexture` on the web, this is **enabled by default**, there is nothing to request at device creation. The only thing to check is that the device supports it before importing. It always does on iOS/macOS; it can be missing on some Android drivers and emulators.
```tsx
-import type { VideoFrame } from "react-native-wgpu";
-
-const FEATURE = "rnwebgpu/shared-texture-memory" as GPUFeatureName;
+import type { NativeVideoFrame } from "react-native-wgpu";
const adapter = await navigator.gpu.requestAdapter();
-const requiredFeatures = adapter!.features.has(FEATURE) ? [FEATURE] : [];
-const device = await adapter!.requestDevice({ requiredFeatures });
+const device = await adapter!.requestDevice();
+
+// On by default when supported; this is the only check you need.
+if (!device.features.has("rnwebgpu/native-texture" as GPUFeatureName)) {
+ return; // rare: some Android drivers/emulators can't import native surfaces
+}
-// `frame` here is a VideoFrame whose .handle is the native surface
-// (IOSurfaceRef / AHardwareBuffer*). VideoFrames are produced by helpers
+// `frame` here is a NativeVideoFrame whose .handle is the native surface
+// (IOSurfaceRef / AHardwareBuffer*). NativeVideoFrames are produced by helpers
// like RNWebGPU.createVideoPlayer or RNWebGPU.createTestVideoFrame, or by
// any third-party module that hands you a compatible native pointer.
const memory = device.importSharedTextureMemory({
@@ -257,6 +259,53 @@ frame.release();
`beginAccess`/`endAccess` bracket the GPU's read window on the shared surface. Pass `initialized: true` when the producer has already written meaningful pixels (the typical video/camera case) and `false` when the next pass will fully overwrite the texture.
+### Importing External Textures
+
+`GPUDevice.importExternalTexture` is the higher-level path for sampling a native surface. You hand it a `NativeVideoFrame` and get back a `GPUExternalTexture` that you bind as a `texture_external` and read with `textureSampleBaseClampToEdge`. It does two things for you on top of `SharedTextureMemory`:
+
+- **Color conversion.** Camera and video surfaces are usually biplanar YUV (NV12), not RGB. An external texture carries the YUV→RGB matrix and the source/destination color-space transfer functions, so on the supported paths the sampler returns ready-to-use RGB in hardware. With raw `SharedTextureMemory` you would sample the luma/chroma planes and do that conversion by hand in WGSL. This is the main reason to prefer it for camera and video frames.
+- **Lifecycle.** It owns the `SharedTextureMemory` + `createTexture` + `beginAccess`/`endAccess` sequence internally, so you just import the frame and `destroy()` the result.
+
+It builds on the same default-on capability as Shared Texture Memory above, so feature-detect the device the same way before importing.
+
+> **Android note:** the hardware YUV→RGB conversion is fully automatic on iOS (NV12 `IOSurface`). On Android, camera frames arrive as an _opaque_ YCbCr `AHardwareBuffer`, and Dawn's Vulkan path forces an identity (`RGB_IDENTITY`) sampler conversion, so the external sample comes back as raw `[Y, Cb, Cr]`. You still get the zero-copy import and the rotation/mirror transform, but you need to apply the YUV→RGB matrix yourself in the shader. See the `CAMERA_PRELUDE` in the [VisionCamera example](/apps/example/src/VisionCamera/shaders.ts) for a ready-made BT.709 decode.
+
+```tsx
+const adapter = await navigator.gpu.requestAdapter();
+const device = await adapter!.requestDevice();
+// Feature-detect as shown above before importing on unsupported hardware.
+
+const render = () => {
+ // A GPUExternalTexture expires once the queue work that used it is submitted,
+ // so re-import one every frame.
+ const externalTexture = device.importExternalTexture({
+ source: frame, // a NativeVideoFrame
+ label: "video-frame",
+ });
+
+ const bindGroup = device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: 0, resource: externalTexture },
+ { binding: 1, resource: sampler },
+ ],
+ });
+
+ // ... encode a pass that samples `externalTexture`, then:
+ device.queue.submit([encoder.finish()]);
+
+ // Release the surface's access window right after the submit that sampled it.
+ externalTexture.destroy();
+ context.present();
+};
+```
+
+Camera frames arrive in the sensor's native orientation, so `importExternalTexture` also accepts non-spec `rotation` (`0` | `90` | `180` | `270`, in degrees) and `mirrored` (horizontal flip) options. Dawn bakes them into the sampling transform, so the shader sees an upright frame. They map directly onto VisionCamera's `frame.orientation` / `frame.isMirrored`.
+
+#### Calling `destroy()`
+
+A `GPUExternalTexture` keeps an open access window on the underlying native surface until the wrapper is destroyed. On the Web `importExternalTexture` is core and the lifetime is handled for you; here the window is tied to the JavaScript object's lifetime. Call `externalTexture.destroy()` right after the `queue.submit()` that sampled it (never before) to release the surface back to its producer immediately. `destroy()` is idempotent, and the surface is also released when the object is garbage-collected, but relying on GC can starve a producer's buffer pool (e.g. an `AVPlayer`'s recycled `IOSurface`s) and pile up GPU resources, so prefer the explicit call in a render loop.
+
### Reanimated Integration
React Native WebGPU supports running WebGPU rendering on the UI thread using [React Native Reanimated](https://docs.swmansion.com/react-native-reanimated/) and [React Native Worklets](https://github.com/margelo/react-native-worklets).
diff --git a/packages/webgpu/android/CMakeLists.txt b/packages/webgpu/android/CMakeLists.txt
index fcab0baa1..3ef33bb7b 100644
--- a/packages/webgpu/android/CMakeLists.txt
+++ b/packages/webgpu/android/CMakeLists.txt
@@ -38,6 +38,7 @@ add_library(${PACKAGE_NAME} SHARED
../cpp/rnwgpu/api/GPUQuerySet.cpp
../cpp/rnwgpu/api/GPUTexture.cpp
../cpp/rnwgpu/api/GPUSharedTextureMemory.cpp
+ ../cpp/rnwgpu/api/GPUExternalTexture.cpp
../cpp/rnwgpu/api/GPURenderBundleEncoder.cpp
../cpp/rnwgpu/api/GPURenderPassEncoder.cpp
../cpp/rnwgpu/api/GPURenderPipeline.cpp
diff --git a/packages/webgpu/android/cpp/AndroidPlatformContext.h b/packages/webgpu/android/cpp/AndroidPlatformContext.h
index cc57eea6a..080b1bb5e 100644
--- a/packages/webgpu/android/cpp/AndroidPlatformContext.h
+++ b/packages/webgpu/android/cpp/AndroidPlatformContext.h
@@ -280,7 +280,8 @@ class AndroidPlatformContext : public PlatformContext {
}
std::unique_ptr
- createVideoPlayer(const std::string & /*path*/) override {
+ createVideoPlayer(const std::string & /*path*/,
+ VideoPixelFormat /*format*/) override {
// TODO: implement using MediaCodec -> ImageReader (AHardwareBuffer mode).
throw std::runtime_error(
"createVideoPlayer is not yet implemented on Android.");
@@ -291,6 +292,37 @@ class AndroidPlatformContext : public PlatformContext {
throw std::runtime_error(
"writeTestVideoFile is not yet implemented on Android.");
}
+
+ VideoFrameHandle wrapNativeBuffer(void *pointer) override {
+ if (!pointer) {
+ throw std::runtime_error("wrapNativeBuffer: pointer is null");
+ }
+ auto *buffer = static_cast(pointer);
+
+ AHardwareBuffer_Desc desc = {};
+ AHardwareBuffer_describe(buffer, &desc);
+
+ AHardwareBuffer_acquire(buffer);
+
+ VideoFrameHandle handle;
+ handle.handle = static_cast(buffer);
+ handle.width = desc.width;
+ handle.height = desc.height;
+ // YUV / opaque formats route through Vulkan's SamplerYcbcrConversion via
+ // Dawn's OpaqueYCbCrAndroidForExternalTexture path. Single-plane RGBA AHBs
+ // take the plain BGRA8 path (sampled as a regular 2D texture).
+ switch (desc.format) {
+ case AHARDWAREBUFFER_FORMAT_Y8Cb8Cr8_420:
+ case AHARDWAREBUFFER_FORMAT_YCbCr_P010:
+ handle.pixelFormat = VideoPixelFormat::NV12;
+ break;
+ default:
+ handle.pixelFormat = VideoPixelFormat::BGRA8;
+ break;
+ }
+ handle.deleter = [buffer]() { AHardwareBuffer_release(buffer); };
+ return handle;
+ }
};
} // namespace rnwgpu
diff --git a/packages/webgpu/apple/ApplePlatformContext.h b/packages/webgpu/apple/ApplePlatformContext.h
index 4d9b29d1f..6536663c4 100644
--- a/packages/webgpu/apple/ApplePlatformContext.h
+++ b/packages/webgpu/apple/ApplePlatformContext.h
@@ -33,9 +33,12 @@ class ApplePlatformContext : public PlatformContext {
uint32_t height) override;
std::unique_ptr
- createVideoPlayer(const std::string &path) override;
+ createVideoPlayer(const std::string &path,
+ VideoPixelFormat format) override;
std::string writeTestVideoFile() override;
+
+ VideoFrameHandle wrapNativeBuffer(void *pointer) override;
};
} // namespace rnwgpu
diff --git a/packages/webgpu/apple/ApplePlatformContext.mm b/packages/webgpu/apple/ApplePlatformContext.mm
index 959b45b8d..594337cfc 100644
--- a/packages/webgpu/apple/ApplePlatformContext.mm
+++ b/packages/webgpu/apple/ApplePlatformContext.mm
@@ -237,14 +237,19 @@ void checkIfUsingSimulatorWithAPIValidation() {
}
std::unique_ptr
-ApplePlatformContext::createVideoPlayer(const std::string &path) {
- return createAppleVideoPlayer(path);
+ApplePlatformContext::createVideoPlayer(const std::string &path,
+ VideoPixelFormat format) {
+ return createAppleVideoPlayer(path, format);
}
std::string ApplePlatformContext::writeTestVideoFile() {
return writeAppleTestVideoFile();
}
+VideoFrameHandle ApplePlatformContext::wrapNativeBuffer(void *pointer) {
+ return wrapCVPixelBuffer(static_cast(pointer));
+}
+
VideoFrameHandle
ApplePlatformContext::createTestVideoFrame(uint32_t width, uint32_t height) {
NSDictionary *attrs = @{
diff --git a/packages/webgpu/apple/AppleVideoPlayer.h b/packages/webgpu/apple/AppleVideoPlayer.h
index 1866f2440..554d69d76 100644
--- a/packages/webgpu/apple/AppleVideoPlayer.h
+++ b/packages/webgpu/apple/AppleVideoPlayer.h
@@ -5,16 +5,27 @@
#include
#include
+#ifdef __OBJC__
+#import
+#endif
+
namespace rnwgpu {
// Factory: creates a new IVideoPlayer backed by AVPlayer +
-// AVPlayerItemVideoOutput.
+// AVPlayerItemVideoOutput. `format` selects the surface layout.
std::unique_ptr
-createAppleVideoPlayer(const std::string &path);
+createAppleVideoPlayer(const std::string &path, VideoPixelFormat format);
// Generate a small procedurally-animated test video and write it to a
// temporary file. Returns the absolute path. Used by the SharedTextureMemory
// example so it doesn't need a bundled .mp4.
std::string writeAppleTestVideoFile();
+#ifdef __OBJC__
+// Build a VideoFrameHandle from an existing CVPixelBuffer. CFRetains the
+// pixel buffer so the caller can release their reference immediately. Reads
+// IOSurface, dimensions, pixel format, and YUV matrix off the buffer.
+VideoFrameHandle wrapCVPixelBuffer(CVPixelBufferRef pixelBuffer);
+#endif
+
} // namespace rnwgpu
diff --git a/packages/webgpu/apple/AppleVideoPlayer.mm b/packages/webgpu/apple/AppleVideoPlayer.mm
index 4c58f215c..a06144dc7 100644
--- a/packages/webgpu/apple/AppleVideoPlayer.mm
+++ b/packages/webgpu/apple/AppleVideoPlayer.mm
@@ -9,6 +9,62 @@
namespace {
+// 3x4 row-major matrices mapping [Y, Cb, Cr, 1] to gamma-encoded R'G'B' (NOT
+// linear): YCbCr is derived from gamma-encoded RGB, so this conversion stays in
+// the encoded domain. The sRGB decode in srcTransferFunctionParameters
+// linearizes afterward (see GPUExternalTexture.cpp). Limited-range (video
+// range) means luma is 16..235, chroma is 16..240 (8-bit).
+// Reference: https://en.wikipedia.org/wiki/YCbCr (BT.601 / BT.709).
+static constexpr float kBT709LimitedToRgb[12] = {
+ 1.164383f, 0.000000f, 1.792741f, -0.972945f, //
+ 1.164383f, -0.213249f, -0.532909f, 0.301517f, //
+ 1.164383f, 2.112402f, 0.000000f, -1.133402f, //
+};
+static constexpr float kBT601LimitedToRgb[12] = {
+ 1.164383f, 0.000000f, 1.596027f, -0.874202f, //
+ 1.164383f, -0.391762f, -0.812968f, 0.531668f, //
+ 1.164383f, 2.017232f, 0.000000f, -1.085631f, //
+};
+static constexpr float kBT2020LimitedToRgb[12] = {
+ 1.164383f, 0.000000f, 1.678674f, -0.915688f, //
+ 1.164383f, -0.187326f, -0.650424f, 0.347459f, //
+ 1.164383f, 2.141772f, 0.000000f, -1.148145f, //
+};
+
+// Pick the right YUV→RGB matrix from the pixel buffer's color attachments.
+// Falls back to BT.709 limited range (the right call for ≥720p H.264, which
+// is what AVPlayer hands us for Big Buck Bunny and most streamed media).
+static void fillYuvMatrix(CVPixelBufferRef pixelBuffer, float out[12]) {
+ CFTypeRef matrixKey = CVBufferGetAttachment(
+ pixelBuffer, kCVImageBufferYCbCrMatrixKey, nullptr);
+ const float *src = kBT709LimitedToRgb;
+ if (matrixKey) {
+ auto matrix = (CFStringRef)matrixKey;
+ if (CFEqual(matrix, kCVImageBufferYCbCrMatrix_ITU_R_601_4) ||
+ CFEqual(matrix, kCVImageBufferYCbCrMatrix_SMPTE_240M_1995)) {
+ src = kBT601LimitedToRgb;
+ } else if (CFEqual(matrix, kCVImageBufferYCbCrMatrix_ITU_R_2020)) {
+ src = kBT2020LimitedToRgb;
+ }
+ }
+ for (int i = 0; i < 12; ++i) {
+ out[i] = src[i];
+ }
+}
+
+// Map a CVPixelBuffer's pixel format to our VideoPixelFormat enum.
+static VideoPixelFormat pixelFormatFromCVPixelBuffer(
+ CVPixelBufferRef pixelBuffer) {
+ OSType type = CVPixelBufferGetPixelFormatType(pixelBuffer);
+ switch (type) {
+ case kCVPixelFormatType_420YpCbCr8BiPlanarVideoRange:
+ case kCVPixelFormatType_420YpCbCr8BiPlanarFullRange:
+ return VideoPixelFormat::NV12;
+ default:
+ return VideoPixelFormat::BGRA8;
+ }
+}
+
class AppleVideoPlayer : public IVideoPlayer {
public:
AppleVideoPlayer(AVPlayer *player, AVPlayerItemVideoOutput *output,
@@ -30,25 +86,23 @@ VideoFrameHandle copyLatestFrame() override {
if (![_output hasNewPixelBufferForItemTime:currentTime]) {
return {};
}
+ // copyPixelBufferForItemTime returns a +1 retained CVPixelBuffer; we then
+ // hand it to wrapCVPixelBuffer which adds another retain. Balance with a
+ // CFRelease here so we don't leak.
CVPixelBufferRef pixelBuffer =
[_output copyPixelBufferForItemTime:currentTime
itemTimeForDisplay:nullptr];
if (!pixelBuffer) {
return {};
}
-
- IOSurfaceRef ioSurface = CVPixelBufferGetIOSurface(pixelBuffer);
- if (!ioSurface) {
+ try {
+ auto handle = wrapCVPixelBuffer(pixelBuffer);
CFRelease(pixelBuffer);
- return {};
+ return handle;
+ } catch (...) {
+ CFRelease(pixelBuffer);
+ throw;
}
-
- VideoFrameHandle handle;
- handle.handle = (void *)ioSurface;
- handle.width = static_cast(CVPixelBufferGetWidth(pixelBuffer));
- handle.height = static_cast(CVPixelBufferGetHeight(pixelBuffer));
- handle.deleter = [pixelBuffer]() { CFRelease(pixelBuffer); };
- return handle;
}
void play() override { [_player play]; }
@@ -62,8 +116,34 @@ VideoFrameHandle copyLatestFrame() override {
} // namespace
+VideoFrameHandle wrapCVPixelBuffer(CVPixelBufferRef pixelBuffer) {
+ if (!pixelBuffer) {
+ throw std::runtime_error("wrapCVPixelBuffer: pointer is null");
+ }
+ IOSurfaceRef ioSurface = CVPixelBufferGetIOSurface(pixelBuffer);
+ if (!ioSurface) {
+ throw std::runtime_error(
+ "wrapCVPixelBuffer: pixel buffer is not IOSurface-backed (was the "
+ "camera/video pipeline configured for Metal/IOSurface output?)");
+ }
+
+ // Retain the pixel buffer so the caller can release theirs immediately.
+ CFRetain(pixelBuffer);
+
+ VideoFrameHandle handle;
+ handle.handle = (void *)ioSurface;
+ handle.width = static_cast(CVPixelBufferGetWidth(pixelBuffer));
+ handle.height = static_cast(CVPixelBufferGetHeight(pixelBuffer));
+ handle.pixelFormat = pixelFormatFromCVPixelBuffer(pixelBuffer);
+ if (handle.pixelFormat == VideoPixelFormat::NV12) {
+ fillYuvMatrix(pixelBuffer, handle.yuvToRgbMatrix);
+ }
+ handle.deleter = [pixelBuffer]() { CFRelease(pixelBuffer); };
+ return handle;
+}
+
std::unique_ptr
-createAppleVideoPlayer(const std::string &path) {
+createAppleVideoPlayer(const std::string &path, VideoPixelFormat format) {
NSString *nsPath = [NSString stringWithUTF8String:path.c_str()];
NSURL *url;
if ([nsPath hasPrefix:@"http://"] || [nsPath hasPrefix:@"https://"] ||
@@ -79,9 +159,15 @@ VideoFrameHandle copyLatestFrame() override {
"AVPlayerItem");
}
+ // NV12 (kCVPixelFormatType_420YpCbCr8BiPlanarVideoRange) lets us hand the
+ // IOSurface straight to Dawn as a multi-planar texture for
+ // importExternalTexture. BGRA is the "decode + convert" path for the
+ // single-plane SharedTextureMemory demo.
+ OSType pixelFormat = format == VideoPixelFormat::NV12
+ ? kCVPixelFormatType_420YpCbCr8BiPlanarVideoRange
+ : kCVPixelFormatType_32BGRA;
NSDictionary *outputSettings = @{
- (NSString *)kCVPixelBufferPixelFormatTypeKey :
- @(kCVPixelFormatType_32BGRA),
+ (NSString *)kCVPixelBufferPixelFormatTypeKey : @(pixelFormat),
(NSString *)kCVPixelBufferIOSurfacePropertiesKey : @{},
(NSString *)kCVPixelBufferMetalCompatibilityKey : @YES,
};
diff --git a/packages/webgpu/cpp/rnwgpu/PlatformContext.h b/packages/webgpu/cpp/rnwgpu/PlatformContext.h
index b34b15a31..fec049256 100644
--- a/packages/webgpu/cpp/rnwgpu/PlatformContext.h
+++ b/packages/webgpu/cpp/rnwgpu/PlatformContext.h
@@ -18,6 +18,16 @@ struct ImageData {
wgpu::TextureFormat format;
};
+// Pixel layout of a VideoFrame. Determines whether the underlying surface is
+// a single RGBA plane or a biplanar Y / CbCr pair.
+enum class VideoPixelFormat {
+ // Single-plane 8-bit BGRA (default; what RGBA-style sampling expects).
+ BGRA8,
+ // Biplanar 4:2:0 8-bit Y + interleaved CbCr (NV12). Used for the
+ // importExternalTexture path; needs the YUV→RGB conversion matrix below.
+ NV12,
+};
+
// A native handle to a video frame that can be imported into a
// GPUSharedTextureMemory.
//
@@ -31,6 +41,12 @@ struct VideoFrameHandle {
void *handle = nullptr;
uint32_t width = 0;
uint32_t height = 0;
+ VideoPixelFormat pixelFormat = VideoPixelFormat::BGRA8;
+ // 3x4 row-major matrix mapping [Y, U, V, 1] → linear [R, G, B]. Pre-computed
+ // at decode time from CVPixelBuffer attachments (kCVImageBufferYCbCrMatrixKey
+ // + range), with a BT.709 limited-range default. Only meaningful when
+ // pixelFormat == NV12.
+ float yuvToRgbMatrix[12] = {};
std::function deleter;
};
@@ -87,8 +103,20 @@ class PlatformContext {
// Open a video file at `path` for playback. The returned player yields
// IOSurface / AHardwareBuffer-backed frames via copyLatestFrame().
+ //
+ // `format` selects the requested pixel layout. BGRA8 is the easiest target
+ // for a regular sampled GPUTexture; NV12 is the right shape for the
+ // importExternalTexture path (zero-copy biplanar YUV).
virtual std::unique_ptr
- createVideoPlayer(const std::string &path) = 0;
+ createVideoPlayer(const std::string &path, VideoPixelFormat format) = 0;
+
+ // Wrap a CVPixelBufferRef (Apple) or AHardwareBuffer* (Android) pointer
+ // obtained from another library (typically VisionCamera's
+ // Frame.getNativeBuffer().pointer) as one of our VideoFrame handles.
+ //
+ // We CFRetain / AHardwareBuffer_acquire on the way in, so callers can
+ // safely release their own reference immediately after.
+ virtual VideoFrameHandle wrapNativeBuffer(void *pointer) = 0;
// Write a small procedurally-generated test video to a temporary location
// and return its absolute path. Lets the SharedTextureMemory example play
diff --git a/packages/webgpu/cpp/rnwgpu/RNWebGPUManager.cpp b/packages/webgpu/cpp/rnwgpu/RNWebGPUManager.cpp
index 38f675d37..f874c14e6 100644
--- a/packages/webgpu/cpp/rnwgpu/RNWebGPUManager.cpp
+++ b/packages/webgpu/cpp/rnwgpu/RNWebGPUManager.cpp
@@ -67,6 +67,12 @@ RNWebGPUManager::RNWebGPUManager(
auto rnWebGPU =
std::make_shared(gpu, _platformContext, _jsCallInvoker);
_gpu = gpu->get();
+
+ // RNWebGPU needs its brand registered in NativeObjectRegistry so the boxing
+ // path can install the prototype on worklet runtimes. installConstructor
+ // does that registration but also sets globalThis.RNWebGPU = ctor, so we
+ // call it FIRST and then overwrite the global with the actual instance.
+ RNWebGPU::installConstructor(*_jsRuntime);
_jsRuntime->global().setProperty(*_jsRuntime, "RNWebGPU",
RNWebGPU::create(*_jsRuntime, rnWebGPU));
diff --git a/packages/webgpu/cpp/rnwgpu/api/Convertors.h b/packages/webgpu/cpp/rnwgpu/api/Convertors.h
index e168afcba..2422d5020 100644
--- a/packages/webgpu/cpp/rnwgpu/api/Convertors.h
+++ b/packages/webgpu/cpp/rnwgpu/api/Convertors.h
@@ -252,13 +252,24 @@ class Convertor {
[[nodiscard]] bool Convert(wgpu::BindGroupLayoutEntry &out,
const GPUBindGroupLayoutEntry &in) {
- return Convert(out.binding, in.binding) &&
- Convert(out.visibility, in.visibility) &&
- Convert(out.buffer, in.buffer) && Convert(out.sampler, in.sampler) &&
- Convert(out.texture, in.texture) &&
- Convert(out.storageTexture, in.storageTexture);
- // no external textures here
- //&& Convert(out.externalTexture, in.externalTexture);
+ out = {};
+ if (!Convert(out.binding, in.binding) ||
+ !Convert(out.visibility, in.visibility) ||
+ !Convert(out.buffer, in.buffer) || !Convert(out.sampler, in.sampler) ||
+ !Convert(out.texture, in.texture) ||
+ !Convert(out.storageTexture, in.storageTexture)) {
+ return false;
+ }
+ if (in.externalTexture.has_value() &&
+ in.externalTexture.value() != nullptr) {
+ // External texture layouts bind via a chained struct rather than a
+ // direct field on BindGroupLayoutEntry. The chained struct must outlive
+ // the BindGroupLayoutEntry until Device::CreateBindGroupLayout returns,
+ // so we allocate it on the Convertor's arena.
+ auto *chain = Allocate();
+ out.nextInChain = chain;
+ }
+ return true;
}
[[nodiscard]] bool Convert(wgpu::BlendComponent &out,
@@ -422,9 +433,11 @@ class Convertor {
}
[[nodiscard]] bool Convert(wgpu::ExternalTextureBindingLayout &out,
- const GPUExternalTextureBindingLayout &in) {
- // no external textures at the moment
- return false;
+ const GPUExternalTextureBindingLayout & /*in*/) {
+ // ExternalTextureBindingLayout carries no fields of its own; its presence
+ // (as a chained struct) is what marks the entry as an external texture.
+ out = {};
+ return true;
}
[[nodiscard]] bool Convert(wgpu::ConstantEntry &out, const std::string &key,
@@ -729,7 +742,16 @@ class Convertor {
out.buffer = buffer->get();
return true;
}
- // Not external textures at the moment
+ if (in.externalTexture != nullptr) {
+ // External textures bind via a chained struct rather than a direct field
+ // on BindGroupEntry. The chained struct must outlive the
+ // BindGroupEntry until Device::CreateBindGroup returns, so we allocate
+ // it on the Convertor's arena.
+ auto *chain = Allocate();
+ chain->externalTexture = in.externalTexture->get();
+ out.nextInChain = chain;
+ return true;
+ }
return false;
}
diff --git a/packages/webgpu/cpp/rnwgpu/api/GPU.cpp b/packages/webgpu/cpp/rnwgpu/api/GPU.cpp
index 764a9aa32..11530f4da 100644
--- a/packages/webgpu/cpp/rnwgpu/api/GPU.cpp
+++ b/packages/webgpu/cpp/rnwgpu/api/GPU.cpp
@@ -20,6 +20,33 @@ GPU::GPU(jsi::Runtime &runtime) : NativeObject(CLASS_NAME) {
wgpu::InstanceLimits limits{.timedWaitAnyMaxCount = 64};
instanceDesc.requiredLimits = &limits;
+
+ // Expose Dawn's experimental adapter features. Several features needed by
+ // our Android external-texture path (YCbCrVulkanSamplers,
+ // OpaqueYCbCrAndroidForExternalTexture) are tagged Experimental in Dawn's
+ // feature table and are otherwise filtered out of adapter.features by
+ // PhysicalDeviceBase::GetSupportedFeatures. The allow_unsafe_apis toggle
+ // disables that filter so the features become visible;
+ // expose_wgsl_experimental_features is the parallel toggle for WGSL language
+ // features.
+ //
+ // Trade-off: these are instance-level Dawn toggles, so they apply to every
+ // adapter/device created from this instance, not just the external-texture
+ // path. There is no finer-grained, per-feature mechanism to surface these.
+ // The exposure is acceptable because the toggle only *un-hides* experimental
+ // features in adapter.features; it does not enable any of them. Nothing
+ // experimental becomes active unless application code explicitly lists that
+ // feature in requiredFeatures at device creation, so the default behavior of
+ // a device that asks for no experimental features is unchanged.
+ static const char *const kEnabledToggles[] = {
+ "allow_unsafe_apis",
+ "expose_wgsl_experimental_features",
+ };
+ wgpu::DawnTogglesDescriptor toggles;
+ toggles.enabledToggleCount = std::size(kEnabledToggles);
+ toggles.enabledToggles = kEnabledToggles;
+ instanceDesc.nextInChain = &toggles;
+
_instance = wgpu::CreateInstance(&instanceDesc);
auto dispatcher = std::make_shared(runtime);
diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUAdapter.cpp b/packages/webgpu/cpp/rnwgpu/api/GPUAdapter.cpp
index 57f77b625..7bdb21840 100644
--- a/packages/webgpu/cpp/rnwgpu/api/GPUAdapter.cpp
+++ b/packages/webgpu/cpp/rnwgpu/api/GPUAdapter.cpp
@@ -1,5 +1,6 @@
#include "GPUAdapter.h"
+#include
#include
#include
#include
@@ -18,6 +19,47 @@ namespace rnwgpu {
async::AsyncTaskHandle GPUAdapter::requestDevice(
std::optional> descriptor) {
+ // Enable the react-native-wgpu "native-texture" umbrella by default, mirroring
+ // the web where importExternalTexture is core and needs no feature request.
+ // We append the umbrella's backing Dawn features to requiredFeatures so the
+ // capability is on without the caller listing it. Two rules keep this safe:
+ // - All-or-nothing: only inject when the adapter supports *every* backing
+ // feature (same semantics as maybeSynthesizeRnNativeTextureFeature). On a
+ // web/fallback adapter the backing set is empty or unsupported, so this is
+ // a no-op and device creation is unaffected.
+ // - Requesting a feature the adapter doesn't support makes RequestDevice
+ // fail, hence the support check below.
+ // Callers can still pass "rnwebgpu/native-texture" explicitly; the dedupe
+ // keeps that idempotent.
+ {
+ auto backing = rnNativeTextureBackingFeatures();
+ if (!backing.empty()) {
+ wgpu::SupportedFeatures supported;
+ _instance.GetFeatures(&supported);
+ std::unordered_set supportedSet(
+ supported.features, supported.features + supported.featureCount);
+ bool allSupported = std::all_of(
+ backing.begin(), backing.end(),
+ [&](wgpu::FeatureName f) { return supportedSet.count(f) > 0; });
+ if (allSupported) {
+ if (!descriptor.has_value()) {
+ descriptor = std::make_shared();
+ }
+ auto &desc = descriptor.value();
+ if (!desc->requiredFeatures.has_value()) {
+ desc->requiredFeatures = std::vector{};
+ }
+ auto &features = desc->requiredFeatures.value();
+ for (auto f : backing) {
+ if (std::find(features.begin(), features.end(), f) ==
+ features.end()) {
+ features.push_back(f);
+ }
+ }
+ }
+ }
+ }
+
wgpu::DeviceDescriptor aDescriptor;
Convertor conv;
if (!conv(aDescriptor, descriptor)) {
@@ -178,7 +220,7 @@ std::unordered_set GPUAdapter::getFeatures() {
result.insert(name);
}
}
- maybeSynthesizeRnSharedTextureMemoryFeature(enabled, result);
+ maybeSynthesizeRnNativeTextureFeature(enabled, result);
return result;
}
diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUDevice.cpp b/packages/webgpu/cpp/rnwgpu/api/GPUDevice.cpp
index f80d7fadf..4d6c92ffa 100644
--- a/packages/webgpu/cpp/rnwgpu/api/GPUDevice.cpp
+++ b/packages/webgpu/cpp/rnwgpu/api/GPUDevice.cpp
@@ -235,8 +235,10 @@ std::shared_ptr GPUDevice::createPipelineLayout(
std::shared_ptr GPUDevice::importExternalTexture(
std::shared_ptr descriptor) {
- throw std::runtime_error(
- "GPUDevice::importExternalTexture(): Not implemented");
+ // The import / begin-access / descriptor-build logic, plus the matching
+ // EndAccess, all live on GPUExternalTexture so the begin/end lifecycle stays
+ // in one translation unit (see GPUExternalTexture.cpp).
+ return GPUExternalTexture::Create(_instance, std::move(descriptor));
}
std::shared_ptr GPUDevice::importSharedTextureMemory(
@@ -442,7 +444,7 @@ std::unordered_set GPUDevice::getFeatures() {
convertEnumToJSUnion(feature, &name);
result.insert(name);
}
- maybeSynthesizeRnSharedTextureMemoryFeature(enabled, result);
+ maybeSynthesizeRnNativeTextureFeature(enabled, result);
return result;
}
diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUExternalTexture.cpp b/packages/webgpu/cpp/rnwgpu/api/GPUExternalTexture.cpp
new file mode 100644
index 000000000..255cd6f80
--- /dev/null
+++ b/packages/webgpu/cpp/rnwgpu/api/GPUExternalTexture.cpp
@@ -0,0 +1,335 @@
+#include "GPUExternalTexture.h"
+
+#include
+#include
+#include
+#include
+
+#include "GPUExternalTextureDescriptor.h"
+
+namespace rnwgpu {
+
+// Identity gamut (BT.709 -> sRGB, same primaries) as a 3x3 column-major matrix.
+static const float kIdentityGamutMatrix[9] = {
+ 1.0f, 0.0f, 0.0f, //
+ 0.0f, 1.0f, 0.0f, //
+ 0.0f, 0.0f, 1.0f, //
+};
+
+// Piecewise gamma transfer-function parameters Dawn expects:
+// for |x| < D: y = sign(x) * (C * |x| + F)
+// else : y = sign(x) * (pow(A * |x| + B, G) + E)
+// sRGB decode (encoded -> linear).
+static const float kSrgbDecodeParams[7] = {
+ 2.4f, // G
+ 1.0f / 1.055f, // A
+ 0.055f / 1.055f, // B
+ 1.0f / 12.92f, // C
+ 0.04045f, // D
+ 0.0f, // E
+ 0.0f, // F
+};
+// sRGB encode (linear -> encoded).
+static const float kSrgbEncodeParams[7] = {
+ 1.0f / 2.4f, // G
+ 1.055f, // A
+ 0.0f, // B
+ 12.92f, // C
+ 0.0031308f, // D
+ -0.055f, // E
+ 0.0f, // F
+};
+
+// Identity transfer (y = x). Used when the sampled surface is already in the
+// render target's color space: a single-plane BGRA IOSurface, or the Android
+// opaque-YCbCr path where the Vulkan sampler already produced RGB. Dawn
+// dereferences the transfer-function arrays unconditionally
+// (ComputeExternalTextureParams), so these must be non-null even when no
+// conversion is wanted.
+static const float kIdentityTransferParams[7] = {
+ 1.0f, // G
+ 1.0f, // A
+ 0.0f, // B
+ 0.0f, // C
+ 0.0f, // D
+ 0.0f, // E
+ 0.0f, // F
+};
+
+// BT.709 limited-range YUV -> R'G'B' as a 3x4 row-major matrix mapping
+// [Y, Cb, Cr, 1] to gamma-encoded R'G'B' (NOT linear; the sRGB decode in
+// srcTransferFunctionParameters linearizes afterwards). Same values the Apple
+// NV12 path computes from the CVPixelBuffer; used for Android buffers that
+// arrive as a *defined* biplanar format (where we split the planes and convert
+// ourselves) rather than an opaque external-format AHB. Camera streams are
+// limited-range BT.709 in the overwhelming majority of cases; full-range /
+// BT.601 would need different coefficients (refine from the buffer's suggested
+// range if it matters).
+[[maybe_unused]] static const float kBT709LimitedToRgb[12] = {
+ 1.164383f, 0.000000f, 1.792741f, -0.972945f, //
+ 1.164383f, -0.213249f, -0.532909f, 0.301517f, //
+ 1.164383f, 2.112402f, 0.000000f, -1.133402f, //
+};
+
+// True for the multi-planar Y + CbCr formats whose planes we can view as
+// Plane0Only (luma) / Plane1Only (chroma) and convert with an explicit matrix.
+// Excludes OpaqueYCbCrAndroid (external format, no plane views) and triplanar
+// formats (would need a third plane). Only referenced on Android.
+[[maybe_unused]] static bool isBiplanarYuvFormat(wgpu::TextureFormat format) {
+ switch (format) {
+ case wgpu::TextureFormat::R8BG8Biplanar420Unorm:
+ case wgpu::TextureFormat::R8BG8Biplanar422Unorm:
+ case wgpu::TextureFormat::R8BG8Biplanar444Unorm:
+ case wgpu::TextureFormat::R10X6BG10X6Biplanar420Unorm:
+ case wgpu::TextureFormat::R10X6BG10X6Biplanar422Unorm:
+ case wgpu::TextureFormat::R10X6BG10X6Biplanar444Unorm:
+ return true;
+ default:
+ return false;
+ }
+}
+
+// Map a rotation in degrees (0 / 90 / 180 / 270) to Dawn's enum. Anything that
+// isn't a clean multiple of 90 snaps to the nearest quadrant; Dawn only
+// supports those four steps for external textures.
+static wgpu::ExternalTextureRotation
+toExternalTextureRotation(double degrees) {
+ int quadrant = static_cast(std::lround(degrees / 90.0)) & 3;
+ switch (quadrant) {
+ case 1:
+ return wgpu::ExternalTextureRotation::Rotate90Degrees;
+ case 2:
+ return wgpu::ExternalTextureRotation::Rotate180Degrees;
+ case 3:
+ return wgpu::ExternalTextureRotation::Rotate270Degrees;
+ default:
+ return wgpu::ExternalTextureRotation::Rotate0Degrees;
+ }
+}
+
+std::shared_ptr GPUExternalTexture::Create(
+ wgpu::Device device,
+ std::shared_ptr descriptor) {
+ if (!descriptor || !descriptor->source) {
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): descriptor.source (VideoFrame) "
+ "is required");
+ }
+ const auto &source = descriptor->source;
+ const auto &frame = source->handle();
+ if (frame.handle == nullptr) {
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): VideoFrame has been released");
+ }
+
+#if defined(__APPLE__)
+ // 1. Import the IOSurface as SharedTextureMemory. For NV12 surfaces this
+ // yields a biplanar texture; for BGRA, a single-plane one.
+ wgpu::SharedTextureMemoryDescriptor memDesc{};
+ std::string label = descriptor->label.value_or("external-texture");
+ if (!label.empty()) {
+ memDesc.label = wgpu::StringView(label.c_str(), label.size());
+ }
+ wgpu::SharedTextureMemoryIOSurfaceDescriptor platformDesc{};
+ platformDesc.ioSurface = frame.handle;
+ // ExternalTexture views are sampled-only; storage binding isn't needed and
+ // for biplanar formats it would fail validation.
+ platformDesc.allowStorageBinding = false;
+ memDesc.nextInChain = &platformDesc;
+ auto memory = device.ImportSharedTextureMemory(&memDesc);
+ if (memory == nullptr) {
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): ImportSharedTextureMemory "
+ "returned null. Is 'shared-texture-memory-iosurface' enabled?");
+ }
+
+ // 2. Create the texture from the surface. We pass the right format
+ // explicitly so Dawn picks the multi-planar variant on NV12.
+ bool isYuv = frame.pixelFormat == VideoPixelFormat::NV12;
+ auto texture = memory.CreateTexture();
+ if (texture == nullptr) {
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): CreateTexture returned null");
+ }
+
+ // 3. Begin access on the underlying memory. The matching EndAccess runs when
+ // the GPUExternalTexture is destroyed (explicitly via destroy() or at GC).
+ wgpu::SharedTextureMemoryBeginAccessDescriptor begin{};
+ begin.initialized = true;
+ begin.concurrentRead = false;
+ if (!memory.BeginAccess(texture, &begin)) {
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): BeginAccess failed");
+ }
+
+ // 4. Build plane views. For NV12 we need plane0 = R8 luma and plane1 = RG8
+ // chroma; for BGRA we only set plane0.
+ wgpu::TextureView plane0;
+ wgpu::TextureView plane1;
+ {
+ wgpu::TextureViewDescriptor v{};
+ v.aspect =
+ isYuv ? wgpu::TextureAspect::Plane0Only : wgpu::TextureAspect::All;
+ plane0 = texture.CreateView(&v);
+ }
+ if (isYuv) {
+ wgpu::TextureViewDescriptor v{};
+ v.aspect = wgpu::TextureAspect::Plane1Only;
+ plane1 = texture.CreateView(&v);
+ }
+
+ // 5. Build the ExternalTextureDescriptor. We hand Dawn explicit YUV→RGB and
+ // sRGB transfer-function parameters so the sampler does the full color
+ // conversion in hardware.
+ wgpu::ExternalTextureDescriptor extDesc{};
+ if (!label.empty()) {
+ extDesc.label = wgpu::StringView(label.c_str(), label.size());
+ }
+ extDesc.plane0 = plane0;
+ extDesc.gamutConversionMatrix = kIdentityGamutMatrix;
+ if (isYuv) {
+ extDesc.plane1 = plane1;
+ extDesc.yuvToRgbConversionMatrix = frame.yuvToRgbMatrix;
+ extDesc.srcTransferFunctionParameters = kSrgbDecodeParams;
+ extDesc.dstTransferFunctionParameters = kSrgbEncodeParams;
+ } else {
+ // BGRA is already RGB in the target color space; pass it through. Dawn
+ // dereferences these arrays unconditionally, so they must be non-null.
+ extDesc.srcTransferFunctionParameters = kIdentityTransferParams;
+ extDesc.dstTransferFunctionParameters = kIdentityTransferParams;
+ }
+ extDesc.cropOrigin = {0, 0};
+ extDesc.cropSize = {frame.width, frame.height};
+ extDesc.apparentSize = {frame.width, frame.height};
+ extDesc.mirrored = descriptor->mirrored.value_or(false);
+ extDesc.rotation =
+ toExternalTextureRotation(descriptor->rotation.value_or(0));
+
+ auto external = device.CreateExternalTexture(&extDesc);
+ if (external == nullptr) {
+ wgpu::SharedTextureMemoryEndAccessState state{};
+ (void)memory.EndAccess(texture, &state);
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): CreateExternalTexture returned "
+ "null");
+ }
+
+ return std::make_shared(
+ std::move(external), std::move(memory), std::move(texture),
+ std::move(descriptor->source), std::move(label));
+#elif defined(__ANDROID__)
+ // 1. Import the AHardwareBuffer as SharedTextureMemory. For YUV AHBs this
+ // yields a Dawn texture in the implementation-defined OpaqueYCbCrAndroid
+ // format; for RGBA AHBs, a regular single-plane texture.
+ wgpu::SharedTextureMemoryDescriptor memDesc{};
+ std::string label = descriptor->label.value_or("external-texture");
+ if (!label.empty()) {
+ memDesc.label = wgpu::StringView(label.c_str(), label.size());
+ }
+ wgpu::SharedTextureMemoryAHardwareBufferDescriptor platformDesc{};
+ platformDesc.handle = frame.handle;
+ memDesc.nextInChain = &platformDesc;
+ auto memory = device.ImportSharedTextureMemory(&memDesc);
+ if (memory == nullptr) {
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): ImportSharedTextureMemory "
+ "returned null. Is 'shared-texture-memory-ahardware-buffer' enabled?");
+ }
+
+ // 2. Create the texture. No descriptor: Dawn picks the right format
+ // (OpaqueYCbCrAndroid for YUV, R8 / RGBA8 / ... for color AHBs).
+ auto texture = memory.CreateTexture();
+ if (texture == nullptr) {
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): CreateTexture returned null");
+ }
+
+ // 3. Begin access. Vulkan requires us to advertise the incoming VkImage
+ // layout (UNDEFINED is fine for the first acquisition of an AHB whose
+ // contents we expect Dawn to read as-is).
+ wgpu::SharedTextureMemoryBeginAccessDescriptor begin{};
+ begin.initialized = true;
+ begin.concurrentRead = false;
+ wgpu::SharedTextureMemoryVkImageLayoutBeginState beginLayout{};
+ begin.nextInChain = &beginLayout;
+ if (!memory.BeginAccess(texture, &begin)) {
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): BeginAccess failed");
+ }
+
+ // 4. Build the ExternalTextureDescriptor. There are two cases depending on
+ // how Dawn imported the AHB (see SharedTextureMemoryVk.cpp):
+ //
+ // a. *External* format (camera buffers whose layout has no Vulkan
+ // equivalent) -> OpaqueYCbCrAndroid, a single opaque plane. Sampling
+ // routes through a Vulkan SamplerYcbcrConversion whose model Dawn
+ // copies verbatim from the AHB's suggestedYcbcrModel. We pass a single
+ // plane + identity transfer and let that conversion (if any) run. NOTE:
+ // when the driver reports RGB_IDENTITY the sample comes back as raw
+ // Y/Cb/Cr; there is no public hook to override the model on this path.
+ //
+ // b. *Defined* biplanar format (e.g. R8BG8Biplanar420Unorm, exposed by the
+ // dawn-multi-planar-formats feature) -> we split Plane0Only (luma) /
+ // Plane1Only (chroma) and hand Dawn an explicit BT.709 matrix + sRGB
+ // transfer, exactly like the iOS NV12 path. This makes numPlanes == 2
+ // so the matrix is actually applied (the single-plane branch in Dawn's
+ // Tint transform ignores yuvToRgbConversionMatrix).
+ //
+ // Either way we must pass non-null gamut/transfer arrays:
+ // ComputeExternalTextureParams dereferences them unconditionally
+ // (kIdentityTransferParams is defined at file scope).
+ const bool isBiplanar = frame.pixelFormat == VideoPixelFormat::NV12 &&
+ isBiplanarYuvFormat(texture.GetFormat());
+
+ wgpu::TextureView plane0;
+ wgpu::TextureView plane1;
+ wgpu::ExternalTextureDescriptor extDesc{};
+ if (!label.empty()) {
+ extDesc.label = wgpu::StringView(label.c_str(), label.size());
+ }
+ extDesc.cropOrigin = {0, 0};
+ extDesc.cropSize = {frame.width, frame.height};
+ extDesc.apparentSize = {frame.width, frame.height};
+ extDesc.gamutConversionMatrix = kIdentityGamutMatrix;
+ if (isBiplanar) {
+ wgpu::TextureViewDescriptor v0{};
+ v0.aspect = wgpu::TextureAspect::Plane0Only;
+ plane0 = texture.CreateView(&v0);
+ wgpu::TextureViewDescriptor v1{};
+ v1.aspect = wgpu::TextureAspect::Plane1Only;
+ plane1 = texture.CreateView(&v1);
+ extDesc.plane0 = plane0;
+ extDesc.plane1 = plane1;
+ extDesc.yuvToRgbConversionMatrix = kBT709LimitedToRgb;
+ extDesc.srcTransferFunctionParameters = kSrgbDecodeParams;
+ extDesc.dstTransferFunctionParameters = kSrgbEncodeParams;
+ } else {
+ plane0 = texture.CreateView();
+ extDesc.plane0 = plane0;
+ extDesc.srcTransferFunctionParameters = kIdentityTransferParams;
+ extDesc.dstTransferFunctionParameters = kIdentityTransferParams;
+ }
+ extDesc.mirrored = descriptor->mirrored.value_or(false);
+ extDesc.rotation =
+ toExternalTextureRotation(descriptor->rotation.value_or(0));
+
+ auto external = device.CreateExternalTexture(&extDesc);
+ if (external == nullptr) {
+ wgpu::SharedTextureMemoryEndAccessState state{};
+ (void)memory.EndAccess(texture, &state);
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): CreateExternalTexture returned "
+ "null");
+ }
+
+ return std::make_shared(
+ std::move(external), std::move(memory), std::move(texture),
+ std::move(descriptor->source), std::move(label));
+#else
+ throw std::runtime_error(
+ "GPUExternalTexture::Create(): not yet implemented on this "
+ "platform");
+#endif
+}
+
+} // namespace rnwgpu
diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUExternalTexture.h b/packages/webgpu/cpp/rnwgpu/api/GPUExternalTexture.h
index 9be5efe6f..217ab45d9 100644
--- a/packages/webgpu/cpp/rnwgpu/api/GPUExternalTexture.h
+++ b/packages/webgpu/cpp/rnwgpu/api/GPUExternalTexture.h
@@ -1,10 +1,13 @@
#pragma once
+#include
#include
+#include
#include "Unions.h"
#include "NativeObject.h"
+#include "VideoFrame.h"
#include "webgpu/webgpu_cpp.h"
@@ -12,16 +15,54 @@ namespace rnwgpu {
namespace jsi = facebook::jsi;
+struct GPUExternalTextureDescriptor;
+
class GPUExternalTexture : public NativeObject {
public:
static constexpr const char *CLASS_NAME = "GPUExternalTexture";
- explicit GPUExternalTexture(wgpu::ExternalTexture instance, std::string label)
- : NativeObject(CLASS_NAME), _instance(instance), _label(label) {}
+ // Import a VideoFrame (via descriptor.source) as a GPUExternalTexture on
+ // `device`: imports the native surface as SharedTextureMemory, begins access,
+ // and wraps the resulting wgpu::ExternalTexture together with the resources
+ // whose lifetime it owns. The matching EndAccess runs in destroy() / the
+ // destructor. Defined in GPUExternalTexture.cpp.
+ static std::shared_ptr
+ Create(wgpu::Device device,
+ std::shared_ptr descriptor);
+
+ // Construct from an already-built wgpu::ExternalTexture plus the underlying
+ // shared-memory resources we need to keep alive. The wrapper takes ownership
+ // of the SharedTextureMemory + Texture and calls EndAccess on destruction so
+ // the producer (e.g. AVPlayer) can reclaim the IOSurface.
+ GPUExternalTexture(wgpu::ExternalTexture instance,
+ wgpu::SharedTextureMemory memory, wgpu::Texture texture,
+ std::shared_ptr source, std::string label)
+ : NativeObject(CLASS_NAME), _instance(std::move(instance)),
+ _memory(std::move(memory)), _texture(std::move(texture)),
+ _source(std::move(source)), _label(std::move(label)) {}
+
+ ~GPUExternalTexture() override { destroy(); }
public:
std::string getBrand() { return CLASS_NAME; }
+ // End the shared-memory access window and release the underlying resources.
+ // Idempotent: safe to call more than once, and the destructor calls it as a
+ // garbage-collection fallback. Call it right after the queue.submit() that
+ // sampled this texture (never before): a GPUExternalTexture's access window
+ // is owned by this wrapper's lifetime, not by submit, so without an explicit
+ // destroy() the producer's surface (e.g. an AVPlayer IOSurface) stays claimed
+ // until GC runs. EndAccess is the designed post-submit call: Dawn keeps the
+ // texture alive for in-flight GPU work via the fences it returns.
+ void destroy() {
+ if (_memory && _texture) {
+ wgpu::SharedTextureMemoryEndAccessState state{};
+ (void)_memory.EndAccess(_texture, &state);
+ }
+ _texture = nullptr;
+ _memory = nullptr;
+ }
+
std::string getLabel() { return _label; }
void setLabel(const std::string &label) {
_label = label;
@@ -33,12 +74,16 @@ class GPUExternalTexture : public NativeObject {
installGetterSetter(runtime, prototype, "label",
&GPUExternalTexture::getLabel,
&GPUExternalTexture::setLabel);
+ installMethod(runtime, prototype, "destroy", &GPUExternalTexture::destroy);
}
inline const wgpu::ExternalTexture get() { return _instance; }
private:
wgpu::ExternalTexture _instance;
+ wgpu::SharedTextureMemory _memory;
+ wgpu::Texture _texture;
+ std::shared_ptr _source;
std::string _label;
};
diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUFeatures.h b/packages/webgpu/cpp/rnwgpu/api/GPUFeatures.h
index c574355ac..a4faa3d98 100644
--- a/packages/webgpu/cpp/rnwgpu/api/GPUFeatures.h
+++ b/packages/webgpu/cpp/rnwgpu/api/GPUFeatures.h
@@ -188,6 +188,9 @@ static void convertEnumToJSUnion(wgpu::FeatureName inEnum,
case wgpu::FeatureName::YCbCrVulkanSamplers:
*outUnion = "ycbcr-vulkan-samplers";
break;
+ case wgpu::FeatureName::OpaqueYCbCrAndroidForExternalTexture:
+ *outUnion = "opaque-ycbcr-android-for-external-texture";
+ break;
case wgpu::FeatureName::ShaderModuleCompilationOptions:
*outUnion = "shader-module-compilation-options";
break;
diff --git a/packages/webgpu/cpp/rnwgpu/api/RNWebGPU.h b/packages/webgpu/cpp/rnwgpu/api/RNWebGPU.h
index c323c0300..4d90b19b5 100644
--- a/packages/webgpu/cpp/rnwgpu/api/RNWebGPU.h
+++ b/packages/webgpu/cpp/rnwgpu/api/RNWebGPU.h
@@ -170,19 +170,29 @@ class RNWebGPU : public NativeObject {
std::shared_ptr loadVideoFrame(std::string path) {
auto frame = _platformContext->loadVideoFrame(path);
- return std::make_shared(frame.handle, frame.width, frame.height,
- std::move(frame.deleter));
+ return std::make_shared(std::move(frame));
}
std::shared_ptr createTestVideoFrame(double width, double height) {
auto frame = _platformContext->createTestVideoFrame(
static_cast(width), static_cast(height));
- return std::make_shared(frame.handle, frame.width, frame.height,
- std::move(frame.deleter));
+ return std::make_shared(std::move(frame));
}
- std::shared_ptr createVideoPlayer(std::string path) {
- auto impl = _platformContext->createVideoPlayer(path);
+ // Wrap a CVPixelBufferRef / AHardwareBuffer* pointer (typed as void* via
+ // BigInt on the JS side) into one of our VideoFrames. The native side
+ // CFRetains / acquires so the caller can release immediately.
+ std::shared_ptr createVideoFrameFromNativeBuffer(void *pointer) {
+ auto handle = _platformContext->wrapNativeBuffer(pointer);
+ return std::make_shared(std::move(handle));
+ }
+
+ std::shared_ptr
+ createVideoPlayer(std::string path, std::optional pixelFormat) {
+ auto format = (pixelFormat && pixelFormat.value() == "nv12")
+ ? VideoPixelFormat::NV12
+ : VideoPixelFormat::BGRA8;
+ auto impl = _platformContext->createVideoPlayer(path, format);
return std::make_shared(std::move(impl));
}
@@ -214,6 +224,8 @@ class RNWebGPU : public NativeObject {
&RNWebGPU::loadVideoFrame);
installMethod(runtime, prototype, "createTestVideoFrame",
&RNWebGPU::createTestVideoFrame);
+ installMethod(runtime, prototype, "createVideoFrameFromNativeBuffer",
+ &RNWebGPU::createVideoFrameFromNativeBuffer);
installMethod(runtime, prototype, "createVideoPlayer",
&RNWebGPU::createVideoPlayer);
installMethod(runtime, prototype, "writeTestVideoFile",
diff --git a/packages/webgpu/cpp/rnwgpu/api/RnFeatures.h b/packages/webgpu/cpp/rnwgpu/api/RnFeatures.h
index 4c60fba23..4577ace66 100644
--- a/packages/webgpu/cpp/rnwgpu/api/RnFeatures.h
+++ b/packages/webgpu/cpp/rnwgpu/api/RnFeatures.h
@@ -12,14 +12,14 @@ namespace rnwgpu {
// platform-specific pair of Dawn features. The prefix is intentional: this
// string is not part of the WebGPU spec, it is our API surface for the
// "import a native surface as a sampleable texture" capability.
-inline constexpr const char *kRnSharedTextureMemoryFeature =
- "rnwebgpu/shared-texture-memory";
+inline constexpr const char *kRnNativeTextureFeature =
+ "rnwebgpu/native-texture";
// Dawn features that back the umbrella on the current platform. Empty on
// platforms where the capability is not available, in which case the umbrella
// behaves as a no-op (it won't appear in adapter.features and asking for it
// in requiredFeatures expands to nothing).
-inline std::vector rnSharedTextureMemoryBackingFeatures() {
+inline std::vector rnNativeTextureBackingFeatures() {
#if defined(__APPLE__)
return {wgpu::FeatureName::SharedTextureMemoryIOSurface,
wgpu::FeatureName::SharedFenceMTLSharedEvent};
@@ -35,10 +35,10 @@ inline std::vector rnSharedTextureMemoryBackingFeatures() {
// umbrella name to `out`. Used by adapter.features / device.features so JS
// callers can see (and call .has on) the same name they pass in.
inline void
-maybeSynthesizeRnSharedTextureMemoryFeature(
+maybeSynthesizeRnNativeTextureFeature(
const std::unordered_set &enabled,
std::unordered_set &out) {
- auto backing = rnSharedTextureMemoryBackingFeatures();
+ auto backing = rnNativeTextureBackingFeatures();
if (backing.empty()) {
return;
}
@@ -47,7 +47,7 @@ maybeSynthesizeRnSharedTextureMemoryFeature(
return;
}
}
- out.insert(kRnSharedTextureMemoryFeature);
+ out.insert(kRnNativeTextureFeature);
}
} // namespace rnwgpu
diff --git a/packages/webgpu/cpp/rnwgpu/api/VideoFrame.h b/packages/webgpu/cpp/rnwgpu/api/VideoFrame.h
index af09cd127..9446e83ae 100644
--- a/packages/webgpu/cpp/rnwgpu/api/VideoFrame.h
+++ b/packages/webgpu/cpp/rnwgpu/api/VideoFrame.h
@@ -7,6 +7,7 @@
#include "JSIConverter.h"
#include "NativeObject.h"
+#include "PlatformContext.h"
namespace rnwgpu {
@@ -19,12 +20,17 @@ namespace jsi = facebook::jsi;
// object stays alive until the JS object is GC'd (or release() is called).
class VideoFrame : public NativeObject {
public:
- static constexpr const char *CLASS_NAME = "VideoFrame";
+ // JS-facing brand. installConstructor() exposes this as globalThis[CLASS_NAME]
+ // and keys the NativeObjectRegistry / worklet serializer off it, so it must
+ // NOT be "VideoFrame": that would install our constructor over the WebCodecs
+ // `VideoFrame` global and shadow it at runtime. "NativeVideoFrame" matches the
+ // public TypeScript type (see src/types.ts) and keeps the two namespaces
+ // distinct. The C++ class stays `VideoFrame` for brevity; only the brand is
+ // namespaced.
+ static constexpr const char *CLASS_NAME = "NativeVideoFrame";
- VideoFrame(void *handle, uint32_t width, uint32_t height,
- std::function deleter)
- : NativeObject(CLASS_NAME), _handle(handle), _width(width),
- _height(height), _deleter(std::move(deleter)) {}
+ explicit VideoFrame(VideoFrameHandle handle)
+ : NativeObject(CLASS_NAME), _handle(std::move(handle)) {}
~VideoFrame() override { release(); }
@@ -32,16 +38,25 @@ class VideoFrame : public NativeObject {
// The native handle (IOSurfaceRef / AHardwareBuffer*) as a uintptr_t value.
// Exposed as a BigInt on the JS side.
- void *getHandle() { return _handle; }
- uint32_t getWidth() { return _width; }
- uint32_t getHeight() { return _height; }
+ void *getHandle() { return _handle.handle; }
+ uint32_t getWidth() { return _handle.width; }
+ uint32_t getHeight() { return _handle.height; }
+
+ // Pixel format as a JS-visible string: "bgra8" | "nv12".
+ std::string getPixelFormat() {
+ return _handle.pixelFormat == VideoPixelFormat::NV12 ? "nv12" : "bgra8";
+ }
+
+ // Direct access to the underlying handle, for use by importExternalTexture /
+ // importSharedTextureMemory inside the C++ layer (not exposed to JS).
+ const VideoFrameHandle &handle() const { return _handle; }
void release() {
- if (_deleter) {
- _deleter();
- _deleter = nullptr;
+ if (_handle.deleter) {
+ _handle.deleter();
+ _handle.deleter = nullptr;
}
- _handle = nullptr;
+ _handle.handle = nullptr;
}
static void definePrototype(jsi::Runtime &runtime, jsi::Object &prototype) {
@@ -49,14 +64,13 @@ class VideoFrame : public NativeObject {
installGetter(runtime, prototype, "handle", &VideoFrame::getHandle);
installGetter(runtime, prototype, "width", &VideoFrame::getWidth);
installGetter(runtime, prototype, "height", &VideoFrame::getHeight);
+ installGetter(runtime, prototype, "pixelFormat",
+ &VideoFrame::getPixelFormat);
installMethod(runtime, prototype, "release", &VideoFrame::release);
}
private:
- void *_handle;
- uint32_t _width;
- uint32_t _height;
- std::function _deleter;
+ VideoFrameHandle _handle;
};
} // namespace rnwgpu
diff --git a/packages/webgpu/cpp/rnwgpu/api/VideoPlayer.h b/packages/webgpu/cpp/rnwgpu/api/VideoPlayer.h
index ee8c2b7af..b97552aca 100644
--- a/packages/webgpu/cpp/rnwgpu/api/VideoPlayer.h
+++ b/packages/webgpu/cpp/rnwgpu/api/VideoPlayer.h
@@ -36,9 +36,7 @@ class VideoPlayer : public NativeObject {
if (handle.handle == nullptr) {
return nullptr;
}
- return std::make_shared(handle.handle, handle.width,
- handle.height,
- std::move(handle.deleter));
+ return std::make_shared(std::move(handle));
}
void play() {
diff --git a/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUBindGroupEntry.h b/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUBindGroupEntry.h
index 73e54dc43..af6ebfda1 100644
--- a/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUBindGroupEntry.h
+++ b/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUBindGroupEntry.h
@@ -21,7 +21,7 @@ struct GPUBindGroupEntry {
std::shared_ptr sampler = nullptr;
std::shared_ptr textureView = nullptr;
std::shared_ptr buffer = nullptr;
- // external textures
+ std::shared_ptr externalTexture = nullptr;
};
} // namespace rnwgpu
@@ -46,6 +46,9 @@ template <> struct JSIConverter> {
} else if (obj.hasNativeState(runtime)) {
result->textureView =
obj.getNativeState(runtime);
+ } else if (obj.hasNativeState(runtime)) {
+ result->externalTexture =
+ obj.getNativeState(runtime);
} else if (obj.hasNativeState(runtime)) {
// Support passing GPUBuffer directly as resource (auto-wrap in
// GPUBufferBinding)
diff --git a/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUDeviceDescriptor.h b/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUDeviceDescriptor.h
index 7e531807c..e18f73e16 100644
--- a/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUDeviceDescriptor.h
+++ b/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUDeviceDescriptor.h
@@ -50,8 +50,8 @@ template <> struct JSIConverter> {
auto str = elementValue.asString(runtime).utf8(runtime);
// Expand react-native-wgpu's umbrella feature into the platform's
// backing Dawn features before they reach RequestDevice.
- if (str == kRnSharedTextureMemoryFeature) {
- for (auto f : rnSharedTextureMemoryBackingFeatures()) {
+ if (str == kRnNativeTextureFeature) {
+ for (auto f : rnNativeTextureBackingFeatures()) {
vector.emplace_back(f);
}
continue;
diff --git a/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUExternalTextureDescriptor.h b/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUExternalTextureDescriptor.h
index fba6721ee..0cb5ecebc 100644
--- a/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUExternalTextureDescriptor.h
+++ b/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUExternalTextureDescriptor.h
@@ -3,35 +3,34 @@
#include
#include
#include
-#include
#include "webgpu/webgpu_cpp.h"
-#include "Convertors.h"
-
#include "JSIConverter.h"
-#include "WGPULogger.h"
+#include "VideoFrame.h"
namespace jsi = facebook::jsi;
namespace rnwgpu {
+// Mirror of GPUExternalTextureDescriptor from the WebGPU spec, but with our
+// VideoFrame as the (only) supported source. We don't expose colorSpace yet;
+// the C++ side picks dst-sRGB and identity gamut, which is the right default
+// for "render this video frame to a regular sRGB framebuffer".
+//
+// `rotation` / `mirrored` are a non-spec extension: camera frames (e.g. from
+// VisionCamera) arrive in the sensor's native orientation, which differs
+// between iOS (CVPixelBuffer) and Android (AHardwareBuffer). Dawn's
+// ExternalTextureDescriptor can bake a rotation + horizontal mirror into the
+// sampling transform, so the shader sees an upright frame without any extra
+// passes. `rotation` is in degrees and must be one of 0 / 90 / 180 / 270.
struct GPUExternalTextureDescriptor {
- // std::variant,
- // std::shared_ptr>
- // source; // | HTMLVideoElement | VideoFrame
- // std::optional colorSpace; // PredefinedColorSpace
- std::optional label; // string
+ std::shared_ptr source;
+ std::optional label;
+ std::optional rotation;
+ std::optional mirrored;
};
-static bool conv(wgpu::ExternalTextureDescriptor &out,
- const std::shared_ptr &in) {
- // TODO: implement
- // return conv(out.source, in->source) && conv(out.colorSpace, in->colorSpace)
- // &&
- // return conv(out.label, in->label);
- return false;
-}
} // namespace rnwgpu
namespace rnwgpu {
@@ -40,23 +39,15 @@ template <>
struct JSIConverter> {
static std::shared_ptr
fromJSI(jsi::Runtime &runtime, const jsi::Value &arg, bool outOfBounds) {
- auto result = std::make_unique();
+ auto result = std::make_shared();
if (!outOfBounds && arg.isObject()) {
auto value = arg.getObject(runtime);
if (value.hasProperty(runtime, "source")) {
auto prop = value.getProperty(runtime, "source");
- // result->source = JSIConverter<
- // std::variant,
- // std::shared_ptr>>::fromJSI(runtime,
- // prop,
- // false);
- }
- if (value.hasProperty(runtime, "colorSpace")) {
- auto prop = value.getProperty(runtime, "colorSpace");
- if (!prop.isUndefined()) {
- // result->colorSpace =
- // JSIConverter>::fromJSI(
- // runtime, prop, false);
+ if (!prop.isUndefined() && !prop.isNull()) {
+ result->source =
+ JSIConverter>::fromJSI(
+ runtime, prop, false);
}
}
if (value.hasProperty(runtime, "label")) {
@@ -66,13 +57,24 @@ struct JSIConverter> {
runtime, prop, false);
}
}
+ if (value.hasProperty(runtime, "rotation")) {
+ auto prop = value.getProperty(runtime, "rotation");
+ if (prop.isNumber()) {
+ result->rotation = prop.asNumber();
+ }
+ }
+ if (value.hasProperty(runtime, "mirrored")) {
+ auto prop = value.getProperty(runtime, "mirrored");
+ if (prop.isBool()) {
+ result->mirrored = prop.getBool();
+ }
+ }
}
-
return result;
}
static jsi::Value
- toJSI(jsi::Runtime &runtime,
- std::shared_ptr arg) {
+ toJSI(jsi::Runtime & /*runtime*/,
+ std::shared_ptr /*arg*/) {
throw std::runtime_error("Invalid GPUExternalTextureDescriptor::toJSI()");
}
};
diff --git a/packages/webgpu/cpp/rnwgpu/api/descriptors/Unions.h b/packages/webgpu/cpp/rnwgpu/api/descriptors/Unions.h
index fef82e9c3..bf3958295 100644
--- a/packages/webgpu/cpp/rnwgpu/api/descriptors/Unions.h
+++ b/packages/webgpu/cpp/rnwgpu/api/descriptors/Unions.h
@@ -527,6 +527,8 @@ inline void convertJSUnionToEnum(const std::string &inUnion,
*outEnum = wgpu::FeatureName::StaticSamplers;
} else if (inUnion == "ycbcr-vulkan-samplers") {
*outEnum = wgpu::FeatureName::YCbCrVulkanSamplers;
+ } else if (inUnion == "opaque-ycbcr-android-for-external-texture") {
+ *outEnum = wgpu::FeatureName::OpaqueYCbCrAndroidForExternalTexture;
} else if (inUnion == "shader-module-compilation-options") {
*outEnum = wgpu::FeatureName::ShaderModuleCompilationOptions;
} else if (inUnion == "dawn-load-resolve-texture") {
@@ -718,6 +720,9 @@ inline void convertEnumToJSUnion(wgpu::FeatureName inEnum,
case wgpu::FeatureName::YCbCrVulkanSamplers:
*outUnion = "ycbcr-vulkan-samplers";
break;
+ case wgpu::FeatureName::OpaqueYCbCrAndroidForExternalTexture:
+ *outUnion = "opaque-ycbcr-android-for-external-texture";
+ break;
case wgpu::FeatureName::ShaderModuleCompilationOptions:
*outUnion = "shader-module-compilation-options";
break;
diff --git a/packages/webgpu/package.json b/packages/webgpu/package.json
index 961528fb3..07a48b53e 100644
--- a/packages/webgpu/package.json
+++ b/packages/webgpu/package.json
@@ -1,6 +1,6 @@
{
"name": "react-native-wgpu",
- "version": "0.5.12",
+ "version": "0.5.13",
"description": "React Native WebGPU",
"main": "lib/commonjs/index",
"module": "lib/module/index",
diff --git a/packages/webgpu/src/__tests__/ExternalTexture.spec.ts b/packages/webgpu/src/__tests__/ExternalTexture.spec.ts
index 47976ec4b..712fec9bb 100644
--- a/packages/webgpu/src/__tests__/ExternalTexture.spec.ts
+++ b/packages/webgpu/src/__tests__/ExternalTexture.spec.ts
@@ -1,5 +1,17 @@
import { checkImage, client, encodeImage } from "./setup";
+interface BitmapData {
+ data: number[];
+ width: number;
+ height: number;
+ format: string;
+}
+
+type EvalResult =
+ | { kind: "skip"; reason: string }
+ | { kind: "fail"; reason: string }
+ | ({ kind: "ok" } & BitmapData);
+
describe("External Textures", () => {
it("Simple (1)", async () => {
const result = await client.eval(
@@ -405,4 +417,152 @@ describe("External Textures", () => {
// This test catches the bug where std::optional was checked incorrectly
checkImage(image, "snapshots/f2.png");
});
+
+ // Regression test for createBindGroupLayout with a `texture_external` entry.
+ // Every other path uses layout: "auto", which never builds an explicit
+ // BindGroupLayoutEntry for an external texture, so the native conversion that
+ // chains ExternalTextureBindingLayout went untested. Here we build the layout
+ // ourselves (externalTexture: {}) and render a GPUExternalTexture through it.
+ // Device-only: needs a native frame, so it skips on the web reference.
+ it("samples a GPUExternalTexture through an explicit bind group layout", async () => {
+ const result = await client.eval, EvalResult>(
+ ({ device, gpu, ctx, canvas }) => {
+ const FEATURE = "rnwebgpu/native-texture";
+ if (!device.features.has(FEATURE as GPUFeatureName)) {
+ return {
+ kind: "skip",
+ reason: `${FEATURE} not enabled on this device`,
+ };
+ }
+ if (typeof RNWebGPU?.createTestVideoFrame !== "function") {
+ return {
+ kind: "skip",
+ reason: "RNWebGPU.createTestVideoFrame is unavailable",
+ };
+ }
+
+ try {
+ const frame = RNWebGPU.createTestVideoFrame(256, 256);
+ const externalTexture = device.importExternalTexture({
+ // createTestVideoFrame returns our NativeVideoFrame; the native
+ // binding accepts it, but the spec type wants a WebCodecs
+ // VideoFrame, so cast to satisfy the signature.
+ source: frame as unknown as VideoFrame,
+ label: "test-frame",
+ });
+
+ const module = device.createShaderModule({
+ code: /* wgsl */ `
+ struct VsOut {
+ @builtin(position) position: vec4f,
+ @location(0) uv: vec2f,
+ };
+
+ @vertex fn vs(@builtin(vertex_index) vid: u32) -> VsOut {
+ var positions = array(
+ vec2f(-1.0, -3.0),
+ vec2f(-1.0, 1.0),
+ vec2f( 3.0, 1.0),
+ );
+ var uvs = array(
+ vec2f(0.0, 2.0),
+ vec2f(0.0, 0.0),
+ vec2f(2.0, 0.0),
+ );
+ var out: VsOut;
+ out.position = vec4f(positions[vid], 0.0, 1.0);
+ out.uv = uvs[vid];
+ return out;
+ }
+
+ @group(0) @binding(0) var srcTex: texture_external;
+ @group(0) @binding(1) var srcSampler: sampler;
+
+ @fragment fn fs(in: VsOut) -> @location(0) vec4f {
+ return textureSampleBaseClampToEdge(srcTex, srcSampler, in.uv);
+ }
+ `,
+ });
+ // The whole point of the test: an explicit external-texture layout.
+ const bindGroupLayout = device.createBindGroupLayout({
+ entries: [
+ {
+ binding: 0,
+ visibility: GPUShaderStage.FRAGMENT,
+ externalTexture: {},
+ },
+ { binding: 1, visibility: GPUShaderStage.FRAGMENT, sampler: {} },
+ ],
+ });
+ const pipeline = device.createRenderPipeline({
+ layout: device.createPipelineLayout({
+ bindGroupLayouts: [bindGroupLayout],
+ }),
+ vertex: { module, entryPoint: "vs" },
+ fragment: {
+ module,
+ entryPoint: "fs",
+ targets: [{ format: gpu.getPreferredCanvasFormat() }],
+ },
+ primitive: { topology: "triangle-list" },
+ });
+ const sampler = device.createSampler({
+ magFilter: "linear",
+ minFilter: "linear",
+ });
+ const bindGroup = device.createBindGroup({
+ layout: bindGroupLayout,
+ entries: [
+ { binding: 0, resource: externalTexture },
+ { binding: 1, resource: sampler },
+ ],
+ });
+
+ const encoder = device.createCommandEncoder();
+ const pass = encoder.beginRenderPass({
+ colorAttachments: [
+ {
+ view: ctx.getCurrentTexture().createView(),
+ clearValue: { r: 0, g: 0, b: 0, a: 1 },
+ loadOp: "clear",
+ storeOp: "store",
+ },
+ ],
+ });
+ pass.setPipeline(pipeline);
+ pass.setBindGroup(0, bindGroup);
+ pass.draw(3);
+ pass.end();
+ device.queue.submit([encoder.finish()]);
+ // End the external texture's shared-memory access window now that the
+ // work sampling it is submitted, rather than waiting for GC.
+ externalTexture.destroy();
+
+ return canvas.getImageData().then((image: BitmapData) => {
+ frame.release();
+ return { kind: "ok" as const, ...image };
+ });
+ } catch (e) {
+ return {
+ kind: "fail",
+ reason: `${(e as Error).message ?? e}`,
+ };
+ }
+ },
+ );
+
+ if (result.kind === "skip") {
+ console.log(
+ `ExternalTexture (explicit layout): skipping (${result.reason})`,
+ );
+ return;
+ }
+ if (result.kind === "fail") {
+ throw new Error(`ExternalTexture (explicit layout): ${result.reason}`);
+ }
+ const image = encodeImage(result);
+ // Same render as the auto-layout path in ImportExternalTexture.spec.ts, so
+ // it must match that snapshot bit-for-bit.
+ checkImage(image, "snapshots/import-external-texture.png");
+ });
});
diff --git a/packages/webgpu/src/__tests__/ImportExternalTexture.spec.ts b/packages/webgpu/src/__tests__/ImportExternalTexture.spec.ts
new file mode 100644
index 000000000..376d26d10
--- /dev/null
+++ b/packages/webgpu/src/__tests__/ImportExternalTexture.spec.ts
@@ -0,0 +1,149 @@
+import { checkImage, client, encodeImage } from "./setup";
+
+interface BitmapData {
+ data: number[];
+ width: number;
+ height: number;
+ format: string;
+}
+
+type EvalResult =
+ | { kind: "skip"; reason: string }
+ | { kind: "fail"; reason: string }
+ | ({ kind: "ok" } & BitmapData);
+
+describe("ImportExternalTexture", () => {
+ it("imports a test frame as an external texture and samples it", async () => {
+ const result = await client.eval, EvalResult>(
+ ({ device, gpu, ctx, canvas }) => {
+ // The umbrella feature backs importExternalTexture's IOSurface /
+ // AHardwareBuffer import. The Chrome reference run and fallback adapters
+ // won't advertise it, so that's the *only* legitimate skip. Anything
+ // past this point is a real failure.
+ const FEATURE = "rnwebgpu/native-texture";
+ if (!device.features.has(FEATURE as GPUFeatureName)) {
+ return {
+ kind: "skip",
+ reason: `${FEATURE} not enabled on this device`,
+ };
+ }
+ if (typeof RNWebGPU?.createTestVideoFrame !== "function") {
+ return {
+ kind: "skip",
+ reason: "RNWebGPU.createTestVideoFrame is unavailable",
+ };
+ }
+
+ try {
+ const frame = RNWebGPU.createTestVideoFrame(256, 256);
+ // Unlike importSharedTextureMemory, there is no createTexture /
+ // beginAccess / endAccess to manage: the GPUExternalTexture owns the
+ // shared-memory access window and keeps the frame alive internally.
+ const externalTexture = device.importExternalTexture({
+ // createTestVideoFrame returns our NativeVideoFrame; the native
+ // importExternalTexture binding accepts it, but the spec type wants
+ // a WebCodecs VideoFrame, so cast to satisfy the signature.
+ source: frame as unknown as VideoFrame,
+ label: "test-frame",
+ });
+
+ const module = device.createShaderModule({
+ code: /* wgsl */ `
+ struct VsOut {
+ @builtin(position) position: vec4f,
+ @location(0) uv: vec2f,
+ };
+
+ @vertex fn vs(@builtin(vertex_index) vid: u32) -> VsOut {
+ var positions = array(
+ vec2f(-1.0, -3.0),
+ vec2f(-1.0, 1.0),
+ vec2f( 3.0, 1.0),
+ );
+ var uvs = array(
+ vec2f(0.0, 2.0),
+ vec2f(0.0, 0.0),
+ vec2f(2.0, 0.0),
+ );
+ var out: VsOut;
+ out.position = vec4f(positions[vid], 0.0, 1.0);
+ out.uv = uvs[vid];
+ return out;
+ }
+
+ @group(0) @binding(0) var srcTex: texture_external;
+ @group(0) @binding(1) var srcSampler: sampler;
+
+ @fragment fn fs(in: VsOut) -> @location(0) vec4f {
+ return textureSampleBaseClampToEdge(srcTex, srcSampler, in.uv);
+ }
+ `,
+ });
+ const pipeline = device.createRenderPipeline({
+ layout: "auto",
+ vertex: { module, entryPoint: "vs" },
+ fragment: {
+ module,
+ entryPoint: "fs",
+ targets: [{ format: gpu.getPreferredCanvasFormat() }],
+ },
+ primitive: { topology: "triangle-list" },
+ });
+ const sampler = device.createSampler({
+ magFilter: "linear",
+ minFilter: "linear",
+ });
+ const bindGroup = device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: 0, resource: externalTexture },
+ { binding: 1, resource: sampler },
+ ],
+ });
+
+ const encoder = device.createCommandEncoder();
+ const pass = encoder.beginRenderPass({
+ colorAttachments: [
+ {
+ view: ctx.getCurrentTexture().createView(),
+ clearValue: { r: 0, g: 0, b: 0, a: 1 },
+ loadOp: "clear",
+ storeOp: "store",
+ },
+ ],
+ });
+ pass.setPipeline(pipeline);
+ pass.setBindGroup(0, bindGroup);
+ pass.draw(3);
+ pass.end();
+ device.queue.submit([encoder.finish()]);
+ // End the external texture's shared-memory access window now that the
+ // work sampling it is submitted, rather than waiting for GC.
+ externalTexture.destroy();
+
+ return canvas.getImageData().then((image: BitmapData) => {
+ // Safe to release now: all GPU work referencing the frame is
+ // submitted and the pixels have been read back.
+ frame.release();
+ return { kind: "ok" as const, ...image };
+ });
+ } catch (e) {
+ return {
+ kind: "fail",
+ reason: `${(e as Error).message ?? e}`,
+ };
+ }
+ },
+ );
+
+ if (result.kind === "skip") {
+ console.log(`ImportExternalTexture: skipping (${result.reason})`);
+ return;
+ }
+ if (result.kind === "fail") {
+ throw new Error(`ImportExternalTexture: ${result.reason}`);
+ }
+ const image = encodeImage(result);
+ checkImage(image, "snapshots/import-external-texture.png");
+ });
+});
diff --git a/packages/webgpu/src/__tests__/SharedTextureMemory.spec.ts b/packages/webgpu/src/__tests__/SharedTextureMemory.spec.ts
index 2e1e3fad4..325f95932 100644
--- a/packages/webgpu/src/__tests__/SharedTextureMemory.spec.ts
+++ b/packages/webgpu/src/__tests__/SharedTextureMemory.spec.ts
@@ -20,7 +20,7 @@ describe("SharedTextureMemory", () => {
// reference run and fallback adapters won't advertise it. Anything
// else past this point is a real failure and must surface as a test
// failure, not a silent skip.
- const FEATURE = "rnwebgpu/shared-texture-memory";
+ const FEATURE = "rnwebgpu/native-texture";
if (!device.features.has(FEATURE as GPUFeatureName)) {
return {
kind: "skip",
@@ -145,4 +145,147 @@ describe("SharedTextureMemory", () => {
const image = encodeImage(result);
checkImage(image, "snapshots/shared-texture-memory.png");
});
+
+ // Same as above, but with an *explicit* bind group layout
+ // (device.createBindGroupLayout + createPipelineLayout) instead of
+ // layout: "auto". This exercises the native BindGroupLayoutEntry conversion
+ // path, which "auto" layouts bypass entirely.
+ it("samples a shared texture through an explicit bind group layout", async () => {
+ const result = await client.eval, EvalResult>(
+ ({ device, gpu, ctx, canvas }) => {
+ const FEATURE = "rnwebgpu/native-texture";
+ if (!device.features.has(FEATURE as GPUFeatureName)) {
+ return {
+ kind: "skip",
+ reason: `${FEATURE} not enabled on this device`,
+ };
+ }
+ if (typeof RNWebGPU?.createTestVideoFrame !== "function") {
+ return {
+ kind: "skip",
+ reason: "RNWebGPU.createTestVideoFrame is unavailable",
+ };
+ }
+
+ try {
+ const frame = RNWebGPU.createTestVideoFrame(256, 256);
+ const memory = device.importSharedTextureMemory({
+ handle: frame.handle,
+ label: "test-frame",
+ });
+ const texture = memory.createTexture();
+ if (!memory.beginAccess(texture, true)) {
+ frame.release();
+ return { kind: "fail", reason: `beginAccess returned false` };
+ }
+
+ const module = device.createShaderModule({
+ code: /* wgsl */ `
+ struct VsOut {
+ @builtin(position) position: vec4f,
+ @location(0) uv: vec2f,
+ };
+
+ @vertex fn vs(@builtin(vertex_index) vid: u32) -> VsOut {
+ var positions = array(
+ vec2f(-1.0, -3.0),
+ vec2f(-1.0, 1.0),
+ vec2f( 3.0, 1.0),
+ );
+ var uvs = array(
+ vec2f(0.0, 2.0),
+ vec2f(0.0, 0.0),
+ vec2f(2.0, 0.0),
+ );
+ var out: VsOut;
+ out.position = vec4f(positions[vid], 0.0, 1.0);
+ out.uv = uvs[vid];
+ return out;
+ }
+
+ @group(0) @binding(0) var srcTex: texture_2d;
+ @group(0) @binding(1) var srcSampler: sampler;
+
+ @fragment fn fs(in: VsOut) -> @location(0) vec4f {
+ return textureSample(srcTex, srcSampler, in.uv);
+ }
+ `,
+ });
+ const bindGroupLayout = device.createBindGroupLayout({
+ entries: [
+ { binding: 0, visibility: GPUShaderStage.FRAGMENT, texture: {} },
+ { binding: 1, visibility: GPUShaderStage.FRAGMENT, sampler: {} },
+ ],
+ });
+ const pipeline = device.createRenderPipeline({
+ layout: device.createPipelineLayout({
+ bindGroupLayouts: [bindGroupLayout],
+ }),
+ vertex: { module, entryPoint: "vs" },
+ fragment: {
+ module,
+ entryPoint: "fs",
+ targets: [{ format: gpu.getPreferredCanvasFormat() }],
+ },
+ primitive: { topology: "triangle-list" },
+ });
+ const sampler = device.createSampler({
+ magFilter: "linear",
+ minFilter: "linear",
+ });
+ const bindGroup = device.createBindGroup({
+ layout: bindGroupLayout,
+ entries: [
+ { binding: 0, resource: texture.createView() },
+ { binding: 1, resource: sampler },
+ ],
+ });
+
+ const encoder = device.createCommandEncoder();
+ const pass = encoder.beginRenderPass({
+ colorAttachments: [
+ {
+ view: ctx.getCurrentTexture().createView(),
+ clearValue: { r: 0, g: 0, b: 0, a: 1 },
+ loadOp: "clear",
+ storeOp: "store",
+ },
+ ],
+ });
+ pass.setPipeline(pipeline);
+ pass.setBindGroup(0, bindGroup);
+ pass.draw(3);
+ pass.end();
+ device.queue.submit([encoder.finish()]);
+
+ return canvas.getImageData().then((image: BitmapData) => {
+ memory.endAccess(texture);
+ texture.destroy();
+ frame.release();
+ return { kind: "ok" as const, ...image };
+ });
+ } catch (e) {
+ return {
+ kind: "fail",
+ reason: `${(e as Error).message ?? e}`,
+ };
+ }
+ },
+ );
+
+ if (result.kind === "skip") {
+ console.log(
+ `SharedTextureMemory (explicit layout): skipping (${result.reason})`,
+ );
+ return;
+ }
+ if (result.kind === "fail") {
+ throw new Error(
+ `SharedTextureMemory (explicit layout): ${result.reason}`,
+ );
+ }
+ const image = encodeImage(result);
+ // Identical render to the auto-layout case above.
+ checkImage(image, "snapshots/shared-texture-memory.png");
+ });
});
diff --git a/packages/webgpu/src/__tests__/snapshots/import-external-texture-web.png b/packages/webgpu/src/__tests__/snapshots/import-external-texture-web.png
new file mode 100644
index 000000000..a15e25e3a
Binary files /dev/null and b/packages/webgpu/src/__tests__/snapshots/import-external-texture-web.png differ
diff --git a/packages/webgpu/src/__tests__/snapshots/import-external-texture.png b/packages/webgpu/src/__tests__/snapshots/import-external-texture.png
new file mode 100644
index 000000000..a15e25e3a
Binary files /dev/null and b/packages/webgpu/src/__tests__/snapshots/import-external-texture.png differ
diff --git a/packages/webgpu/src/index.tsx b/packages/webgpu/src/index.tsx
index 4672dfc8b..6ac631a48 100644
--- a/packages/webgpu/src/index.tsx
+++ b/packages/webgpu/src/index.tsx
@@ -5,13 +5,16 @@ import type {
NativeCanvas,
RNCanvasContext,
VideoPlayer,
- VideoFrame,
+ NativeVideoFrame,
+ NativeVideoPixelFormat,
} from "./types";
export * from "./main";
export type {
- VideoFrame,
+ NativeVideoFrame,
VideoPlayer,
+ NativeVideoPixelFormat,
+ CreateVideoPlayerOptions,
GPUSharedTextureMemory,
GPUSharedTextureMemoryDescriptor,
} from "./types";
@@ -34,9 +37,16 @@ declare global {
) => RNCanvasContext;
DecodeToUTF8: (buffer: NodeJS.ArrayBufferView | ArrayBuffer) => string;
createImageBitmap: typeof createImageBitmap;
- loadVideoFrame: (path: string) => VideoFrame;
- createTestVideoFrame: (width: number, height: number) => VideoFrame;
- createVideoPlayer: (path: string) => VideoPlayer;
+ loadVideoFrame: (path: string) => NativeVideoFrame;
+ createTestVideoFrame: (width: number, height: number) => NativeVideoFrame;
+ // Wrap a NativeBuffer.pointer (CVPixelBufferRef on iOS / AHardwareBuffer*
+ // on Android) into a NativeVideoFrame. Matches the shape used by libraries
+ // that emit NativeBuffer (e.g. react-native-vision-camera).
+ createVideoFrameFromNativeBuffer: (pointer: bigint) => NativeVideoFrame;
+ createVideoPlayer: (
+ path: string,
+ pixelFormat?: NativeVideoPixelFormat,
+ ) => VideoPlayer;
writeTestVideoFile: () => string;
};
@@ -46,6 +56,27 @@ declare global {
): GPUSharedTextureMemory;
}
+ // Non-spec extension: camera frames arrive in the sensor's native
+ // orientation, which differs between iOS and Android. `rotation` (degrees,
+ // one of 0/90/180/270) and `mirrored` (horizontal flip) are baked into the
+ // sampling transform by Dawn, so the shader sees an upright frame. Maps
+ // directly onto VisionCamera's `frame.orientation` / `frame.isMirrored`.
+ interface GPUExternalTextureDescriptor {
+ rotation?: 0 | 90 | 180 | 270;
+ mirrored?: boolean;
+ }
+
+ // Non-spec extension: a GPUExternalTexture imported from a native surface
+ // holds an open shared-memory access window on that surface until this
+ // wrapper is destroyed. Call destroy() right after the queue.submit() that
+ // sampled it (never before) to release the surface back to its producer
+ // immediately, instead of waiting for garbage collection. Forgetting to call
+ // it is not fatal (GC still cleans up), but it can starve a producer's buffer
+ // pool (e.g. a camera/video player) and pile up GPU resources.
+ interface GPUExternalTexture {
+ destroy(): void;
+ }
+
// Extend createImageBitmap to accept ArrayBuffer/TypedArray (encoded image bytes)
function createImageBitmap(
image: ArrayBuffer | ArrayBufferView,
diff --git a/packages/webgpu/src/types.ts b/packages/webgpu/src/types.ts
index b015747c5..a982763a5 100644
--- a/packages/webgpu/src/types.ts
+++ b/packages/webgpu/src/types.ts
@@ -18,18 +18,32 @@ export interface CanvasRef {
getNativeSurface: () => NativeCanvas;
}
+// Pixel layout of a NativeVideoFrame. NOT the WebCodecs `VideoPixelFormat`
+// enum — these are the two native surface layouts we support, lower-cased to
+// avoid being mistaken for the spec values ("NV12", "BGRA", …).
+export type NativeVideoPixelFormat = "bgra8" | "nv12";
+
// A native, GPU-shareable handle to a single video frame.
//
+// NOT the WebCodecs `VideoFrame`: there is no `close()`/`format`/`timestamp`,
+// the surface is referenced by a raw native pointer, and disposal is
+// `release()`. Named with a `Native` prefix so it doesn't shadow the global
+// WebCodecs type or imply spec semantics it doesn't have.
+//
// - handle is the raw pointer (IOSurfaceRef on Apple, AHardwareBuffer* on
// Android) encoded as a BigInt. Pass it to
// GPUDevice.importSharedTextureMemory.
+// - pixelFormat describes the surface layout: 'bgra8' for a sampled
+// GPUTexture; 'nv12' (biplanar Y + CbCr) for the importExternalTexture
+// path.
// - release() drops the underlying backing object (a CVPixelBuffer on Apple).
// The frame is also released when the JS wrapper is garbage-collected; call
// release() eagerly when you know you're done.
-export interface VideoFrame {
+export interface NativeVideoFrame {
readonly handle: bigint;
readonly width: number;
readonly height: number;
+ readonly pixelFormat: NativeVideoPixelFormat;
release(): void;
}
@@ -37,17 +51,26 @@ export interface VideoFrame {
// to obtain the most recently decoded frame as an IOSurface/AHardwareBuffer
// (returns null between frames so callers can skip the import work).
export interface VideoPlayer {
- copyLatestFrame(): VideoFrame | null;
+ copyLatestFrame(): NativeVideoFrame | null;
play(): void;
pause(): void;
release(): void;
}
+export interface CreateVideoPlayerOptions {
+ // 'bgra8' (default): emit a single-plane BGRA surface, suitable for
+ // SharedTextureMemory and a regular sampled GPUTexture.
+ // 'nv12': emit biplanar Y + CbCr surfaces, suitable for
+ // GPUDevice.importExternalTexture.
+ pixelFormat?: NativeVideoPixelFormat;
+}
+
export interface GPUSharedTextureMemoryDescriptor {
// Raw native handle (IOSurfaceRef on Apple, AHardwareBuffer* on Android),
// encoded as a BigInt. The caller is responsible for keeping the underlying
// object alive for as long as the shared memory (and any textures derived
- // from it) are in use. Using VideoFrame.handle handles this automatically.
+ // from it) are in use. Using NativeVideoFrame.handle handles this
+ // automatically.
handle: bigint;
label?: string;
}