diff --git a/apps/example/src/App.tsx b/apps/example/src/App.tsx index 354a1c572..917429ce0 100644 --- a/apps/example/src/App.tsx +++ b/apps/example/src/App.tsx @@ -37,7 +37,9 @@ import { AsyncStarvation } from "./Diagnostics/AsyncStarvation"; import { DeviceLostHang } from "./Diagnostics/DeviceLostHang"; import { StorageBufferVertices } from "./StorageBufferVertices"; import { SharedTextureMemory } from "./SharedTextureMemory"; -import { Camera } from "./Camera"; +import { ExternalTexture } from "./ExternalTexture"; +import { VisionCamera } from "./VisionCamera"; +import { ChromeSphere } from "./ChromeSphere"; // The two lines below are needed by three.js import "fast-text-encoding"; @@ -48,9 +50,9 @@ const Stack = createStackNavigator(); function App() { const assets = useAssets(); - if (assets === null) { - return null; - } + // if (assets === null) { + // return null; + // } return ( @@ -89,7 +91,7 @@ function App() { - {(props) => } + {(props) => (assets ? : null)} @@ -103,7 +105,9 @@ 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/ChromeSphere/ChromeSphere.tsx b/apps/example/src/ChromeSphere/ChromeSphere.tsx new file mode 100644 index 000000000..287962d6f --- /dev/null +++ b/apps/example/src/ChromeSphere/ChromeSphere.tsx @@ -0,0 +1,980 @@ +/* eslint-disable prefer-destructuring */ +import React, { useEffect } 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 "../VisionCamera/blurShaders"; + +import { + generateSphere, + NORMAL_OFFSET, + POSITION_OFFSET, + VERTEX_STRIDE_BYTES, +} from "./geometry"; +import { BACKDROP_SHADER, SHADER } from "./shader"; + +// All matrix math runs inside the Vision Camera frame-processor worklet, so it +// has to be implemented with worklet-friendly helpers. wgpu-matrix calls +// `mat4.create` / `vec3.create` internally and those aren't marked as +// worklets, so we inline what we need here. Conventions match wgpu-matrix: +// column-major 4x4 stored as Float32Array(16), index = col * 4 + row, +// right-handed view space, perspective maps z to [0, 1] (WebGPU clip space). + +const setIdentity = (out: Float32Array) => { + "worklet"; + out[0] = 1; + out[1] = 0; + out[2] = 0; + out[3] = 0; + out[4] = 0; + out[5] = 1; + out[6] = 0; + out[7] = 0; + out[8] = 0; + out[9] = 0; + out[10] = 1; + out[11] = 0; + out[12] = 0; + out[13] = 0; + out[14] = 0; + out[15] = 1; +}; + +// dst = m * Ry(angle). Safe with dst === m. +const applyRotateY = (m: Float32Array, angle: number, dst: Float32Array) => { + "worklet"; + const c = Math.cos(angle); + const s = Math.sin(angle); + const m00 = m[0], + m10 = m[1], + m20 = m[2], + m30 = m[3]; + const m02 = m[8], + m12 = m[9], + m22 = m[10], + m32 = m[11]; + if (dst !== m) { + dst[4] = m[4]; + dst[5] = m[5]; + dst[6] = m[6]; + dst[7] = m[7]; + dst[12] = m[12]; + dst[13] = m[13]; + dst[14] = m[14]; + dst[15] = m[15]; + } + dst[0] = c * m00 - s * m02; + dst[1] = c * m10 - s * m12; + dst[2] = c * m20 - s * m22; + dst[3] = c * m30 - s * m32; + dst[8] = s * m00 + c * m02; + dst[9] = s * m10 + c * m12; + dst[10] = s * m20 + c * m22; + dst[11] = s * m30 + c * m32; +}; + +// WebGPU-style perspective: right-handed, output z mapped to [0, 1]. fovy is +// the vertical field of view in radians. +const setPerspective = ( + out: Float32Array, + fovy: number, + aspect: number, + near: number, + far: number, +) => { + "worklet"; + const f = 1 / Math.tan(fovy * 0.5); + const rangeInv = 1 / (near - far); + out[0] = f / aspect; + out[1] = 0; + out[2] = 0; + out[3] = 0; + out[4] = 0; + out[5] = f; + out[6] = 0; + out[7] = 0; + out[8] = 0; + out[9] = 0; + out[10] = far * rangeInv; + out[11] = -1; + out[12] = 0; + out[13] = 0; + out[14] = far * near * rangeInv; + out[15] = 0; +}; + +// View matrix (RH). Camera looks down -z in its local frame. Same layout +// wgpu-matrix's lookAt produces, with the translation column carrying the +// negative basis-dot-eye terms. +const setLookAt = ( + out: Float32Array, + ex: number, + ey: number, + ez: number, + tx: number, + ty: number, + tz: number, + ux: number, + uy: number, + uz: number, +) => { + "worklet"; + let zx = ex - tx; + let zy = ey - ty; + let zz = ez - tz; + const zLen = Math.hypot(zx, zy, zz); + zx /= zLen; + zy /= zLen; + zz /= zLen; + + let xx = uy * zz - uz * zy; + let xy = uz * zx - ux * zz; + let xz = ux * zy - uy * zx; + const xLen = Math.hypot(xx, xy, xz); + xx /= xLen; + xy /= xLen; + xz /= xLen; + + const yx = zy * xz - zz * xy; + const yy = zz * xx - zx * xz; + const yz = zx * xy - zy * xx; + + out[0] = xx; + out[1] = yx; + out[2] = zx; + out[3] = 0; + out[4] = xy; + out[5] = yy; + out[6] = zy; + out[7] = 0; + out[8] = xz; + out[9] = yz; + out[10] = zz; + out[11] = 0; + out[12] = -(xx * ex + xy * ey + xz * ez); + out[13] = -(yx * ex + yy * ey + yz * ez); + out[14] = -(zx * ex + zy * ey + zz * ez); + out[15] = 1; +}; + +// dst = a * b. Safe with dst === a or dst === b (snapshots both fully first). +const multiplyMat4 = (a: Float32Array, b: Float32Array, dst: Float32Array) => { + "worklet"; + const a00 = a[0], + a10 = a[1], + a20 = a[2], + a30 = a[3]; + const a01 = a[4], + a11 = a[5], + a21 = a[6], + a31 = a[7]; + const a02 = a[8], + a12 = a[9], + a22 = a[10], + a32 = a[11]; + const a03 = a[12], + a13 = a[13], + a23 = a[14], + a33 = a[15]; + const b00 = b[0], + b10 = b[1], + b20 = b[2], + b30 = b[3]; + const b01 = b[4], + b11 = b[5], + b21 = b[6], + b31 = b[7]; + const b02 = b[8], + b12 = b[9], + b22 = b[10], + b32 = b[11]; + const b03 = b[12], + b13 = b[13], + b23 = b[14], + b33 = b[15]; + + dst[0] = a00 * b00 + a01 * b10 + a02 * b20 + a03 * b30; + dst[1] = a10 * b00 + a11 * b10 + a12 * b20 + a13 * b30; + dst[2] = a20 * b00 + a21 * b10 + a22 * b20 + a23 * b30; + dst[3] = a30 * b00 + a31 * b10 + a32 * b20 + a33 * b30; + dst[4] = a00 * b01 + a01 * b11 + a02 * b21 + a03 * b31; + dst[5] = a10 * b01 + a11 * b11 + a12 * b21 + a13 * b31; + dst[6] = a20 * b01 + a21 * b11 + a22 * b21 + a23 * b31; + dst[7] = a30 * b01 + a31 * b11 + a32 * b21 + a33 * b31; + dst[8] = a00 * b02 + a01 * b12 + a02 * b22 + a03 * b32; + dst[9] = a10 * b02 + a11 * b12 + a12 * b22 + a13 * b32; + dst[10] = a20 * b02 + a21 * b12 + a22 * b22 + a23 * b32; + dst[11] = a30 * b02 + a31 * b12 + a32 * b22 + a33 * b32; + dst[12] = a00 * b03 + a01 * b13 + a02 * b23 + a03 * b33; + dst[13] = a10 * b03 + a11 * b13 + a12 * b23 + a13 * b33; + dst[14] = a20 * b03 + a21 * b13 + a22 * b23 + a23 * b33; + dst[15] = a30 * b03 + a31 * b13 + a32 * b23 + a33 * b33; +}; + +// The 3D variant of the VisionCamera demo. Reuses the same shared-texture-memory +// pipeline to import camera frames as GPUExternalTextures, but instead of +// applying 2D effects, it samples the camera as a spherical environment map on +// a chrome sphere, an orbiting cube, and a torus. + +const REQUIRED_FEATURES: GPUFeatureName[] = [ + "rnwebgpu/shared-texture-memory" as GPUFeatureName, + "dawn-multi-planar-formats" as GPUFeatureName, +]; + +// Android-only feature; same probe as VisionCamera.tsx. Without it Dawn can't +// wrap a YCbCr AHB as a GPUExternalTexture. +const OPAQUE_YCBCR_EXT = + "opaque-ycbcr-android-for-external-texture" as GPUFeatureName; + +const DEPTH_FORMAT: GPUTextureFormat = "depth24plus"; + +// Backdrop blur tuning. Matches VisionCamera's "Strong" preset: prepass to a +// 1/4-res rgba8unorm, then 3 H-V iterations of the tile-based box blur. The +// final result is sampled by BACKDROP_SHADER as a fullscreen backdrop. +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; +const BLUR_ITERATIONS = 3; + +// Scene UBO layout. mat4(64) + vec4(16) + vec4(16) = 96 bytes; pad to 16-byte +// alignment for safety. +const SCENE_UBO_SIZE = 96; +const SCENE_UBO_FLOATS = SCENE_UBO_SIZE / 4; +const OBJECT_UBO_SIZE = 64; // mat4 + +type Shape = { + vertexBuffer: GPUBuffer; + indexBuffer: GPUBuffer; + indexCount: number; + uniformBuffer: GPUBuffer; + // Returns a model matrix for time t (seconds). + modelAt: (t: number, out: Float32Array) => void; +}; + +export const ChromeSphere = () => { + 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 SceneView = () => { + const ref = useCanvasRef(); + const [gpu, setGpu] = React.useState<{ + adapter: GPUAdapter; + device: GPUDevice; + } | null>(null); + const [deviceError, setDeviceError] = React.useState(null); + React.useEffect(() => { + let cancelled = false; + (async () => { + try { + const adapter = await navigator.gpu.requestAdapter(); + if (!adapter) { + throw new Error("requestAdapter returned null"); + } + 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.", + ); + } + const featuresToRequest: GPUFeatureName[] = [ + ...REQUIRED_FEATURES, + ...(Platform.OS === "android" ? [OPAQUE_YCBCR_EXT] : []), + ]; + const device = await adapter.requestDevice({ + requiredFeatures: featuresToRequest, + }); + if (cancelled) { + return; + } + setGpu({ adapter, device }); + } catch (e) { + if (cancelled) { + return; + } + setDeviceError(String(e)); + } + })(); + return () => { + cancelled = true; + }; + }, []); + const device = gpu?.device ?? null; + const adapter = gpu?.adapter ?? null; + const devices = useCameraDevices(); + const cameraDevice = React.useMemo( + () => + devices.find((d) => d.position === "back") ?? + devices.find((d) => d.position === "front") ?? + devices[0], + [devices], + ); + + const [pipelineState, setPipelineState] = React.useState<{ + pipeline: GPURenderPipeline; + sampler: GPUSampler; + sceneUniformBuffer: GPUBuffer; + depthView: GPUTextureView; + context: RNCanvasContext; + canvasWidth: number; + canvasHeight: number; + shapes: Shape[]; + // Pre-allocated scratch so the worklet doesn't allocate per frame. These + // get mutated each tick, which is safe because the worklet sees its own + // copy after closure serialization. + view: Float32Array; + proj: Float32Array; + viewProj: Float32Array; + sceneData: Float32Array; + modelScratch: Float32Array; + // Blurred-camera backdrop infrastructure. + backdropPipeline: GPURenderPipeline; + backdropBindGroup: GPUBindGroup; + prepassPipeline: GPURenderPipeline; + prepassUniformBuffer: GPUBuffer; + blurPipeline: GPUComputePipeline; + blurConstants: GPUBindGroup; + blurBindGroup0: GPUBindGroup; + blurBindGroup1: GPUBindGroup; + blurBindGroup2: GPUBindGroup; + blurSrcTexture: GPUTexture; + blurWidth: number; + blurHeight: number; + } | null>(null); + const [error, setError] = React.useState(null); + + 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: SHADER }); + const pipeline = device.createRenderPipeline({ + layout: "auto", + vertex: { + module, + entryPoint: "vs_main", + buffers: [ + { + arrayStride: VERTEX_STRIDE_BYTES, + attributes: [ + { + shaderLocation: 0, + offset: POSITION_OFFSET, + format: "float32x3", + }, + { + shaderLocation: 1, + offset: NORMAL_OFFSET, + format: "float32x3", + }, + ], + }, + ], + }, + fragment: { + module, + entryPoint: "fs_main", + targets: [{ format: presentationFormat }], + }, + primitive: { topology: "triangle-list", cullMode: "back" }, + depthStencil: { + depthCompare: "less", + depthWriteEnabled: true, + format: DEPTH_FORMAT, + }, + }); + const sampler = device.createSampler({ + magFilter: "linear", + minFilter: "linear", + }); + const sceneUniformBuffer = device.createBuffer({ + size: SCENE_UBO_SIZE, + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + }); + const depthTexture = device.createTexture({ + size: [canvas.width, canvas.height], + format: DEPTH_FORMAT, + usage: GPUTextureUsage.RENDER_ATTACHMENT, + }); + + const sphereMesh = generateSphere(1.0, 48, 64); + + const buildShape = ( + mesh: { vertices: Float32Array; indices: Uint16Array }, + modelAt: (t: number, out: Float32Array) => void, + ): Shape => { + const vertexBuffer = device.createBuffer({ + size: mesh.vertices.byteLength, + usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_DST, + }); + device.queue.writeBuffer(vertexBuffer, 0, mesh.vertices); + // Index buffers must be sized to a multiple of 4 bytes; pad Uint16Array + // by one extra index if needed. + const indexByteLength = (mesh.indices.byteLength + 3) & ~3; + const indexBuffer = device.createBuffer({ + size: indexByteLength, + usage: GPUBufferUsage.INDEX | GPUBufferUsage.COPY_DST, + }); + device.queue.writeBuffer(indexBuffer, 0, mesh.indices); + const uniformBuffer = device.createBuffer({ + size: OBJECT_UBO_SIZE, + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + }); + // The full bind group is built per-frame in the worklet because + // GPUExternalTexture is single-use; scene + object buffers above stay + // alive across frames and feed back into that per-frame build. + return { + vertexBuffer, + indexBuffer, + indexCount: mesh.indices.length, + uniformBuffer, + modelAt, + }; + }; + + // Single chrome sphere center stage with a slow Y-rotation so the + // reflection drifts even when the orbit camera is between key positions. + // Runs inside the frame worklet. + const shapes: Shape[] = [ + buildShape(sphereMesh, (t, out) => { + "worklet"; + setIdentity(out); + applyRotateY(out, t * 0.25, out); + }), + ]; + + // ----- Backdrop blur infrastructure (same as VisionCamera "Strong") --- + 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: 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(); + + // Backdrop pipeline: shares the render pass with the chrome sphere, so + // it must declare a matching depth-stencil layout. depthCompare always / + // depthWriteEnabled false means it draws unconditionally and never + // disturbs depth for the subsequent sphere draw. + const backdropModule = device.createShaderModule({ code: BACKDROP_SHADER }); + const backdropPipeline = device.createRenderPipeline({ + layout: "auto", + vertex: { module: backdropModule, entryPoint: "vs_main" }, + fragment: { + module: backdropModule, + entryPoint: "fs_main", + targets: [{ format: presentationFormat }], + }, + primitive: { topology: "triangle-list" }, + depthStencil: { + depthCompare: "always", + depthWriteEnabled: false, + format: DEPTH_FORMAT, + }, + }); + const backdropBindGroup = device.createBindGroup({ + layout: backdropPipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: blurredView }, + { binding: 1, resource: sampler }, + ], + }); + + setPipelineState({ + pipeline, + sampler, + sceneUniformBuffer, + depthView: depthTexture.createView(), + context, + canvasWidth: canvas.width, + canvasHeight: canvas.height, + shapes, + view: new Float32Array(16), + proj: new Float32Array(16), + viewProj: new Float32Array(16), + sceneData: new Float32Array(SCENE_UBO_FLOATS), + modelScratch: new Float32Array(16), + backdropPipeline, + backdropBindGroup, + prepassPipeline, + prepassUniformBuffer, + blurPipeline, + blurConstants, + blurBindGroup0, + blurBindGroup1, + blurBindGroup2, + blurSrcTexture, + blurWidth, + blurHeight, + }); + }, [device, adapter, ref, pipelineState]); + + const startTimeRef = React.useRef(performance.now()); + const logBox = React.useMemo(() => ({ seen: false }), []); + + const frameOutput = useFrameOutput({ + pixelFormat: "native", + onFrame: (frame) => { + "worklet"; + if (!logBox.seen) { + logBox.seen = true; + console.log( + "[ChromeSphere] worklet first frame, hasPipeline=" + + String(pipelineState != null) + + " frame=" + + String(frame.width) + + "x" + + String(frame.height), + ); + } + if (!pipelineState || !device) { + frame.dispose(); + return; + } + const { + pipeline, + sampler, + sceneUniformBuffer, + depthView, + context, + canvasWidth, + canvasHeight, + shapes, + view, + proj, + viewProj, + sceneData, + modelScratch, + backdropPipeline, + backdropBindGroup, + prepassPipeline, + prepassUniformBuffer, + blurPipeline, + blurConstants, + blurBindGroup0, + blurBindGroup1, + blurBindGroup2, + blurSrcTexture, + blurWidth, + blurHeight, + } = pipelineState; + const nativeBuffer = frame.getNativeBuffer(); + try { + const videoFrame = device.createVideoFrameFromNativeBuffer( + nativeBuffer.pointer, + ); + try { + const t = (performance.now() - startTimeRef.current) / 1000; + + // Orbit the eye around the origin, looking at it. Small bob in y + // so reflections shift along the polar axis too. + const orbitR = 5.2; + const ex = Math.cos(t * 0.15) * orbitR; + const ey = 1.2 + Math.sin(t * 0.2) * 0.4; + const ez = Math.sin(t * 0.15) * orbitR; + setLookAt(view, ex, ey, ez, 0.0, 0.0, 0.0, 0, 1, 0); + setPerspective( + proj, + Math.PI / 4, + canvasWidth / canvasHeight, + 0.1, + 100, + ); + multiplyMat4(proj, view, viewProj); + + sceneData.set(viewProj, 0); + sceneData[16] = ex; + sceneData[17] = ey; + sceneData[18] = ez; + sceneData[19] = 0; + // Key light rises and slowly drifts so the chrome specular sweeps + // across the silhouettes. + const ldRawX = Math.cos(t * 0.3) * 0.6; + const ldRawY = 0.8; + const ldRawZ = Math.sin(t * 0.3) * 0.6; + const ldLen = Math.hypot(ldRawX, ldRawY, ldRawZ); + sceneData[20] = ldRawX / ldLen; + sceneData[21] = ldRawY / ldLen; + sceneData[22] = ldRawZ / ldLen; + sceneData[23] = 0; + device.queue.writeBuffer(sceneUniformBuffer, 0, sceneData); + + for (const shape of shapes) { + shape.modelAt(t, modelScratch); + device.queue.writeBuffer(shape.uniformBuffer, 0, modelScratch); + } + + let externalTex; + try { + externalTex = device.importExternalTexture({ + source: videoFrame, + label: "chrome-env", + }); + } catch (e) { + console.warn( + "[ChromeSphere] importExternalTexture threw: " + String(e), + ); + throw e; + } + + const encoder = device.createCommandEncoder(); + + // ---- Backdrop blur (prepass + 3 H-V iterations at 1/4 res) ---- + device.queue.writeBuffer( + prepassUniformBuffer, + 0, + new Float32Array([ + videoFrame.width, + videoFrame.height, + 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(); + + 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 < BLUR_ITERATIONS - 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(); + + // ---- Main scene pass: backdrop first (no depth write), then sphere + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: context.getCurrentTexture().createView(), + clearValue: { r: 0, g: 0, b: 0, a: 1 }, + loadOp: "clear", + storeOp: "store", + }, + ], + depthStencilAttachment: { + view: depthView, + depthClearValue: 1.0, + depthLoadOp: "clear", + depthStoreOp: "store", + }, + }); + + // Backdrop pipeline has depthCompare: "always" and writes nothing + // to depth, so the subsequent sphere draw still sees a clean depth + // buffer with the canvas-clear far value. + pass.setPipeline(backdropPipeline); + pass.setBindGroup(0, backdropBindGroup); + pass.draw(3); + + pass.setPipeline(pipeline); + for (const shape of shapes) { + // The external texture is bound per-shape (one bind group per + // shape, the only difference being the per-object uniform), so + // we rebuild the bind group each frame to splice in this frame's + // externalTex alongside the cached scene + object buffers. + const bindGroup = device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: sceneUniformBuffer } }, + { binding: 1, resource: { buffer: shape.uniformBuffer } }, + { binding: 2, resource: externalTex }, + { binding: 3, resource: sampler }, + ], + }); + pass.setBindGroup(0, bindGroup); + pass.setVertexBuffer(0, shape.vertexBuffer); + pass.setIndexBuffer(shape.indexBuffer, "uint16"); + pass.drawIndexed(shape.indexCount); + } + pass.end(); + device.queue.submit([encoder.finish()]); + context.present(); + } finally { + videoFrame.release(); + } + } finally { + nativeBuffer.release(); + frame.dispose(); + } + }, + }); + + 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/ChromeSphere/geometry.ts b/apps/example/src/ChromeSphere/geometry.ts new file mode 100644 index 000000000..1ede29fcd --- /dev/null +++ b/apps/example/src/ChromeSphere/geometry.ts @@ -0,0 +1,50 @@ +// Sphere mesh generator for the chrome scene. Returns interleaved +// position+normal vertices and a triangle index list. The shader uses the +// upper-left 3x3 of the model matrix to transform normals, so the sphere is +// built with unit-length normals (no non-uniform scale at generation time). + +export type Mesh = { + vertices: Float32Array; // [px, py, pz, nx, ny, nz] * N + indices: Uint16Array; +}; + +const FLOATS_PER_VERTEX = 6; + +export const VERTEX_STRIDE_BYTES = FLOATS_PER_VERTEX * 4; +export const POSITION_OFFSET = 0; +export const NORMAL_OFFSET = 12; + +export function generateSphere( + radius: number, + latBands: number, + lonBands: number, +): Mesh { + const verts: number[] = []; + const idx: number[] = []; + for (let lat = 0; lat <= latBands; lat++) { + const theta = (lat * Math.PI) / latBands; + const sinTheta = Math.sin(theta); + const cosTheta = Math.cos(theta); + for (let lon = 0; lon <= lonBands; lon++) { + const phi = (lon * 2 * Math.PI) / lonBands; + const sinPhi = Math.sin(phi); + const cosPhi = Math.cos(phi); + const nx = cosPhi * sinTheta; + const ny = cosTheta; + const nz = sinPhi * sinTheta; + verts.push(nx * radius, ny * radius, nz * radius, nx, ny, nz); + } + } + const stride = lonBands + 1; + for (let lat = 0; lat < latBands; lat++) { + for (let lon = 0; lon < lonBands; lon++) { + const a = lat * stride + lon; + const b = a + stride; + idx.push(a, b, a + 1, b, b + 1, a + 1); + } + } + return { + vertices: new Float32Array(verts), + indices: new Uint16Array(idx), + }; +} diff --git a/apps/example/src/ChromeSphere/index.ts b/apps/example/src/ChromeSphere/index.ts new file mode 100644 index 000000000..5f32bab2e --- /dev/null +++ b/apps/example/src/ChromeSphere/index.ts @@ -0,0 +1 @@ +export * from "./ChromeSphere"; diff --git a/apps/example/src/ChromeSphere/shader.ts b/apps/example/src/ChromeSphere/shader.ts new file mode 100644 index 000000000..425606aaf --- /dev/null +++ b/apps/example/src/ChromeSphere/shader.ts @@ -0,0 +1,137 @@ +// Backdrop shader. Fullscreen triangle that samples the pre-blurred camera +// image (cover-fit baked in by the prepass), dims it ~50%, and adds a soft +// vignette so the chrome sphere stays the focal point of the scene. +export const BACKDROP_SHADER = /* wgsl */ ` +@group(0) @binding(0) var src: texture_2d; +@group(0) @binding(1) var samp: sampler; + +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 { + var c = textureSampleLevel(src, samp, in.uv, 0.0).rgb; + // Subdue the room: half-brightness + soft vignette darkening the edges by + // another ~40%. Keeps a sense of ambient lighting without competing with + // the chrome reflection. + c = c * 0.55; + let d = distance(in.uv, vec2f(0.5)); + let v = 1.0 - smoothstep(0.4, 0.95, d) * 0.45; + return vec4f(c * v, 1.0); +} +`; + +// Chrome-style env-map reflection driven by the live camera feed. +// +// For each fragment we compute the reflection vector around the world-space +// normal, convert it to spherical (theta, phi) → uv, and sample the external +// camera texture. textureSampleBaseClampToEdge is the only sampling call +// allowed on texture_external, but it's all we need: spherical uvs land +// inside [0, 1]^2 so no wrap mode tricks are required. +// +// The visible result is the classic CGI chrome ball: your face wraps around +// the sphere as if the camera image were the ambient environment, and the +// cube + torus reflect distorted copies of the same scene. + +export const SHADER = /* wgsl */ ` +struct Scene { + viewProj: mat4x4f, + cameraPos: vec4f, // xyz = world-space eye, w unused + lightDir: vec4f, // xyz = unit vector toward light, w unused +}; + +struct Object { + model: mat4x4f, +}; + +@group(0) @binding(0) var scene: Scene; +@group(0) @binding(1) var obj: Object; +@group(0) @binding(2) var srcTex: texture_external; +@group(0) @binding(3) var srcSampler: sampler; + +struct VsIn { + @location(0) position: vec3f, + @location(1) normal: vec3f, +}; + +struct VsOut { + @builtin(position) clipPos: vec4f, + @location(0) worldPos: vec3f, + @location(1) worldNormal: vec3f, +}; + +@vertex +fn vs_main(in: VsIn) -> VsOut { + let world = obj.model * vec4f(in.position, 1.0); + // Rotation-only transforms: normal matrix is the upper-left 3x3 of model. + // No non-uniform scale anywhere in the scene, so inverse-transpose collapses + // to this. + let normalMat = mat3x3f( + obj.model[0].xyz, + obj.model[1].xyz, + obj.model[2].xyz, + ); + var out: VsOut; + out.clipPos = scene.viewProj * world; + out.worldPos = world.xyz; + out.worldNormal = normalize(normalMat * in.normal); + return out; +} + +const PI: f32 = 3.14159265359; + +fn sphericalUv(dir: vec3f) -> vec2f { + // dir is treated as a direction from the chrome surface to the + // environment. theta wraps around the y-axis, phi runs pole to pole. + let theta = atan2(dir.z, dir.x); + let phi = acos(clamp(dir.y, -1.0, 1.0)); + return vec2f((theta + PI) / (2.0 * PI), phi / PI); +} + +@fragment +fn fs_main(in: VsOut) -> @location(0) vec4f { + let n = normalize(in.worldNormal); + let v = normalize(scene.cameraPos.xyz - in.worldPos); + // reflect() expects the incident vector (from light/source to surface), so + // negate the view direction. + let r = normalize(reflect(-v, n)); + + let uv = sphericalUv(r); + let env = textureSampleBaseClampToEdge(srcTex, srcSampler, uv).rgb; + + // Schlick-ish Fresnel: chrome reflectance is high everywhere, but boost a + // bit at grazing angles so silhouettes catch the light. + let cosTheta = max(dot(n, v), 0.0); + let fresnel = pow(1.0 - cosTheta, 3.0); + let tint = vec3f(0.96, 0.97, 1.00); // subtle cool cast, very polished chrome + var color = env * tint + vec3f(fresnel) * 0.18; + + // Single directional specular highlight so the chrome doesn't look like a + // pure billboard. Halfway vector with a tight exponent. + let h = normalize(scene.lightDir.xyz + v); + let spec = pow(max(dot(n, h), 0.0), 96.0); + color = color + vec3f(spec) * 1.2; + + return vec4f(color, 1.0); +} +`; diff --git a/apps/example/src/ExternalTexture/ExternalTexture.tsx b/apps/example/src/ExternalTexture/ExternalTexture.tsx new file mode 100644 index 000000000..2b2504b4c --- /dev/null +++ b/apps/example/src/ExternalTexture/ExternalTexture.tsx @@ -0,0 +1,251 @@ +import React, { useEffect, useRef, useState } from "react"; +import { PixelRatio, StyleSheet, Text, View } from "react-native"; +import { + Canvas, + useCanvasRef, + useDevice, + type NativeCanvas, + type NativeVideoFrame, +} from "react-native-wgpu"; + +// importExternalTexture is the spec-mandated path for "I have a YUV-encoded +// video/camera frame and I want to sample it in a shader without copying or +// hand-rolling YUV math". The WGSL side uses texture_external + +// textureSampleBaseClampToEdge; the driver does the planar fetch, YUV→RGB +// matrix multiply, sRGB transfer, and gamut conversion in the sampler. +// +// Bind groups for texture_external use auto layout slots like any other +// resource. WGSL doesn't expose the underlying plane textures directly. +const SHADER = /* wgsl */ ` +struct VsOut { + @builtin(position) position: vec4f, + @location(0) uv: vec2f, +}; + +struct Uniforms { + 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 { + 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); +} +`; + +// rnwebgpu/shared-texture-memory is our umbrella that expands to the +// platform's shared-memory + shared-fence pair (the IOSurface / AHB still +// flows through SharedTextureMemory under the hood). Plus +// dawn-multi-planar-formats so Dawn can interpret the NV12 surface as a +// biplanar texture. +const REQUIRED_FEATURES: GPUFeatureName[] = [ + "rnwebgpu/shared-texture-memory" as GPUFeatureName, + "dawn-multi-planar-formats" as GPUFeatureName, +]; + +const VIDEO_URL = + "https://test-videos.co.uk/vids/bigbuckbunny/mp4/h264/1080/Big_Buck_Bunny_1080_10s_5MB.mp4"; + +export const ExternalTexture = () => { + const ref = useCanvasRef(); + const [error, setError] = useState(null); + const rafRef = useRef(null); + + const { device, adapter } = useDevice(undefined, { + requiredFeatures: REQUIRED_FEATURES, + }); + + useEffect(() => { + if (!device) { + return; + } + const missing = REQUIRED_FEATURES.filter((f) => !device.features.has(f)); + if (missing.length > 0) { + setError( + `Device is missing required features [${missing.join(", ")}]. 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", + }); + + // Same Big Buck Bunny URL as the SharedTextureMemory demo, but ask AVPlayer + // for native NV12 instead of BGRA. Each VideoFrame now carries the YUV + // matrix + plane info that importExternalTexture needs. + const player = RNWebGPU.createVideoPlayer(VIDEO_URL, "nv12"); + player.play(); + + 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; we rewrite its uvScale whenever a new + // frame's dimensions differ from the last one. + const uniformBuffer = device.createBuffer({ + size: 16, + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + }); + let lastUvScale: [number, number] | null = null; + const writeUvScale = (texW: number, texH: number) => { + const canvasAR = canvas.width / canvas.height; + const texAR = texW / texH; + const next: [number, number] = + texAR > canvasAR ? [canvasAR / texAR, 1] : [1, texAR / canvasAR]; + if ( + !lastUvScale || + lastUvScale[0] !== next[0] || + lastUvScale[1] !== next[1] + ) { + device.queue.writeBuffer(uniformBuffer, 0, new Float32Array(next)); + lastUvScale = next; + } + }; + + // The video plays at ~24fps but we tick at the display's 60Hz, so most rAF + // ticks have no new frame from AVPlayer. Hold the latest VideoFrame across + // ticks and re-import an ExternalTexture from it on the "no new frame" + // ticks — this is what stops the canvas from flashing black ~2/3 of the + // time. AVPlayer's pool is several buffers deep so holding one back like + // this doesn't stall decoding. + let currentFrame: NativeVideoFrame | null = null; + + const render = () => { + const newFrame = player.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", + }, + ], + }); + + if (currentFrame) { + // GPUExternalTexture expires after each submit, so we rebuild one + // every tick — even when sampling the same VideoFrame as last tick. + let externalTex: GPUExternalTexture | null = null; + try { + externalTex = device.importExternalTexture({ + source: currentFrame, + label: "video-external", + }); + } catch (e) { + console.warn("[ExternalTexture] importExternalTexture failed:", e); + } + + if (externalTex) { + writeUvScale(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()]); + 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(); + player.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/ExternalTexture/index.ts b/apps/example/src/ExternalTexture/index.ts new file mode 100644 index 000000000..d019021b8 --- /dev/null +++ b/apps/example/src/ExternalTexture/index.ts @@ -0,0 +1 @@ +export * from "./ExternalTexture"; diff --git a/apps/example/src/Home.tsx b/apps/example/src/Home.tsx index 74c153c0f..d43f0822e 100644 --- a/apps/example/src/Home.tsx +++ b/apps/example/src/Home.tsx @@ -132,8 +132,16 @@ export const examples = [ title: "🎞️ Shared Texture Memory", }, { - screen: "Camera", - title: "📷 Camera", + screen: "ExternalTexture", + title: "🟨 External Texture (YUV)", + }, + { + screen: "VisionCamera", + title: "📷 VisionCamera integration", + }, + { + screen: "ChromeSphere", + title: "🪩 Chrome Sphere (camera env map)", }, ]; diff --git a/apps/example/src/Route.ts b/apps/example/src/Route.ts index 5c43b2a36..57f09c992 100644 --- a/apps/example/src/Route.ts +++ b/apps/example/src/Route.ts @@ -30,5 +30,7 @@ export type Routes = { DeviceLostHang: undefined; StorageBufferVertices: undefined; SharedTextureMemory: undefined; - Camera: undefined; + ExternalTexture: undefined; + VisionCamera: undefined; + ChromeSphere: undefined; }; diff --git a/apps/example/src/SharedTextureMemory/SharedTextureMemory.tsx b/apps/example/src/SharedTextureMemory/SharedTextureMemory.tsx index 7805cb7ff..e91c69367 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 */ ` @@ -103,7 +103,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 +117,7 @@ export const SharedTextureMemory = () => { release: () => player.release(), }; } else { - let pending: VideoFrame | null = RNWebGPU.createTestVideoFrame( + let pending: NativeVideoFrame | null = RNWebGPU.createTestVideoFrame( 1024, 1024, ); @@ -168,7 +168,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 +190,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/ThreeJS/CameraHelmet.tsx b/apps/example/src/ThreeJS/CameraHelmet.tsx new file mode 100644 index 000000000..0d48d0c47 --- /dev/null +++ b/apps/example/src/ThreeJS/CameraHelmet.tsx @@ -0,0 +1,711 @@ +import React, { useEffect, useRef, useState } from "react"; +import { + Linking, + Platform, + StyleSheet, + Text, + TouchableOpacity, + View, +} from "react-native"; +import type { CanvasRef } from "react-native-wgpu"; +import { Canvas } from "react-native-wgpu"; +import { + CommonResolutions, + NativePreviewView, + useCameraDevices, + useCameraPermission, + useFrameOutput, + usePreviewOutput, + VisionCamera as VisionCameraFactory, +} from "react-native-vision-camera"; +import type { + CameraController, + CameraSession, +} from "react-native-vision-camera"; +import * as THREE from "three"; + +import { useGLTF } from "./assets/AssetManager"; +import { makeWebGPURenderer } from "./components/makeWebGPURenderer"; +import { CAMERA_ENV_SHADER } from "./cameraEnvShader"; + +// Live camera as a three.js environment map. The GLTF helmet renders with +// three.js' WebGPURenderer; its env map is a THREE.ExternalTexture wrapping a +// GPUTexture we own. A Vision Camera frame-processor worklet writes each +// camera frame into that GPUTexture via its own render pass. Three.js and the +// worklet share a single GPUDevice (the one three.js creates internally), so +// the queue ordering between "write env" and "sample env" is automatic. +// +// The front-camera frame is wrapped around the inside of a large back-side +// sphere centered on the helmet, so the CubeCamera at the helmet's middle +// samples camera content in every direction. The helmet's reflections +// therefore fully cover the helmet and never reveal the frame's edges. The +// trade-off is the obvious "face wrapped around the world" look on grazing +// reflections; that's chosen deliberately, since full coverage matters more +// than the seam artifact for this demo. + +// 9:16 portrait — front cam delivers 16:9 landscape and the shader rotates it +// 90° to selfie-upright, so this aspect lets the rotated frame fill the env +// texture with no stretching. Cube face dimension matches ENV_HEIGHT. Each +// frame we do (env write + 6 cube faces + optional mipmap chain), so this +// knob drives most of the per-frame GPU cost. +const ENV_WIDTH = 540; +const ENV_HEIGHT = 960; + +// Cube face size is mode-dependent. Chrome is a pure mirror reflection, so it +// wants the env's full resolution. PBR samples roughness-selected mips (the +// helmet's metal-roughness keeps most surfaces in the 0.3-0.6 range, i.e. +// mip 2-3), so a 128 cube + its short mip chain is visually indistinguishable +// from 512 while cutting cubemap fill rate ~16x and mip-regen cost with it. +const CUBE_SIZE_PBR = 128; +const CUBE_SIZE_CHROME = 512; + +// PBR mode is runtime-toggleable from a button in the JSX below. PBR uses +// the GLTF's MeshStandardMaterial textures (albedo / normal / metalRoughness +// / AO) plus a cubemap envMap for the live reflection; "chrome" mode swaps +// every mesh for a single MeshBasicMaterial that just samples the cubemap. +// Chrome is roughly 5-10x cheaper on fragments and also skips the per-frame +// cubemap mipmap regen the PBR roughness lookup needs. + +// Vision Camera + react-native-wgpu both want these features for the external +// texture path. dawn-multi-planar-formats lets Dawn interpret NV12 buffers. +const REQUIRED_FEATURES: GPUFeatureName[] = [ + "rnwebgpu/shared-texture-memory" as GPUFeatureName, + "dawn-multi-planar-formats" as GPUFeatureName, +]; + +const OPAQUE_YCBCR_EXT = + "opaque-ycbcr-android-for-external-texture" as GPUFeatureName; + +export const CameraHelmet = () => { + 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 Scene = () => { + useEffect(() => { + console.log("[CameraHelmet] Scene mounted"); + return () => console.log("[CameraHelmet] Scene unmounted"); + }, []); + const ref = useRef(null); + const gltf = useGLTF(require("./assets/helmet/DamagedHelmet.gltf")); + // Live camera preview, rendered as a native view behind the WebGPU canvas. + // The worklet still writes the camera into our env texture for the helmet + // reflection, but the *backdrop* now comes straight from this native + // preview — no detour through equirect/cubemap, no quality loss. + const previewOutput = usePreviewOutput(); + + // Two cameras at once: the back camera feeds the native preview backdrop, + // the front camera feeds the helmet's environment map (so you see yourself + // reflected in the chrome). Requires multi-cam capable hardware (iPhone + // XS+ / most modern Android flagships). + const devices = useCameraDevices(); + const backDevice = React.useMemo( + () => devices.find((d) => d.position === "back"), + [devices], + ); + const frontDevice = React.useMemo( + () => devices.find((d) => d.position === "front"), + [devices], + ); + + const [pipelineState, setPipelineState] = useState<{ + device: GPUDevice; + cameraPipeline: GPURenderPipeline; + cameraSampler: GPUSampler; + envTexture: GPUTexture; + envTextureView: GPUTextureView; + } | null>(null); + const [error, setError] = useState(null); + const [device, setDevice] = useState(null); + + // PBR is the default. Toggling to chrome swaps every helmet material for + // a single MeshBasicMaterial that just samples the cubemap. usePBRRef + // shadows the state so the (one-shot) setup effect can read the *current* + // value when it first applies materials, even if the user has already + // toggled before three.js finished initializing. + const [usePBR, setUsePBR] = useState(true); + const usePBRRef = useRef(true); + const applyPBRFnRef = useRef<((pbr: boolean) => void) | null>(null); + const togglePBR = () => { + const next = !usePBRRef.current; + usePBRRef.current = next; + setUsePBR(next); + applyPBRFnRef.current?.(next); + }; + + // Acquire the GPU device on its own effect. By the time the async adapter + + // device requests resolve, the Canvas component has been rendered and its + // ref populated, so the main setup effect (gated on `device`) can grab the + // GPUCanvasContext synchronously. Same two-effect pattern as + // VisionCamera.tsx / ChromeSphere.tsx. + useEffect(() => { + console.log("[CameraHelmet] device-acquisition effect fired"); + let cancelled = false; + (async () => { + try { + const adapter = await navigator.gpu.requestAdapter(); + if (!adapter) { + throw new Error("requestAdapter returned null"); + } + const requiredFeatures = [...adapter.features] as GPUFeatureName[]; + const missing = REQUIRED_FEATURES.filter( + (f) => !adapter.features.has(f), + ); + const needsAndroidExt = + Platform.OS === "android" && !adapter.features.has(OPAQUE_YCBCR_EXT); + if (missing.length > 0 || needsAndroidExt) { + throw new Error( + "Adapter doesn't advertise the features the Vision Camera " + + "external-texture path needs: " + + `${[...missing, needsAndroidExt ? OPAQUE_YCBCR_EXT : null] + .filter(Boolean) + .join(", ")}.`, + ); + } + const d = await adapter.requestDevice({ requiredFeatures }); + console.log( + "[CameraHelmet] device acquired, features: " + + [...d.features].sort().join(", "), + ); + if (cancelled) { + d.destroy(); + return; + } + setDevice(d); + } catch (e) { + if (cancelled) { + return; + } + console.warn("[CameraHelmet] device acquisition failed: " + String(e)); + setError(String(e)); + } + })(); + return () => { + cancelled = true; + }; + }, []); + + // Note: pipelineState is intentionally not in the deps array. Including it + // would re-run the effect when we call setPipelineState below — React would + // run the cleanup (which calls setAnimationLoop(null)) and then the effect + // would bail on the pipelineState guard, leaving us with no render loop. + // The effect only needs to fire once, when `device` transitions to set. + + useEffect(() => { + console.log( + "[CameraHelmet] setup effect fired, device=" + + String(device != null) + + " gltf=" + + String(gltf != null), + ); + if (!device || !gltf) { + return; + } + const context = ref.current?.getContext("webgpu"); + if (!context) { + console.log( + "[CameraHelmet] no webgpu context yet (ref.current=" + + String(ref.current != null) + + ") — bailing this effect run", + ); + return; + } + let cancelled = false; + let renderer: THREE.WebGPURenderer | null = null; + + console.log("[CameraHelmet] context acquired, building three.js scene"); + (async () => { + try { + const { width, height } = context.canvas; + console.log( + "[CameraHelmet] canvas size = " + + String(width) + + "x" + + String(height), + ); + + // alpha:true configures the canvas with premultiplied alpha mode, so + // pixels outside the helmet stay transparent and the native camera + // preview behind the canvas shows through. + renderer = makeWebGPURenderer(context, { device, alpha: true }); + renderer.toneMapping = THREE.ACESFilmicToneMapping; + renderer.setClearColor(0x000000, 0); + await renderer.init(); + console.log("[CameraHelmet] three.js renderer init complete"); + if (cancelled) { + return; + } + + // Env GPUTexture: render target on our side, sampleable on three's + // side. rgba8unorm + RENDER_ATTACHMENT|TEXTURE_BINDING lets the + // single resource pivot between the two roles via implicit barriers. + const envTexture = device.createTexture({ + size: [ENV_WIDTH, ENV_HEIGHT], + format: "rgba8unorm", + usage: + GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.TEXTURE_BINDING, + }); + + // Camera prepass pipeline. Output format matches the env texture so + // it can be the render target. + const module = device.createShaderModule({ code: CAMERA_ENV_SHADER }); + const cameraPipeline = device.createRenderPipeline({ + layout: "auto", + vertex: { module, entryPoint: "vs_main" }, + fragment: { + module, + entryPoint: "fs_main", + targets: [{ format: "rgba8unorm" }], + }, + primitive: { topology: "triangle-list" }, + }); + const cameraSampler = device.createSampler({ + magFilter: "linear", + minFilter: "linear", + }); + + // THREE.ExternalTexture bridges our GPUTexture into three.js as a + // sampleable 2D texture. Used below as the .map of a billboarded + // plane that represents the viewer's screen inside the env layer. + const envExternalTexture = new THREE.ExternalTexture(envTexture); + envExternalTexture.colorSpace = THREE.SRGBColorSpace; + (envExternalTexture as unknown as { image: unknown }).image = { + width: ENV_WIDTH, + height: ENV_HEIGHT, + }; + envExternalTexture.needsUpdate = true; + + // The cubemap is rendered each frame by THREE.CubeCamera (six + // perspective passes into the six faces) instead of being blitted + // from an equirect. CubeCamera renders the actual scene, so the + // reflection picks up any other geometry we add — not just the sky + // sphere that carries the camera feed. + // Two cube targets, one per mode. We swap which one CubeCamera writes + // into on toggle (resizing a single target via setSize doesn't fully + // reallocate cleanly on the WebGPU path, the reflection comes back + // blurry). Only the active one gets updated each frame so the cost + // stays paid for one mode at a time. + const makeCubeRT = (size: number) => { + const rt = new THREE.CubeRenderTarget(size); + rt.texture.mapping = THREE.CubeReflectionMapping; + rt.texture.colorSpace = THREE.SRGBColorSpace; + // Mipmaps only matter for PBR (roughness-aware sample). Chrome + // samples mip 0 only, so we skip the regen cost on the 512 target. + const wantsMips = size <= CUBE_SIZE_PBR; + rt.texture.generateMipmaps = wantsMips; + rt.texture.minFilter = wantsMips + ? THREE.LinearMipmapLinearFilter + : THREE.LinearFilter; + rt.texture.magFilter = THREE.LinearFilter; + return rt; + }; + const cubeRTPbr = makeCubeRT(CUBE_SIZE_PBR); + const cubeRTChrome = makeCubeRT(CUBE_SIZE_CHROME); + const activeCubeRT = () => + usePBRRef.current ? cubeRTPbr : cubeRTChrome; + + const scene = new THREE.Scene(); + // No scene.background — the canvas is alpha-cleared and the native + // camera preview View sits behind it (see JSX below). + + // Layer split: main camera sees layer 0 (helmet only) so the native + // preview View remains visible everywhere else; CubeCamera sees + // layer 1 (the reflection screen + gradient backdrop) so the helmet + // never reflects itself. + const ENV_LAYER = 1; + + // Live camera frame wrapped around the inside of a large back-side + // sphere. The CubeCamera at the helmet's center sees the camera in + // every direction, so reflections fully cover the helmet without + // any visible frame edges. + const envSphere = new THREE.Mesh( + new THREE.SphereGeometry(50, 64, 32), + new THREE.MeshBasicMaterial({ + map: envExternalTexture, + side: THREE.BackSide, + toneMapped: false, + }), + ); + envSphere.layers.set(ENV_LAYER); + scene.add(envSphere); + + const cubeCamera = new THREE.CubeCamera(0.1, 100, activeCubeRT()); + cubeCamera.layers.set(ENV_LAYER); + scene.add(cubeCamera); + + // PBR path: keep the GLTF's MeshStandardMaterial intact (albedo / + // normal / metalRoughness / AO from the original textures) and plug + // our live cubemap into each material's envMap. We capture the + // originals so the chrome toggle can swap back and forth without + // losing them. + const pbrMaterials = new Map< + THREE.Mesh, + THREE.Material | THREE.Material[] + >(); + gltf.scene.traverse((child) => { + const mesh = child as THREE.Mesh; + if (!mesh.isMesh) { + return; + } + pbrMaterials.set(mesh, mesh.material); + const mats = Array.isArray(mesh.material) + ? mesh.material + : [mesh.material]; + for (const m of mats) { + const std = m as THREE.MeshStandardMaterial; + std.envMap = cubeRTPbr.texture; + std.envMapIntensity = 1.0; + std.needsUpdate = true; + } + }); + + // Chrome path: every mesh shares a single MeshBasicMaterial that + // just samples the cubemap. No surface detail, but ~5-10x cheaper + // fragment cost, a useful A/B against PBR on the same scene. + // MeshBasicMaterial samples the envMap at full intensity (no Fresnel + // / roughness attenuation like PBR), then ACES tone-maps the result. + // Bright camera regions land in ACES's compressed range and read as + // washed-out. Dimming the base color multiplies the env sample down + // before tone mapping; ~70% matches a real chrome surface's + // reflectance and brings the brightness back in line with PBR. + const chromeMaterial = new THREE.MeshBasicMaterial({ + envMap: cubeRTChrome.texture, + color: 0xb0b0b0, + }); + + const applyPBR = (pbr: boolean) => { + ( + cubeCamera as unknown as { renderTarget: THREE.CubeRenderTarget } + ).renderTarget = pbr ? cubeRTPbr : cubeRTChrome; + for (const [mesh, original] of pbrMaterials) { + mesh.material = pbr ? original : chromeMaterial; + } + }; + applyPBR(usePBRRef.current); + applyPBRFnRef.current = applyPBR; + + scene.add(gltf.scene); + + // Drive the perspective from min(width, height) so the helmet keeps + // a consistent on-screen size in both orientations. three.js' + // PerspectiveCamera takes a *vertical* FOV; on portrait canvases we + // derive vFov from a fixed horizontal FOV so the wider dimension + // never under-frames the helmet. + const aspect = width / height; + const baseFov = 45; + let vFov = baseFov; + if (aspect < 1) { + const hFovRad = (baseFov * Math.PI) / 180; + const vFovRad = 2 * Math.atan(Math.tan(hFovRad / 2) / aspect); + vFov = (vFovRad * 180) / Math.PI; + } + const camera = new THREE.PerspectiveCamera(vFov, aspect, 0.25, 20); + camera.position.set(0, 0, 3); + + const clock = new THREE.Clock(); + const distance = 3; + let frameCount = 0; + const animate = () => { + // Slow time-based orbit around the helmet, matching the three.js + // env-map reference demo. + const elapsed = clock.getElapsedTime(); + camera.position.x = Math.sin(elapsed * 0.4) * distance; + camera.position.z = Math.cos(elapsed * 0.4) * distance; + camera.position.y = 0; + camera.lookAt(0, 0, 0); + + // Refresh the cubemap by rendering the env-layer (camera sphere) + // from six perspectives. Costlier than an equirect blit but lets + // us add other layer-1 props later that would also show up in + // reflections. + cubeCamera.update(renderer!, scene); + renderer!.render(scene, camera); + context.present(); + frameCount++; + if (frameCount === 1) { + console.log("[CameraHelmet] first three.js frame rendered"); + } + }; + renderer.setAnimationLoop(animate); + console.log("[CameraHelmet] animation loop started"); + + setPipelineState({ + device, + cameraPipeline, + cameraSampler, + envTexture, + envTextureView: envTexture.createView(), + }); + console.log("[CameraHelmet] pipelineState set, camera will activate"); + } catch (e) { + if (cancelled) { + return; + } + console.warn("[CameraHelmet] setup failed: " + String(e)); + setError(String(e)); + } + })(); + + return () => { + console.log("[CameraHelmet] setup-effect cleanup"); + cancelled = true; + applyPBRFnRef.current = null; + if (renderer) { + renderer.setAnimationLoop(null); + } + }; + }, [device, gltf]); + + // Frame processor worklet: copy the camera frame into envTexture each tick. + // The single device.queue is shared with three.js, so the helmet pass on + // the next rAF tick samples this frame's write. + const logBox = React.useMemo(() => ({ count: 0 }), []); + const frameOutput = useFrameOutput({ + pixelFormat: "native", + // 720p front-cam frames are plenty for the helmet's reflection — it's a + // small on-screen area. Keeping this low matters more in a multi-cam + // session, where both cameras share AVFoundation's bandwidth budget. + targetResolution: CommonResolutions.HD_16_9, + onFrame: (frame) => { + "worklet"; + logBox.count += 1; + if (logBox.count === 1) { + console.log( + "[CameraHelmet] worklet first frame, hasPipeline=" + + String(pipelineState != null) + + " frame=" + + String(frame.width) + + "x" + + String(frame.height), + ); + } + if (!pipelineState) { + frame.dispose(); + return; + } + const { + device: gpuDevice, + cameraPipeline, + cameraSampler, + envTextureView, + } = pipelineState; + const nativeBuffer = frame.getNativeBuffer(); + try { + const videoFrame = gpuDevice.createVideoFrameFromNativeBuffer( + nativeBuffer.pointer, + ); + try { + const externalTex = gpuDevice.importExternalTexture({ + source: videoFrame, + label: "camera-helmet-env", + }); + const bindGroup = gpuDevice.createBindGroup({ + layout: cameraPipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: externalTex }, + { binding: 1, resource: cameraSampler }, + ], + }); + const encoder = gpuDevice.createCommandEncoder(); + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: envTextureView, + clearValue: { r: 0, g: 0, b: 0, a: 1 }, + loadOp: "clear", + storeOp: "store", + }, + ], + }); + pass.setPipeline(cameraPipeline); + pass.setBindGroup(0, bindGroup); + pass.draw(3); + pass.end(); + gpuDevice.queue.submit([encoder.finish()]); + } finally { + videoFrame.release(); + } + } finally { + nativeBuffer.release(); + frame.dispose(); + } + }, + }); + + // ---- Multi-cam session ------------------------------------------------ + // useCamera always sets enableMultiCamSupport=false, so we drop down to + // the imperative API to drive two camera connections from a single + // session: front → frameOutput (helmet env), back → previewOutput + // (backdrop View). + const [session, setSession] = useState(null); + useEffect(() => { + if (!VisionCameraFactory.supportsMultiCamSessions) { + setError( + "This device doesn't support multi-cam sessions. Need an iPhone XS " + + "or newer / a comparable Android flagship.", + ); + return; + } + let cancelled = false; + let created: CameraSession | null = null; + (async () => { + const s = await VisionCameraFactory.createCameraSession(true); + if (cancelled) { + s.dispose(); + return; + } + created = s; + setSession(s); + })(); + return () => { + cancelled = true; + created?.stop(); + created?.dispose(); + }; + }, []); + + // Configure the session with two connections once everything is ready. + // We wait on pipelineState too because the worklet (which receives the + // front cam frames) only has somewhere to write once the env texture + + // camera-copy pipeline exist. + useEffect(() => { + if (!session || !backDevice || !frontDevice || !pipelineState) { + return; + } + console.log("[CameraHelmet] configuring multi-cam session"); + let cancelled = false; + let controllers: CameraController[] = []; + (async () => { + try { + controllers = await session.configure( + [ + { + input: backDevice, + outputs: [{ output: previewOutput, mirrorMode: "auto" }], + constraints: [], + }, + { + input: frontDevice, + outputs: [{ output: frameOutput, mirrorMode: "auto" }], + constraints: [], + }, + ], + {}, + ); + if (cancelled) { + controllers.forEach((c) => c.dispose()); + return; + } + console.log("[CameraHelmet] session configured, starting"); + session.start(); + } catch (e) { + if (cancelled) { + return; + } + console.warn("[CameraHelmet] session configure failed: " + String(e)); + setError(String(e)); + } + })(); + return () => { + cancelled = true; + session.stop(); + controllers.forEach((c) => c.dispose()); + }; + }, [ + session, + backDevice, + frontDevice, + previewOutput, + frameOutput, + pipelineState, + ]); + + if (error) { + return ( + + {error} + + ); + } + if (backDevice == null || frontDevice == null) { + return ( + + + Need both a back and a front camera. The iOS Simulator has none, and + some devices expose only one. + + + ); + } + return ( + + + + + {usePBR ? "PBR" : "Chrome"} + + + ); +}; + +const styles = StyleSheet.create({ + root: { flex: 1, backgroundColor: "black" }, + // Transparent canvas overlaid on the native camera preview view. + canvas: { ...StyleSheet.absoluteFillObject, backgroundColor: "transparent" }, + 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" }, + toggleButton: { + position: "absolute", + top: 60, + right: 16, + backgroundColor: "rgba(0, 0, 0, 0.6)", + paddingHorizontal: 16, + paddingVertical: 10, + borderRadius: 20, + minWidth: 88, + alignItems: "center", + }, + toggleButtonText: { color: "white", fontSize: 14, fontWeight: "600" }, +}); diff --git a/apps/example/src/ThreeJS/CameraSpheres.tsx b/apps/example/src/ThreeJS/CameraSpheres.tsx new file mode 100644 index 000000000..f3d7dedf3 --- /dev/null +++ b/apps/example/src/ThreeJS/CameraSpheres.tsx @@ -0,0 +1,489 @@ +import React, { useEffect, useRef, useState } from "react"; +import { + Linking, + Platform, + StyleSheet, + Text, + TouchableOpacity, + View, +} from "react-native"; +import type { CanvasRef } from "react-native-wgpu"; +import { Canvas } from "react-native-wgpu"; +import { + CommonResolutions, + NativePreviewView, + useCameraDevices, + useCameraPermission, + useFrameOutput, + usePreviewOutput, + VisionCamera as VisionCameraFactory, +} from "react-native-vision-camera"; +import type { + CameraController, + CameraSession, +} from "react-native-vision-camera"; +import * as THREE from "three"; + +import { makeWebGPURenderer } from "./components/makeWebGPURenderer"; +import { CAMERA_ENV_SHADER } from "./cameraEnvShader"; + +// Sibling of CameraHelmet but with three procedural chrome spheres in place +// of the GLTF helmet. Multi-cam setup: back camera renders behind the canvas +// as a native preview view, front camera feeds the cubemap that the spheres +// reflect. No PBR (no GLTF assets), so we skip mipmap generation entirely — +// the chrome look stays sharp on all surfaces. + +const ENV_WIDTH = 1024; +const ENV_HEIGHT = 512; + +const REQUIRED_FEATURES: GPUFeatureName[] = [ + "rnwebgpu/shared-texture-memory" as GPUFeatureName, + "dawn-multi-planar-formats" as GPUFeatureName, +]; + +const OPAQUE_YCBCR_EXT = + "opaque-ycbcr-android-for-external-texture" as GPUFeatureName; + +// Three big chrome spheres swirling around the origin, inspired by three.js' +// stereo-effects demo (which uses InstancedMesh + per-frame matrix updates). +// Even at 3 instances the InstancedMesh path is still nice because all +// three render in a single draw call. +const BEAD_COUNT = 3; +const BEAD_RADIUS = 0.55; +// XY radius of the swirl. Spheres orbit evenly spaced around the origin. +const SWIRL_RADIUS = 1.8; + +export const CameraSpheres = () => { + 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 Scene = () => { + useEffect(() => { + console.log("[CameraSpheres] Scene mounted"); + return () => console.log("[CameraSpheres] Scene unmounted"); + }, []); + const ref = useRef(null); + const previewOutput = usePreviewOutput(); + + const devices = useCameraDevices(); + const backDevice = React.useMemo( + () => devices.find((d) => d.position === "back"), + [devices], + ); + const frontDevice = React.useMemo( + () => devices.find((d) => d.position === "front"), + [devices], + ); + + const [pipelineState, setPipelineState] = useState<{ + device: GPUDevice; + cameraPipeline: GPURenderPipeline; + cameraSampler: GPUSampler; + envTexture: GPUTexture; + envTextureView: GPUTextureView; + } | null>(null); + const [error, setError] = useState(null); + const [device, setDevice] = useState(null); + + useEffect(() => { + let cancelled = false; + (async () => { + try { + const adapter = await navigator.gpu.requestAdapter(); + if (!adapter) { + throw new Error("requestAdapter returned null"); + } + const requiredFeatures = [...adapter.features] as GPUFeatureName[]; + const missing = REQUIRED_FEATURES.filter( + (f) => !adapter.features.has(f), + ); + const needsAndroidExt = + Platform.OS === "android" && !adapter.features.has(OPAQUE_YCBCR_EXT); + if (missing.length > 0 || needsAndroidExt) { + throw new Error( + "Adapter doesn't advertise the features the Vision Camera " + + "external-texture path needs: " + + `${[...missing, needsAndroidExt ? OPAQUE_YCBCR_EXT : null] + .filter(Boolean) + .join(", ")}.`, + ); + } + const d = await adapter.requestDevice({ requiredFeatures }); + if (cancelled) { + d.destroy(); + return; + } + setDevice(d); + } catch (e) { + if (cancelled) { + return; + } + console.warn("[CameraSpheres] device acquisition failed: " + String(e)); + setError(String(e)); + } + })(); + return () => { + cancelled = true; + }; + }, []); + + useEffect(() => { + if (!device) { + return; + } + const context = ref.current?.getContext("webgpu"); + if (!context) { + return; + } + let cancelled = false; + let renderer: THREE.WebGPURenderer | null = null; + + (async () => { + try { + const { width, height } = context.canvas; + + renderer = makeWebGPURenderer(context, { device, alpha: true }); + renderer.toneMapping = THREE.ACESFilmicToneMapping; + renderer.setClearColor(0x000000, 0); + await renderer.init(); + if (cancelled) { + return; + } + + const envTexture = device.createTexture({ + size: [ENV_WIDTH, ENV_HEIGHT], + format: "rgba8unorm", + usage: + GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.TEXTURE_BINDING, + }); + + const module = device.createShaderModule({ code: CAMERA_ENV_SHADER }); + const cameraPipeline = device.createRenderPipeline({ + layout: "auto", + vertex: { module, entryPoint: "vs_main" }, + fragment: { + module, + entryPoint: "fs_main", + targets: [{ format: "rgba8unorm" }], + }, + primitive: { topology: "triangle-list" }, + }); + const cameraSampler = device.createSampler({ + magFilter: "linear", + minFilter: "linear", + }); + + const envExternalTexture = new THREE.ExternalTexture(envTexture); + envExternalTexture.mapping = THREE.EquirectangularReflectionMapping; + envExternalTexture.colorSpace = THREE.SRGBColorSpace; + (envExternalTexture as unknown as { image: unknown }).image = { + width: ENV_WIDTH, + height: ENV_HEIGHT, + }; + envExternalTexture.needsUpdate = true; + + // Cube target refreshed per frame from the equirect — same dynamic + // env trick as CameraHelmet. No mipmap chain: the spheres use + // MeshBasicMaterial which samples mip 0 unconditionally, so the + // auto-mipmap regeneration would be pure waste. + const cubeRT = new THREE.CubeRenderTarget(ENV_HEIGHT); + cubeRT.texture.mapping = THREE.CubeReflectionMapping; + cubeRT.texture.colorSpace = THREE.SRGBColorSpace; + + const scene = new THREE.Scene(); + // Single InstancedMesh: all three spheres in one draw call. Mesh + // tessellation matters more here than in the swarm version because + // each sphere is bigger on screen, so 48x32 segments instead of + // 16x8 — smoother silhouettes against the panorama backdrop. + const chrome = new THREE.MeshBasicMaterial({ envMap: cubeRT.texture }); + const geometry = new THREE.SphereGeometry(BEAD_RADIUS, 48, 32); + const beads = new THREE.InstancedMesh(geometry, chrome, BEAD_COUNT); + beads.instanceMatrix.setUsage(THREE.DynamicDrawUsage); + scene.add(beads); + + // Seed each instance's matrix to identity. The animate loop + // overwrites the translation column each frame; scale and rotation + // stay as identity (= sphere radius set by BEAD_RADIUS above). + const dummy = new THREE.Object3D(); + for (let i = 0; i < BEAD_COUNT; i++) { + dummy.updateMatrix(); + beads.setMatrixAt(i, dummy.matrix); + } + beads.instanceMatrix.needsUpdate = true; + const beadPos = new THREE.Vector3(); + + const aspect = width / height; + const baseFov = 60; + let vFov = baseFov; + if (aspect < 1) { + const hFovRad = (baseFov * Math.PI) / 180; + const vFovRad = 2 * Math.atan(Math.tan(hFovRad / 2) / aspect); + vFov = (vFovRad * 180) / Math.PI; + } + const camera = new THREE.PerspectiveCamera(vFov, aspect, 0.25, 20); + camera.position.set(0, 0, 4); + + const clock = new THREE.Clock(); + const animate = () => { + // Slow rotation of the three spheres around the origin in the XY + // plane. Phase offsets are evenly spaced (2π/3 apart) so the + // spheres form a rotating equilateral triangle, never overlapping. + const elapsed = clock.getElapsedTime() * 0.4; + for (let i = 0; i < BEAD_COUNT; i++) { + const angle = elapsed + (i * 2 * Math.PI) / BEAD_COUNT; + beadPos.set( + SWIRL_RADIUS * Math.cos(angle), + SWIRL_RADIUS * Math.sin(angle), + 0, + ); + beads.getMatrixAt(i, dummy.matrix); + dummy.matrix.setPosition(beadPos); + beads.setMatrixAt(i, dummy.matrix); + } + beads.instanceMatrix.needsUpdate = true; + + cubeRT.fromEquirectangularTexture(renderer!, envExternalTexture); + renderer!.render(scene, camera); + context.present(); + }; + renderer.setAnimationLoop(animate); + + setPipelineState({ + device, + cameraPipeline, + cameraSampler, + envTexture, + envTextureView: envTexture.createView(), + }); + } catch (e) { + if (cancelled) { + return; + } + console.warn("[CameraSpheres] setup failed: " + String(e)); + setError(String(e)); + } + })(); + + return () => { + cancelled = true; + if (renderer) { + renderer.setAnimationLoop(null); + } + }; + }, [device]); + + const logBox = React.useMemo(() => ({ count: 0 }), []); + const frameOutput = useFrameOutput({ + pixelFormat: "native", + targetResolution: CommonResolutions.HD_16_9, + onFrame: (frame) => { + "worklet"; + logBox.count += 1; + if (logBox.count === 1) { + console.log( + "[CameraSpheres] worklet first frame, frame=" + + String(frame.width) + + "x" + + String(frame.height), + ); + } + if (!pipelineState) { + frame.dispose(); + return; + } + const { + device: gpuDevice, + cameraPipeline, + cameraSampler, + envTextureView, + } = pipelineState; + const nativeBuffer = frame.getNativeBuffer(); + try { + const videoFrame = gpuDevice.createVideoFrameFromNativeBuffer( + nativeBuffer.pointer, + ); + try { + const externalTex = gpuDevice.importExternalTexture({ + source: videoFrame, + label: "camera-spheres-env", + }); + const bindGroup = gpuDevice.createBindGroup({ + layout: cameraPipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: externalTex }, + { binding: 1, resource: cameraSampler }, + ], + }); + const encoder = gpuDevice.createCommandEncoder(); + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: envTextureView, + clearValue: { r: 0, g: 0, b: 0, a: 1 }, + loadOp: "clear", + storeOp: "store", + }, + ], + }); + pass.setPipeline(cameraPipeline); + pass.setBindGroup(0, bindGroup); + pass.draw(3); + pass.end(); + gpuDevice.queue.submit([encoder.finish()]); + } finally { + videoFrame.release(); + } + } finally { + nativeBuffer.release(); + frame.dispose(); + } + }, + }); + + const [session, setSession] = useState(null); + useEffect(() => { + if (!VisionCameraFactory.supportsMultiCamSessions) { + setError( + "This device doesn't support multi-cam sessions. Need an iPhone XS " + + "or newer / a comparable Android flagship.", + ); + return; + } + let cancelled = false; + let created: CameraSession | null = null; + (async () => { + const s = await VisionCameraFactory.createCameraSession(true); + if (cancelled) { + s.dispose(); + return; + } + created = s; + setSession(s); + })(); + return () => { + cancelled = true; + created?.stop(); + created?.dispose(); + }; + }, []); + + useEffect(() => { + if (!session || !backDevice || !frontDevice || !pipelineState) { + return; + } + let cancelled = false; + let controllers: CameraController[] = []; + (async () => { + try { + controllers = await session.configure( + [ + { + input: backDevice, + outputs: [{ output: previewOutput, mirrorMode: "auto" }], + constraints: [], + }, + { + input: frontDevice, + outputs: [{ output: frameOutput, mirrorMode: "auto" }], + constraints: [], + }, + ], + {}, + ); + if (cancelled) { + controllers.forEach((c) => c.dispose()); + return; + } + session.start(); + } catch (e) { + if (cancelled) { + return; + } + console.warn("[CameraSpheres] session configure failed: " + String(e)); + setError(String(e)); + } + })(); + return () => { + cancelled = true; + session.stop(); + controllers.forEach((c) => c.dispose()); + }; + }, [ + session, + backDevice, + frontDevice, + previewOutput, + frameOutput, + pipelineState, + ]); + + if (error) { + return ( + + {error} + + ); + } + if (backDevice == null || frontDevice == null) { + return ( + + + Need both a back and a front camera. The iOS Simulator has none, and + some devices expose only one. + + + ); + } + return ( + + + + + ); +}; + +const styles = StyleSheet.create({ + root: { flex: 1, backgroundColor: "black" }, + canvas: { ...StyleSheet.absoluteFillObject, backgroundColor: "transparent" }, + 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/ThreeJS/List.tsx b/apps/example/src/ThreeJS/List.tsx index 9601abb71..ad3547fdb 100644 --- a/apps/example/src/ThreeJS/List.tsx +++ b/apps/example/src/ThreeJS/List.tsx @@ -23,6 +23,14 @@ export const examples = [ screen: "Helmet", title: "⛑️ Helmet", }, + { + screen: "CameraHelmet", + title: "📷 Camera Env Sphere", + }, + { + screen: "CameraSpheres", + title: "📷 Camera Env Spheres", + }, { screen: "PostProcessing", title: "🪄 Post Processing Effects", diff --git a/apps/example/src/ThreeJS/Routes.ts b/apps/example/src/ThreeJS/Routes.ts index cb76ac65f..68971cef5 100644 --- a/apps/example/src/ThreeJS/Routes.ts +++ b/apps/example/src/ThreeJS/Routes.ts @@ -2,6 +2,8 @@ export type Routes = { List: undefined; Cube: undefined; Helmet: undefined; + CameraHelmet: undefined; + CameraSpheres: undefined; Backdrop: undefined; InstancedMesh: undefined; Fiber: undefined; diff --git a/apps/example/src/ThreeJS/cameraEnvShader.ts b/apps/example/src/ThreeJS/cameraEnvShader.ts new file mode 100644 index 000000000..2222ef444 --- /dev/null +++ b/apps/example/src/ThreeJS/cameraEnvShader.ts @@ -0,0 +1,49 @@ +// Tiny "copy camera frame into an rgba8unorm texture" shader. The output +// texture is then wrapped in a THREE.ExternalTexture and mapped onto a +// billboarded plane that three.js' CubeCamera bakes into the helmet's +// envMap — i.e. it acts as a virtual screen at the viewer's location +// rather than a 360° panorama. The destination texture's aspect (9:16) is +// chosen to match the camera frame's post-rotation aspect so no stretching +// happens here; the fullscreen triangle just rotates+mirrors the source to +// selfie-upright. + +export const CAMERA_ENV_SHADER = /* wgsl */ ` +struct VsOut { + @builtin(position) position: vec4f, + @location(0) uv: vec2f, +}; + +@group(0) @binding(0) var srcTex: texture_external; +@group(0) @binding(1) var srcSampler: sampler; + +@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 { + // Front camera: iOS delivers landscape-orientation frames with the + // horizontal axis already mirrored (selfie convention). To bring those + // upright in the equirect we (a) compensate for the horizontal mirror + // by sampling at (1-x) and (b) rotate 90° CCW with V flipped, giving + // (1-v, 1-u). Equivalent to the 90° CW back-cam mapping (v, 1-u) with + // its U axis pre-flipped to undo the mirror. + let rotatedUv = vec2f(1.0 - in.uv.y, 1.0 - in.uv.x); + let c = textureSampleBaseClampToEdge(srcTex, srcSampler, rotatedUv); + return vec4f(c.rgb, 1.0); +} +`; diff --git a/apps/example/src/ThreeJS/components/makeWebGPURenderer.ts b/apps/example/src/ThreeJS/components/makeWebGPURenderer.ts index 39c4bf399..f6729b3c9 100644 --- a/apps/example/src/ThreeJS/components/makeWebGPURenderer.ts +++ b/apps/example/src/ThreeJS/components/makeWebGPURenderer.ts @@ -60,12 +60,22 @@ export class ReactNativeCanvas { export const makeWebGPURenderer = ( context: GPUCanvasContext, - { antialias = true }: { antialias?: boolean } = {}, + { + antialias = true, + device, + alpha = false, + }: { antialias?: boolean; device?: GPUDevice; alpha?: boolean } = {}, ) => new THREE.WebGPURenderer({ antialias, + alpha, // eslint-disable-next-line @typescript-eslint/ban-ts-comment // @ts-expect-error canvas: new ReactNativeCanvas(context.canvas), context, + // When supplied, three.js skips its own adapter/device acquisition and + // uses this device. Lets callers request custom features (e.g. Dawn's + // shared-texture-memory) that three.js doesn't include in its default + // GPUFeatureName enum walk. + ...(device ? { device } : {}), }); diff --git a/apps/example/src/ThreeJS/index.tsx b/apps/example/src/ThreeJS/index.tsx index 244c43d71..e679d88e4 100644 --- a/apps/example/src/ThreeJS/index.tsx +++ b/apps/example/src/ThreeJS/index.tsx @@ -6,6 +6,8 @@ import { Cube } from "./Cube"; import type { Routes } from "./Routes"; import { List } from "./List"; import { Helmet } from "./Helmet"; +import { CameraHelmet } from "./CameraHelmet"; +import { CameraSpheres } from "./CameraSpheres"; import { Backdrop } from "./Backdrop"; import { InstancedMesh } from "./InstancedMesh"; import { Fiber } from "./Fiber"; @@ -73,6 +75,20 @@ export const ThreeJS = () => { title: "⛑️ Helmet", }} /> + + 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..db629f12a --- /dev/null +++ b/apps/example/src/VisionCamera/VisionCamera.tsx @@ -0,0 +1,827 @@ +import React, { useEffect } 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"; + +// 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. + +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 textureSampleBaseClampToEdge(srcTex, srcSampler, 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); +} +`; + +const REQUIRED_FEATURES: GPUFeatureName[] = [ + "rnwebgpu/shared-texture-memory" 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. Mirrors the ExternalTexture demo: 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] = React.useState<{ + adapter: GPUAdapter; + device: GPUDevice; + } | null>(null); + const [deviceError, setDeviceError] = React.useState(null); + React.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; + 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 = React.useMemo( + () => + devices.find((d) => d.position === "back") ?? + devices.find((d) => d.position === "front") ?? + devices[0], + [devices], + ); + + const [pipelineState, setPipelineState] = React.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] = React.useState(null); + const [modes, setModes] = React.useState(INITIAL_MODES); + const cycle = React.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: 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 (matches ExternalTexture's "Blur" chain) ---- + 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: 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 = React.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), + ); + } + 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 device, not on the + // RNWebGPU global — `device` is already box-able across worklet + // runtimes via the WebGPU custom serializer (proven by the + // Reanimated demo); RNWebGPU is a main-runtime-only global. + videoFrame = device.createVideoFrameFromNativeBuffer( + nativeBuffer.pointer, + ); + } catch (e) { + console.warn( + "[VisionCamera] createVideoFrameFromNativeBuffer threw: " + + String(e), + ); + throw e; + } + try { + // 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 = videoFrame.width / videoFrame.height; + 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", + }); + } 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([ + videoFrame.width, + videoFrame.height, + 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()]); + 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..e589bff07 --- /dev/null +++ b/apps/example/src/VisionCamera/blurShaders.ts @@ -0,0 +1,141 @@ +// 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; + let c = textureSampleBaseClampToEdge( + srcTex, + srcSampler, + 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..7e8ed57a8 --- /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: 1, + 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/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..64e713fd5 100644 --- a/packages/webgpu/apple/AppleVideoPlayer.mm +++ b/packages/webgpu/apple/AppleVideoPlayer.mm @@ -9,6 +9,59 @@ namespace { +// 3x4 row-major matrices mapping [Y, U, V, 1] to linear RGB. +// 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 kBT709LimitedToLinearRGB[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 kBT601LimitedToLinearRGB[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 kBT2020LimitedToLinearRGB[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 = kBT709LimitedToLinearRGB; + if (matrixKey) { + auto matrix = (CFStringRef)matrixKey; + if (CFEqual(matrix, kCVImageBufferYCbCrMatrix_ITU_R_601_4) || + CFEqual(matrix, kCVImageBufferYCbCrMatrix_SMPTE_240M_1995)) { + src = kBT601LimitedToLinearRGB; + } else if (CFEqual(matrix, kCVImageBufferYCbCrMatrix_ITU_R_2020)) { + src = kBT2020LimitedToLinearRGB; + } + } + 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 +83,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 +113,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 +156,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..cdab10da4 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; }; @@ -54,6 +70,14 @@ class PlatformContext { PlatformContext() = default; virtual ~PlatformContext() = default; + // Singleton-style accessor so leaf classes (e.g. GPUDevice) can reach the + // platform context without threading it through every constructor. Set by + // RNWebGPUManager at startup. + static std::shared_ptr &global() { + static std::shared_ptr instance; + return instance; + } + virtual wgpu::Surface makeSurface(wgpu::Instance instance, void *surface, int width, int height) = 0; virtual ImageData createImageBitmap(std::string blobId, double offset, @@ -87,8 +111,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..6b8a28cb3 100644 --- a/packages/webgpu/cpp/rnwgpu/RNWebGPUManager.cpp +++ b/packages/webgpu/cpp/rnwgpu/RNWebGPUManager.cpp @@ -63,10 +63,21 @@ RNWebGPUManager::RNWebGPUManager( // Register main runtime for RuntimeAwareCache BaseRuntimeAwareCache::setMainJsRuntime(_jsRuntime); + // Expose the platform context for leaf classes that need it (e.g. + // GPUDevice::createVideoFrameFromNativeBuffer) without threading it through + // every constructor. + PlatformContext::global() = _platformContext; + auto gpu = std::make_shared(*_jsRuntime); 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..8d4e29bd1 100644 --- a/packages/webgpu/cpp/rnwgpu/api/Convertors.h +++ b/packages/webgpu/cpp/rnwgpu/api/Convertors.h @@ -729,7 +729,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..b9d2a46bd 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPU.cpp +++ b/packages/webgpu/cpp/rnwgpu/api/GPU.cpp @@ -20,6 +20,24 @@ 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; application code + // still has to list each one in requiredFeatures. expose_wgsl_experimental_features + // is the parallel toggle for WGSL language features. + 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/GPUDevice.cpp b/packages/webgpu/cpp/rnwgpu/api/GPUDevice.cpp index f80d7fadf..4a0fdba0c 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUDevice.cpp +++ b/packages/webgpu/cpp/rnwgpu/api/GPUDevice.cpp @@ -8,6 +8,7 @@ #include "Convertors.h" #include "JSIConverter.h" +#include "PlatformContext.h" #include "GPUFeatures.h" #include "GPUInternalError.h" @@ -233,10 +234,239 @@ std::shared_ptr GPUDevice::createPipelineLayout( _instance.CreatePipelineLayout(&desc), descriptor->label.value_or("")); } +// 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 +}; + std::shared_ptr GPUDevice::importExternalTexture( std::shared_ptr descriptor) { + if (!descriptor || !descriptor->source) { + throw std::runtime_error( + "GPUDevice::importExternalTexture(): descriptor.source (VideoFrame) " + "is required"); + } + const auto &source = descriptor->source; + const auto &frame = source->handle(); + if (frame.handle == nullptr) { + throw std::runtime_error( + "GPUDevice::importExternalTexture(): 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 = _instance.ImportSharedTextureMemory(&memDesc); + if (memory == nullptr) { + throw std::runtime_error( + "GPUDevice::importExternalTexture(): 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( + "GPUDevice::importExternalTexture(): CreateTexture returned null"); + } + + // 3. Begin access on the underlying memory. The matching EndAccess runs in + // the GPUExternalTexture destructor. + wgpu::SharedTextureMemoryBeginAccessDescriptor begin{}; + begin.initialized = true; + begin.concurrentRead = false; + if (!memory.BeginAccess(texture, &begin)) { + throw std::runtime_error( + "GPUDevice::importExternalTexture(): 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; + if (isYuv) { + extDesc.plane1 = plane1; + extDesc.yuvToRgbConversionMatrix = frame.yuvToRgbMatrix; + extDesc.srcTransferFunctionParameters = kSrgbDecodeParams; + extDesc.dstTransferFunctionParameters = kSrgbEncodeParams; + extDesc.gamutConversionMatrix = kIdentityGamutMatrix; + } + extDesc.cropOrigin = {0, 0}; + extDesc.cropSize = {frame.width, frame.height}; + extDesc.apparentSize = {frame.width, frame.height}; + + auto external = _instance.CreateExternalTexture(&extDesc); + if (external == nullptr) { + wgpu::SharedTextureMemoryEndAccessState state{}; + (void)memory.EndAccess(texture, &state); + throw std::runtime_error( + "GPUDevice::importExternalTexture(): 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 = _instance.ImportSharedTextureMemory(&memDesc); + if (memory == nullptr) { + throw std::runtime_error( + "GPUDevice::importExternalTexture(): 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( + "GPUDevice::importExternalTexture(): 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( + "GPUDevice::importExternalTexture(): BeginAccess failed"); + } + + // 4. Build the ExternalTextureDescriptor. Unlike iOS we do *not* split + // planes or pass an explicit YUV→RGB matrix: when the underlying texture + // is OpaqueYCbCrAndroid, Dawn routes sampling through a Vulkan + // SamplerYcbcrConversion that does the conversion implicitly, driven by + // the AHB's own format metadata. We still must pass noop gamut/transfer + // arrays: Dawn's ComputeExternalTextureParams unconditionally dereferences + // gamutConversionMatrix / src/dstTransferFunctionParameters (see + // externals/dawn/.../ExternalTexture.cpp), so leaving them null produces a + // silent black sample. Identity transfer = TransferFunctionToArray of + // kEOTF_Identity ({g=1,a=1,rest=0}); identity gamut = 3x3 identity. + 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 + }; + wgpu::ExternalTextureDescriptor extDesc{}; + if (!label.empty()) { + extDesc.label = wgpu::StringView(label.c_str(), label.size()); + } + extDesc.plane0 = texture.CreateView(); + extDesc.cropOrigin = {0, 0}; + extDesc.cropSize = {frame.width, frame.height}; + extDesc.apparentSize = {frame.width, frame.height}; + extDesc.gamutConversionMatrix = kIdentityGamutMatrix; + extDesc.srcTransferFunctionParameters = kIdentityTransferParams; + extDesc.dstTransferFunctionParameters = kIdentityTransferParams; + + auto external = _instance.CreateExternalTexture(&extDesc); + if (external == nullptr) { + wgpu::SharedTextureMemoryEndAccessState state{}; + (void)memory.EndAccess(texture, &state); + throw std::runtime_error( + "GPUDevice::importExternalTexture(): 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( - "GPUDevice::importExternalTexture(): Not implemented"); + "GPUDevice::importExternalTexture(): not yet implemented on this " + "platform"); +#endif +} + +std::shared_ptr +GPUDevice::createVideoFrameFromNativeBuffer(uint64_t pointer) { + auto platformContext = PlatformContext::global(); + if (!platformContext) { + throw std::runtime_error( + "GPUDevice::createVideoFrameFromNativeBuffer(): PlatformContext is " + "not initialized"); + } + auto handle = + platformContext->wrapNativeBuffer(reinterpret_cast(pointer)); + return std::make_shared(std::move(handle)); } std::shared_ptr GPUDevice::importSharedTextureMemory( diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUDevice.h b/packages/webgpu/cpp/rnwgpu/api/GPUDevice.h index 2ab1ddd14..28add6e0c 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUDevice.h +++ b/packages/webgpu/cpp/rnwgpu/api/GPUDevice.h @@ -53,6 +53,7 @@ #include "GPUTexture.h" #include "GPUTextureDescriptor.h" #include "GPUUncapturedErrorEvent.h" +#include "VideoFrame.h" namespace rnwgpu { @@ -120,6 +121,12 @@ class GPUDevice : public NativeObject { std::shared_ptr descriptor); std::shared_ptr importSharedTextureMemory( std::shared_ptr descriptor); + // Wrap a CVPixelBufferRef / AHardwareBuffer* pointer (as a BigInt) into a + // VideoFrame. Mirrors RNWebGPU.createVideoFrameFromNativeBuffer but is + // reachable from worklet runtimes since GPUDevice is already serialized + // across the worklet boundary via the WebGPU custom serializer. + std::shared_ptr + createVideoFrameFromNativeBuffer(uint64_t pointer); std::shared_ptr createBindGroupLayout( std::shared_ptr descriptor); std::shared_ptr @@ -175,6 +182,8 @@ class GPUDevice : public NativeObject { &GPUDevice::importExternalTexture); installMethod(runtime, prototype, "importSharedTextureMemory", &GPUDevice::importSharedTextureMemory); + installMethod(runtime, prototype, "createVideoFrameFromNativeBuffer", + &GPUDevice::createVideoFrameFromNativeBuffer); installMethod(runtime, prototype, "createBindGroupLayout", &GPUDevice::createBindGroupLayout); installMethod(runtime, prototype, "createPipelineLayout", diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUExternalTexture.h b/packages/webgpu/cpp/rnwgpu/api/GPUExternalTexture.h index 9be5efe6f..71e23f3b7 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" @@ -16,8 +19,23 @@ 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) {} + // 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 { + if (_memory && _texture) { + wgpu::SharedTextureMemoryEndAccessState state{}; + (void)_memory.EndAccess(_texture, &state); + } + } public: std::string getBrand() { return CLASS_NAME; } @@ -39,6 +57,9 @@ class GPUExternalTexture : public NativeObject { 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/VideoFrame.h b/packages/webgpu/cpp/rnwgpu/api/VideoFrame.h index af09cd127..f001285e0 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 { @@ -21,10 +22,8 @@ class VideoFrame : public NativeObject { public: static constexpr const char *CLASS_NAME = "VideoFrame"; - 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 +31,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 +57,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/GPUExternalTextureDescriptor.h b/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUExternalTextureDescriptor.h index fba6721ee..da7affc9b 100644 --- a/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUExternalTextureDescriptor.h +++ b/packages/webgpu/cpp/rnwgpu/api/descriptors/GPUExternalTextureDescriptor.h @@ -3,35 +3,25 @@ #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". 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; }; -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 +30,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")) { @@ -67,12 +49,11 @@ struct JSIConverter> { } } } - 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/src/index.tsx b/packages/webgpu/src/index.tsx index 4672dfc8b..09d2877c3 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; }; @@ -44,6 +54,10 @@ declare global { importSharedTextureMemory( descriptor: GPUSharedTextureMemoryDescriptor, ): GPUSharedTextureMemory; + // Wrap a NativeBuffer.pointer into a NativeVideoFrame. Reachable from + // worklet runtimes (e.g. Vision Camera frame processors) because GPUDevice + // is serialized across worklet boundaries via the WebGPU custom serializer. + createVideoFrameFromNativeBuffer(pointer: bigint): NativeVideoFrame; } // Extend createImageBitmap to accept ArrayBuffer/TypedArray (encoded image bytes) 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; }