From 3d364374a04d71e6bcef58a0b06e3053b793678c Mon Sep 17 00:00:00 2001 From: Marc Rousavy Date: Fri, 21 Aug 2026 12:33:06 +0200 Subject: [PATCH 1/4] feat: GPUBuffer.readSync() - synchronous small-buffer readback Adds a non-spec extension for the GPU-compute readback pattern that mapAsync cannot serve: render/worklet loops that must consume a compute result in the SAME frame (hand/pose landmarks, histogram ranges, GPU picking ids, counters). Awaiting mapAsync from such a loop forces at least one frame of staleness; readSync blocks the calling thread until previously submitted work touching the buffer completes and returns an owned copy of the bytes. Implementation: MapAsync with CallbackMode::WaitAnyOnly + a 2s Instance::WaitAny - the instance already enables TimedWaitAny. External instances without the feature (e.g. Skia-provided) fail the wait and throw instead of hanging. Capped at 1 MiB: this is a primitive for tiny results, not bulk transfers; the async path remains the right tool there. Requires MAP_READ usage, so the usage rules push callers to the correct copy-to-staging pattern by construction. Typed via the existing non-spec declare-global block, with tests covering the same-tick compute readback, offset/size, repeated reads, and the size cap. --- packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp | 64 +++++++++++++ packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h | 3 + .../webgpu/src/__tests__/ReadSync.spec.ts | 93 +++++++++++++++++++ packages/webgpu/src/index.tsx | 15 +++ 4 files changed, 175 insertions(+) create mode 100644 packages/webgpu/src/__tests__/ReadSync.spec.ts diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp index db698fff7..2c2039c7b 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp +++ b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp @@ -1,5 +1,6 @@ #include "GPUBuffer.h" +#include #include #include @@ -35,6 +36,69 @@ GPUBuffer::getMappedRange(std::optional o, std::optional size) { return array_buffer; } +namespace { +// readSync returns a copy that outlives the mapping, so it needs an +// ArrayBuffer that owns (and frees) its backing store - the base class +// wraps foreign memory and never frees. +struct OwnedArrayBuffer : ArrayBuffer { + explicit OwnedArrayBuffer(size_t size) + : ArrayBuffer(malloc(size), size, 1) {} + ~OwnedArrayBuffer() override { free(_data); } +}; +} // namespace + +std::shared_ptr +GPUBuffer::readSync(std::optional o, std::optional sizeIn) { + // Synchronous small-buffer readback: blocks the calling thread until all + // previously submitted GPU work using this buffer completes, then returns + // a copy of the mapped bytes. Built for tiny compute results (landmarks, + // ranges, counters) that must be consumed in the SAME frame - the async + // mapAsync path forces at least one frame of staleness in render loops + // that cannot await. Requires MAP_READ usage (pair with COPY_DST and copy + // into this buffer from your storage buffer). The wait uses + // Instance::WaitAny, which this library's instance enables via the + // TimedWaitAny feature at creation; external (Skia-provided) instances + // without it fail the wait and throw rather than hang. + size_t offset = o.value_or(0); + size_t size = sizeIn.has_value() + ? sizeIn.value() + : static_cast(_instance.GetSize() - offset); + constexpr size_t kMaxReadSyncBytes = 1 << 20; + if (size > kMaxReadSyncBytes) { + throw std::runtime_error( + "readSync is intended for small readbacks (<= 1 MiB); use mapAsync " + "for large buffers"); + } + wgpu::MapAsyncStatus mapStatus = wgpu::MapAsyncStatus::Error; + std::string mapMessage = "callback never ran"; + auto future = _instance.MapAsync( + wgpu::MapMode::Read, offset, size, wgpu::CallbackMode::WaitAnyOnly, + [&mapStatus, &mapMessage](wgpu::MapAsyncStatus status, + wgpu::StringView message) { + mapStatus = status; + mapMessage = std::string(message); + }); + constexpr uint64_t kTimeoutNs = 2'000'000'000; // 2s: a hung GPU, not a wait + auto waitStatus = _async->instance().WaitAny(future, kTimeoutNs); + if (waitStatus != wgpu::WaitStatus::Success) { + throw std::runtime_error( + "readSync: WaitAny did not complete (timeout, or the instance lacks " + "the TimedWaitAny feature)"); + } + if (mapStatus != wgpu::MapAsyncStatus::Success) { + throw std::runtime_error("readSync: mapping failed: " + mapMessage); + } + const void *ptr = _instance.GetConstMappedRange(offset, size); + if (ptr == nullptr) { + _instance.Unmap(); + throw std::runtime_error("readSync: GetConstMappedRange failed"); + } + auto result = std::make_shared(size); + memcpy(result->data(), ptr, size); + _instance.Unmap(); + return result; +} + void GPUBuffer::destroy() { _instance.Destroy(); } async::AsyncTaskHandle GPUBuffer::mapAsync(jsi::Runtime &runtime, diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h index 1b5cf123e..d996db160 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h +++ b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h @@ -38,6 +38,8 @@ class GPUBuffer : public NativeObject { std::optional size); std::shared_ptr getMappedRange(std::optional offset, std::optional size); + std::shared_ptr readSync(std::optional offset, + std::optional size); void unmap(); void destroy(); @@ -57,6 +59,7 @@ class GPUBuffer : public NativeObject { &GPUBuffer::mapAsync); installMethod(runtime, prototype, "getMappedRange", &GPUBuffer::getMappedRange); + installMethod(runtime, prototype, "readSync", &GPUBuffer::readSync); installMethod(runtime, prototype, "unmap", &GPUBuffer::unmap); installMethod(runtime, prototype, "destroy", &GPUBuffer::destroy); installGetter(runtime, prototype, "size", &GPUBuffer::getSize); diff --git a/packages/webgpu/src/__tests__/ReadSync.spec.ts b/packages/webgpu/src/__tests__/ReadSync.spec.ts new file mode 100644 index 000000000..1e1f4d0bb --- /dev/null +++ b/packages/webgpu/src/__tests__/ReadSync.spec.ts @@ -0,0 +1,93 @@ +import { client } from "./setup"; + +describe("readSync", () => { + it("reads back compute results synchronously, same tick", async () => { + const result = await client.eval(({ device }) => { + const storage = device.createBuffer({ + size: 16, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + const staging = device.createBuffer({ + size: 16, + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, + }); + const module = device.createShaderModule({ + code: ` +@group(0) @binding(0) var out: array; +@compute @workgroup_size(1) +fn main() { + out[0] = 1u; + out[1] = 2u; + out[2] = 3u; + out[3] = 42u; +}`, + }); + const pipeline = device.createComputePipeline({ + layout: "auto", + compute: { module, entryPoint: "main" }, + }); + const bindGroup = device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [{ binding: 0, resource: { buffer: storage } }], + }); + const encoder = device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bindGroup); + pass.dispatchWorkgroups(1); + pass.end(); + encoder.copyBufferToBuffer(storage, 0, staging, 0, 16); + device.queue.submit([encoder.finish()]); + // No await between submit and read: the readback is synchronous. + return Array.from(new Uint32Array(staging.readSync())); + }); + expect(result).toEqual([1, 2, 3, 42]); + }); + it("respects offset and size", async () => { + const result = await client.eval(({ device }) => { + const staging = device.createBuffer({ + size: 16, + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, + }); + device.queue.writeBuffer( + staging, + 0, + new Uint32Array([10, 20, 30, 40]).buffer, + ); + return Array.from(new Uint32Array(staging.readSync(8, 8))); + }); + expect(result).toEqual([30, 40]); + }); + it("is repeatable on the same buffer", async () => { + const result = await client.eval(({ device }) => { + const staging = device.createBuffer({ + size: 4, + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, + }); + const reads: number[] = []; + for (let i = 0; i < 3; i++) { + device.queue.writeBuffer(staging, 0, new Uint32Array([i]).buffer); + reads.push(new Uint32Array(staging.readSync())[0]!); + } + return reads; + }); + expect(result).toEqual([0, 1, 2]); + }); + it("rejects reads above the size cap", async () => { + const result = await client.eval(({ device }) => { + const staging = device.createBuffer({ + size: 2 * 1024 * 1024, + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, + }); + try { + staging.readSync(); + return "no error"; + } catch (e) { + return e instanceof Error && e.message.includes("small readbacks") + ? "capped" + : "wrong error"; + } + }); + expect(result).toBe("capped"); + }); +}); diff --git a/packages/webgpu/src/index.tsx b/packages/webgpu/src/index.tsx index 5db6b3d01..e3711f892 100644 --- a/packages/webgpu/src/index.tsx +++ b/packages/webgpu/src/index.tsx @@ -80,6 +80,21 @@ declare global { readonly nativePointer: bigint; } + // Non-spec RN extension: synchronous small-buffer readback. Blocks the + // calling thread until all previously submitted GPU work using this buffer + // completes, then returns an owned copy of the mapped bytes. Built for + // tiny compute results (landmarks, histogram ranges, counters, picking + // ids) that a render/worklet loop must consume in the SAME frame - such + // loops cannot await mapAsync without accepting a frame of staleness. + // Requires MAP_READ usage; the only valid companion is COPY_DST, so the + // pattern is: copy from your storage buffer into this staging buffer, + // submit, readSync(). Capped at 1 MiB (use mapAsync for bulk data); throws + // instead of hanging if the instance lacks TimedWaitAny (external/Skia + // instances) or the GPU hangs past 2s. + interface GPUBuffer { + readSync(offset?: number, size?: number): ArrayBuffer; + } + interface GPUDevice { importSharedTextureMemory( descriptor: GPUSharedTextureMemoryDescriptor, From 08b1808a08158fef1dc2ac7601b131a3e20c0ccb Mon Sep 17 00:00:00 2001 From: Marc Rousavy Date: Fri, 21 Aug 2026 15:58:12 +0200 Subject: [PATCH 2/4] fix: address GPUBuffer.readSync review feedback --- .../content/api/gpu-device-extensions.mdx | 36 +++++- packages/webgpu/cpp/rnwgpu/ArrayBuffer.h | 20 +++- packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp | 108 ++++++++++++------ packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h | 9 +- .../webgpu/src/__tests__/ReadSync.spec.ts | 30 ++++- packages/webgpu/src/index.tsx | 29 +++-- 6 files changed, 178 insertions(+), 54 deletions(-) diff --git a/apps/docs/content/api/gpu-device-extensions.mdx b/apps/docs/content/api/gpu-device-extensions.mdx index 6bab25583..659ce36a2 100644 --- a/apps/docs/content/api/gpu-device-extensions.mdx +++ b/apps/docs/content/api/gpu-device-extensions.mdx @@ -1,10 +1,44 @@ --- title: Native Extensions -description: Native interop extensions to GPUDevice and the WebGPU globals. +description: React Native extensions to WebGPU objects and globals. --- React Native WebGPU adds a handful of native extensions on top of the WebGPU spec. +## GPUBuffer.readSync + +`readSync()` synchronously reads a small `MAP_READ` buffer into an owned `ArrayBuffer`. It blocks the calling JavaScript or worklet thread until previously submitted GPU work completes, copies the mapped bytes, and unmaps the GPU buffer before returning. + + + Synchronous readback is not faster than `mapAsync()`. It stalls the calling + thread and removes CPU/GPU parallelism. Keep dependent work on the GPU when + possible, and use `mapAsync()` when the result can be consumed in a later + JavaScript turn. Use `readSync()` only when a small CPU result is genuinely + required in the current turn. + + +Create a `MAP_READ | COPY_DST` staging buffer, copy the compute result into it, submit the commands, and read it: + +```tsx +const staging = device.createBuffer({ + size: 16, + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, +}); + +encoder.copyBufferToBuffer(storage, 0, staging, 0, 16); +device.queue.submit([encoder.finish()]); + +const values = new Uint32Array(staging.readSync(0, 16, 500)); +``` + +| Parameter | Default | Description | +| ----------- | --------------------- | ---------------------------------------- | +| `offset` | `0` | Byte offset where the read begins | +| `size` | Remaining buffer size | Number of bytes to copy, capped at 1 MiB | +| `timeoutMs` | `2000` | Maximum time to wait for the GPU | + +`readSync()` throws if the range is invalid, the read exceeds 1 MiB, mapping fails, the timeout elapses, or the Dawn instance does not support timed waits. The last case can occur with an externally provided Dawn instance. + ## importSharedTextureMemory Imports a native buffer (IOSurface / AHardwareBuffer) for zero-copy sampling. diff --git a/packages/webgpu/cpp/rnwgpu/ArrayBuffer.h b/packages/webgpu/cpp/rnwgpu/ArrayBuffer.h index 3e6d38cb7..f177b9216 100644 --- a/packages/webgpu/cpp/rnwgpu/ArrayBuffer.h +++ b/packages/webgpu/cpp/rnwgpu/ArrayBuffer.h @@ -13,17 +13,27 @@ namespace jsi = facebook::jsi; struct ArrayBuffer : jsi::MutableBuffer { ArrayBuffer(void *data, size_t size, size_t bytesPerElement) - : _data(data), _size(size), _bytesPerElement(bytesPerElement) {} + : _data(data), _size(size), _bytesPerElement(bytesPerElement), + _ownsData(false) {} - ~ArrayBuffer() override {} + ArrayBuffer(size_t size, size_t bytesPerElement) + : _ownedData(size == 0 ? nullptr : new uint8_t[size]), + _data(_ownedData.get()), _size(size), _bytesPerElement(bytesPerElement), + _ownsData(true) {} + + ~ArrayBuffer() override = default; size_t size() const override { return _size; } uint8_t *data() override { return static_cast(_data); } + bool ownsData() const { return _ownsData; } + + std::unique_ptr _ownedData; void *_data; size_t _size; size_t _bytesPerElement; + bool _ownsData; }; static std::shared_ptr @@ -110,7 +120,11 @@ template <> struct JSIConverter> { static jsi::Value toJSI(jsi::Runtime &runtime, std::shared_ptr arg) { - return jsi::ArrayBuffer(runtime, arg); + jsi::ArrayBuffer arrayBuffer(runtime, arg); + if (arg->ownsData()) { + arrayBuffer.setExternalMemoryPressure(runtime, arg->size()); + } + return jsi::Value(runtime, arrayBuffer); } }; diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp index 2c2039c7b..8a3d204e5 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp +++ b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp @@ -1,6 +1,8 @@ #include "GPUBuffer.h" +#include #include +#include #include #include @@ -36,19 +38,10 @@ GPUBuffer::getMappedRange(std::optional o, std::optional size) { return array_buffer; } -namespace { -// readSync returns a copy that outlives the mapping, so it needs an -// ArrayBuffer that owns (and frees) its backing store - the base class -// wraps foreign memory and never frees. -struct OwnedArrayBuffer : ArrayBuffer { - explicit OwnedArrayBuffer(size_t size) - : ArrayBuffer(malloc(size), size, 1) {} - ~OwnedArrayBuffer() override { free(_data); } -}; -} // namespace - std::shared_ptr -GPUBuffer::readSync(std::optional o, std::optional sizeIn) { +GPUBuffer::readSync(jsi::Runtime &runtime, std::optional offsetIn, + std::optional sizeIn, + std::optional timeoutMsIn) { // Synchronous small-buffer readback: blocks the calling thread until all // previously submitted GPU work using this buffer completes, then returns // a copy of the mapped bytes. Built for tiny compute results (landmarks, @@ -59,41 +52,88 @@ GPUBuffer::readSync(std::optional o, std::optional sizeIn) { // Instance::WaitAny, which this library's instance enables via the // TimedWaitAny feature at creation; external (Skia-provided) instances // without it fail the wait and throw rather than hang. - size_t offset = o.value_or(0); - size_t size = sizeIn.has_value() - ? sizeIn.value() - : static_cast(_instance.GetSize() - offset); + auto toByteSize = [&runtime](const char *name, double value) -> size_t { + constexpr double kMaxSafeInteger = 9'007'199'254'740'991.0; + if (!std::isfinite(value) || value < 0 || std::floor(value) != value || + value > kMaxSafeInteger || + value > static_cast(std::numeric_limits::max())) { + throw jsi::JSError(runtime, std::string("GPUBuffer.readSync ") + name + + " must be a non-negative safe integer"); + } + return static_cast(value); + }; + + const size_t offset = + offsetIn.has_value() ? toByteSize("offset", *offsetIn) : 0; + const uint64_t bufferSize = _instance.GetSize(); + if (offset > bufferSize) { + throw jsi::JSError(runtime, + "GPUBuffer.readSync offset exceeds the buffer size"); + } + const size_t size = sizeIn.has_value() + ? toByteSize("size", *sizeIn) + : static_cast(bufferSize - offset); + if (size > bufferSize - offset) { + throw jsi::JSError(runtime, + "GPUBuffer.readSync range exceeds the buffer size"); + } + constexpr size_t kMaxReadSyncBytes = 1 << 20; if (size > kMaxReadSyncBytes) { - throw std::runtime_error( - "readSync is intended for small readbacks (<= 1 MiB); use mapAsync " - "for large buffers"); + throw jsi::JSError( + runtime, + "GPUBuffer.readSync is limited to 1 MiB; use mapAsync for larger " + "readbacks"); } - wgpu::MapAsyncStatus mapStatus = wgpu::MapAsyncStatus::Error; - std::string mapMessage = "callback never ran"; + + constexpr double kDefaultTimeoutMs = 2'000.0; + constexpr double kNanosecondsPerMillisecond = 1'000'000.0; + const double timeoutMs = timeoutMsIn.value_or(kDefaultTimeoutMs); + const double maxTimeoutMs = + static_cast(std::numeric_limits::max()) / + kNanosecondsPerMillisecond; + if (!std::isfinite(timeoutMs) || timeoutMs < 0 || timeoutMs > maxTimeoutMs) { + throw jsi::JSError( + runtime, + "GPUBuffer.readSync timeoutMs must be a finite, non-negative number"); + } + const uint64_t timeoutNs = + static_cast(timeoutMs * kNanosecondsPerMillisecond); + + struct MapResult { + wgpu::MapAsyncStatus status = wgpu::MapAsyncStatus::Error; + std::string message = "callback never ran"; + }; + auto mapResult = std::make_shared(); auto future = _instance.MapAsync( wgpu::MapMode::Read, offset, size, wgpu::CallbackMode::WaitAnyOnly, - [&mapStatus, &mapMessage](wgpu::MapAsyncStatus status, - wgpu::StringView message) { - mapStatus = status; - mapMessage = std::string(message); + [mapResult](wgpu::MapAsyncStatus status, wgpu::StringView message) { + mapResult->status = status; + mapResult->message = std::string(message); }); - constexpr uint64_t kTimeoutNs = 2'000'000'000; // 2s: a hung GPU, not a wait - auto waitStatus = _async->instance().WaitAny(future, kTimeoutNs); + auto waitStatus = _async->instance().WaitAny(future, timeoutNs); if (waitStatus != wgpu::WaitStatus::Success) { - throw std::runtime_error( - "readSync: WaitAny did not complete (timeout, or the instance lacks " - "the TimedWaitAny feature)"); + // Cancels the pending map request. The callback owns MapResult so a late + // completion cannot access stack memory after this method returns. + _instance.Unmap(); + throw jsi::JSError( + runtime, + "GPUBuffer.readSync did not complete before timeoutMs, or the Dawn " + "instance does not support timed waits"); } - if (mapStatus != wgpu::MapAsyncStatus::Success) { - throw std::runtime_error("readSync: mapping failed: " + mapMessage); + if (mapResult->status != wgpu::MapAsyncStatus::Success) { + throw jsi::JSError(runtime, "GPUBuffer.readSync mapping failed: " + + mapResult->message); } const void *ptr = _instance.GetConstMappedRange(offset, size); if (ptr == nullptr) { _instance.Unmap(); - throw std::runtime_error("readSync: GetConstMappedRange failed"); + throw jsi::JSError(runtime, + "GPUBuffer.readSync could not access the mapped range"); } - auto result = std::make_shared(size); + // Allocate owned native storage. The JSI converter wraps this memory without + // copying it again and reports its external memory pressure to the runtime. + auto result = std::make_shared(size, 1); memcpy(result->data(), ptr, size); _instance.Unmap(); return result; diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h index d996db160..5a744a150 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h +++ b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h @@ -38,8 +38,10 @@ class GPUBuffer : public NativeObject { std::optional size); std::shared_ptr getMappedRange(std::optional offset, std::optional size); - std::shared_ptr readSync(std::optional offset, - std::optional size); + std::shared_ptr readSync(jsi::Runtime &runtime, + std::optional offset, + std::optional size, + std::optional timeoutMs); void unmap(); void destroy(); @@ -59,7 +61,8 @@ class GPUBuffer : public NativeObject { &GPUBuffer::mapAsync); installMethod(runtime, prototype, "getMappedRange", &GPUBuffer::getMappedRange); - installMethod(runtime, prototype, "readSync", &GPUBuffer::readSync); + installMethodWithRuntime(runtime, prototype, "readSync", + &GPUBuffer::readSync); installMethod(runtime, prototype, "unmap", &GPUBuffer::unmap); installMethod(runtime, prototype, "destroy", &GPUBuffer::destroy); installGetter(runtime, prototype, "size", &GPUBuffer::getSize); diff --git a/packages/webgpu/src/__tests__/ReadSync.spec.ts b/packages/webgpu/src/__tests__/ReadSync.spec.ts index 1e1f4d0bb..d1b73f9bc 100644 --- a/packages/webgpu/src/__tests__/ReadSync.spec.ts +++ b/packages/webgpu/src/__tests__/ReadSync.spec.ts @@ -83,11 +83,39 @@ fn main() { staging.readSync(); return "no error"; } catch (e) { - return e instanceof Error && e.message.includes("small readbacks") + return e instanceof Error && e.message.includes("limited to 1 MiB") ? "capped" : "wrong error"; } }); expect(result).toBe("capped"); }); + it("accepts a configurable timeout", async () => { + const result = await client.eval(({ device }) => { + const staging = device.createBuffer({ + size: 4, + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, + }); + device.queue.writeBuffer(staging, 0, new Uint32Array([42]).buffer); + return new Uint32Array(staging.readSync(undefined, undefined, 5_000))[0]; + }); + expect(result).toBe(42); + }); + it("rejects an invalid timeout", async () => { + const result = await client.eval(({ device }) => { + const staging = device.createBuffer({ + size: 4, + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, + }); + try { + staging.readSync(undefined, undefined, -1); + return "no error"; + } catch (e) { + return e instanceof Error && e.message.includes("timeoutMs") + ? "invalid timeout" + : "wrong error"; + } + }); + expect(result).toBe("invalid timeout"); + }); }); diff --git a/packages/webgpu/src/index.tsx b/packages/webgpu/src/index.tsx index e3711f892..f19f3e88e 100644 --- a/packages/webgpu/src/index.tsx +++ b/packages/webgpu/src/index.tsx @@ -80,19 +80,24 @@ declare global { readonly nativePointer: bigint; } - // Non-spec RN extension: synchronous small-buffer readback. Blocks the - // calling thread until all previously submitted GPU work using this buffer - // completes, then returns an owned copy of the mapped bytes. Built for - // tiny compute results (landmarks, histogram ranges, counters, picking - // ids) that a render/worklet loop must consume in the SAME frame - such - // loops cannot await mapAsync without accepting a frame of staleness. - // Requires MAP_READ usage; the only valid companion is COPY_DST, so the - // pattern is: copy from your storage buffer into this staging buffer, - // submit, readSync(). Capped at 1 MiB (use mapAsync for bulk data); throws - // instead of hanging if the instance lacks TimedWaitAny (external/Skia - // instances) or the GPU hangs past 2s. interface GPUBuffer { - readSync(offset?: number, size?: number): ArrayBuffer; + /** + * Blocks the calling thread until pending GPU work completes, then returns + * a copy of this buffer's mapped bytes. + * + * This React Native extension is intended for results up to 1 MiB that + * must be consumed in the current JavaScript turn. Prefer `mapAsync()` + * whenever a later turn is acceptable because synchronous readback stalls + * the calling thread and reduces CPU/GPU parallelism. + * + * The buffer must have `GPUBufferUsage.MAP_READ` usage. The returned + * `ArrayBuffer` owns its storage and remains valid after this method + * unmaps the GPU buffer. + * + * @throws If the range is invalid, exceeds 1 MiB, mapping fails, the wait + * times out, or the Dawn instance does not support timed waits. + */ + readSync(offset?: number, size?: number, timeoutMs?: number): ArrayBuffer; } interface GPUDevice { From 21eaa0e3f3bcf1f9ed808f05f65bee006582a975 Mon Sep 17 00:00:00 2001 From: Marc Rousavy Date: Fri, 21 Aug 2026 16:15:45 +0200 Subject: [PATCH 3/4] refactor: rename readSync to readbackSync --- .../content/api/gpu-device-extensions.mdx | 10 +++---- packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp | 29 ++++++++++--------- packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h | 12 ++++---- ...{ReadSync.spec.ts => ReadbackSync.spec.ts} | 16 +++++----- packages/webgpu/src/index.tsx | 6 +++- 5 files changed, 41 insertions(+), 32 deletions(-) rename packages/webgpu/src/__tests__/{ReadSync.spec.ts => ReadbackSync.spec.ts} (89%) diff --git a/apps/docs/content/api/gpu-device-extensions.mdx b/apps/docs/content/api/gpu-device-extensions.mdx index 659ce36a2..bd751894c 100644 --- a/apps/docs/content/api/gpu-device-extensions.mdx +++ b/apps/docs/content/api/gpu-device-extensions.mdx @@ -5,15 +5,15 @@ description: React Native extensions to WebGPU objects and globals. React Native WebGPU adds a handful of native extensions on top of the WebGPU spec. -## GPUBuffer.readSync +## GPUBuffer.readbackSync -`readSync()` synchronously reads a small `MAP_READ` buffer into an owned `ArrayBuffer`. It blocks the calling JavaScript or worklet thread until previously submitted GPU work completes, copies the mapped bytes, and unmaps the GPU buffer before returning. +`readbackSync()` synchronously copies a small `MAP_READ` buffer from the GPU into an owned CPU-accessible `ArrayBuffer`. It blocks the calling JavaScript or worklet thread until previously submitted GPU work completes, copies the mapped bytes, and unmaps the GPU buffer before returning. Synchronous readback is not faster than `mapAsync()`. It stalls the calling thread and removes CPU/GPU parallelism. Keep dependent work on the GPU when possible, and use `mapAsync()` when the result can be consumed in a later - JavaScript turn. Use `readSync()` only when a small CPU result is genuinely + JavaScript turn. Use `readbackSync()` only when a small CPU result is genuinely required in the current turn. @@ -28,7 +28,7 @@ const staging = device.createBuffer({ encoder.copyBufferToBuffer(storage, 0, staging, 0, 16); device.queue.submit([encoder.finish()]); -const values = new Uint32Array(staging.readSync(0, 16, 500)); +const values = new Uint32Array(staging.readbackSync(0, 16, 500)); ``` | Parameter | Default | Description | @@ -37,7 +37,7 @@ const values = new Uint32Array(staging.readSync(0, 16, 500)); | `size` | Remaining buffer size | Number of bytes to copy, capped at 1 MiB | | `timeoutMs` | `2000` | Maximum time to wait for the GPU | -`readSync()` throws if the range is invalid, the read exceeds 1 MiB, mapping fails, the timeout elapses, or the Dawn instance does not support timed waits. The last case can occur with an externally provided Dawn instance. +`readbackSync()` throws if the range is invalid, the read exceeds 1 MiB, mapping fails, the timeout elapses, or the Dawn instance does not support timed waits. The last case can occur with an externally provided Dawn instance. ## importSharedTextureMemory diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp index 8a3d204e5..b5655316f 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp +++ b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp @@ -39,9 +39,9 @@ GPUBuffer::getMappedRange(std::optional o, std::optional size) { } std::shared_ptr -GPUBuffer::readSync(jsi::Runtime &runtime, std::optional offsetIn, - std::optional sizeIn, - std::optional timeoutMsIn) { +GPUBuffer::readbackSync(jsi::Runtime &runtime, std::optional offsetIn, + std::optional sizeIn, + std::optional timeoutMsIn) { // Synchronous small-buffer readback: blocks the calling thread until all // previously submitted GPU work using this buffer completes, then returns // a copy of the mapped bytes. Built for tiny compute results (landmarks, @@ -57,7 +57,8 @@ GPUBuffer::readSync(jsi::Runtime &runtime, std::optional offsetIn, if (!std::isfinite(value) || value < 0 || std::floor(value) != value || value > kMaxSafeInteger || value > static_cast(std::numeric_limits::max())) { - throw jsi::JSError(runtime, std::string("GPUBuffer.readSync ") + name + + throw jsi::JSError(runtime, std::string("GPUBuffer.readbackSync ") + + name + " must be a non-negative safe integer"); } return static_cast(value); @@ -68,21 +69,21 @@ GPUBuffer::readSync(jsi::Runtime &runtime, std::optional offsetIn, const uint64_t bufferSize = _instance.GetSize(); if (offset > bufferSize) { throw jsi::JSError(runtime, - "GPUBuffer.readSync offset exceeds the buffer size"); + "GPUBuffer.readbackSync offset exceeds the buffer size"); } const size_t size = sizeIn.has_value() ? toByteSize("size", *sizeIn) : static_cast(bufferSize - offset); if (size > bufferSize - offset) { throw jsi::JSError(runtime, - "GPUBuffer.readSync range exceeds the buffer size"); + "GPUBuffer.readbackSync range exceeds the buffer size"); } - constexpr size_t kMaxReadSyncBytes = 1 << 20; - if (size > kMaxReadSyncBytes) { + constexpr size_t kMaxReadbackSyncBytes = 1 << 20; + if (size > kMaxReadbackSyncBytes) { throw jsi::JSError( runtime, - "GPUBuffer.readSync is limited to 1 MiB; use mapAsync for larger " + "GPUBuffer.readbackSync is limited to 1 MiB; use mapAsync for larger " "readbacks"); } @@ -95,7 +96,8 @@ GPUBuffer::readSync(jsi::Runtime &runtime, std::optional offsetIn, if (!std::isfinite(timeoutMs) || timeoutMs < 0 || timeoutMs > maxTimeoutMs) { throw jsi::JSError( runtime, - "GPUBuffer.readSync timeoutMs must be a finite, non-negative number"); + "GPUBuffer.readbackSync timeoutMs must be a finite, non-negative " + "number"); } const uint64_t timeoutNs = static_cast(timeoutMs * kNanosecondsPerMillisecond); @@ -118,18 +120,19 @@ GPUBuffer::readSync(jsi::Runtime &runtime, std::optional offsetIn, _instance.Unmap(); throw jsi::JSError( runtime, - "GPUBuffer.readSync did not complete before timeoutMs, or the Dawn " + "GPUBuffer.readbackSync did not complete before timeoutMs, or the Dawn " "instance does not support timed waits"); } if (mapResult->status != wgpu::MapAsyncStatus::Success) { - throw jsi::JSError(runtime, "GPUBuffer.readSync mapping failed: " + + throw jsi::JSError(runtime, "GPUBuffer.readbackSync mapping failed: " + mapResult->message); } const void *ptr = _instance.GetConstMappedRange(offset, size); if (ptr == nullptr) { _instance.Unmap(); throw jsi::JSError(runtime, - "GPUBuffer.readSync could not access the mapped range"); + "GPUBuffer.readbackSync could not access the mapped " + "range"); } // Allocate owned native storage. The JSI converter wraps this memory without // copying it again and reports its external memory pressure to the runtime. diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h index 5a744a150..41598e4b9 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h +++ b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h @@ -38,10 +38,10 @@ class GPUBuffer : public NativeObject { std::optional size); std::shared_ptr getMappedRange(std::optional offset, std::optional size); - std::shared_ptr readSync(jsi::Runtime &runtime, - std::optional offset, - std::optional size, - std::optional timeoutMs); + std::shared_ptr readbackSync(jsi::Runtime &runtime, + std::optional offset, + std::optional size, + std::optional timeoutMs); void unmap(); void destroy(); @@ -61,8 +61,8 @@ class GPUBuffer : public NativeObject { &GPUBuffer::mapAsync); installMethod(runtime, prototype, "getMappedRange", &GPUBuffer::getMappedRange); - installMethodWithRuntime(runtime, prototype, "readSync", - &GPUBuffer::readSync); + installMethodWithRuntime(runtime, prototype, "readbackSync", + &GPUBuffer::readbackSync); installMethod(runtime, prototype, "unmap", &GPUBuffer::unmap); installMethod(runtime, prototype, "destroy", &GPUBuffer::destroy); installGetter(runtime, prototype, "size", &GPUBuffer::getSize); diff --git a/packages/webgpu/src/__tests__/ReadSync.spec.ts b/packages/webgpu/src/__tests__/ReadbackSync.spec.ts similarity index 89% rename from packages/webgpu/src/__tests__/ReadSync.spec.ts rename to packages/webgpu/src/__tests__/ReadbackSync.spec.ts index d1b73f9bc..fbded86dc 100644 --- a/packages/webgpu/src/__tests__/ReadSync.spec.ts +++ b/packages/webgpu/src/__tests__/ReadbackSync.spec.ts @@ -1,6 +1,6 @@ import { client } from "./setup"; -describe("readSync", () => { +describe("readbackSync", () => { it("reads back compute results synchronously, same tick", async () => { const result = await client.eval(({ device }) => { const storage = device.createBuffer({ @@ -39,7 +39,7 @@ fn main() { encoder.copyBufferToBuffer(storage, 0, staging, 0, 16); device.queue.submit([encoder.finish()]); // No await between submit and read: the readback is synchronous. - return Array.from(new Uint32Array(staging.readSync())); + return Array.from(new Uint32Array(staging.readbackSync())); }); expect(result).toEqual([1, 2, 3, 42]); }); @@ -54,7 +54,7 @@ fn main() { 0, new Uint32Array([10, 20, 30, 40]).buffer, ); - return Array.from(new Uint32Array(staging.readSync(8, 8))); + return Array.from(new Uint32Array(staging.readbackSync(8, 8))); }); expect(result).toEqual([30, 40]); }); @@ -67,7 +67,7 @@ fn main() { const reads: number[] = []; for (let i = 0; i < 3; i++) { device.queue.writeBuffer(staging, 0, new Uint32Array([i]).buffer); - reads.push(new Uint32Array(staging.readSync())[0]!); + reads.push(new Uint32Array(staging.readbackSync())[0]!); } return reads; }); @@ -80,7 +80,7 @@ fn main() { usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, }); try { - staging.readSync(); + staging.readbackSync(); return "no error"; } catch (e) { return e instanceof Error && e.message.includes("limited to 1 MiB") @@ -97,7 +97,9 @@ fn main() { usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, }); device.queue.writeBuffer(staging, 0, new Uint32Array([42]).buffer); - return new Uint32Array(staging.readSync(undefined, undefined, 5_000))[0]; + return new Uint32Array( + staging.readbackSync(undefined, undefined, 5_000), + )[0]; }); expect(result).toBe(42); }); @@ -108,7 +110,7 @@ fn main() { usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, }); try { - staging.readSync(undefined, undefined, -1); + staging.readbackSync(undefined, undefined, -1); return "no error"; } catch (e) { return e instanceof Error && e.message.includes("timeoutMs") diff --git a/packages/webgpu/src/index.tsx b/packages/webgpu/src/index.tsx index f19f3e88e..a95603645 100644 --- a/packages/webgpu/src/index.tsx +++ b/packages/webgpu/src/index.tsx @@ -97,7 +97,11 @@ declare global { * @throws If the range is invalid, exceeds 1 MiB, mapping fails, the wait * times out, or the Dawn instance does not support timed waits. */ - readSync(offset?: number, size?: number, timeoutMs?: number): ArrayBuffer; + readbackSync( + offset?: number, + size?: number, + timeoutMs?: number, + ): ArrayBuffer; } interface GPUDevice { From 21ccf701a5a4018cc8abe0571c6a5e29aa8220cc Mon Sep 17 00:00:00 2001 From: Marc Rousavy Date: Fri, 21 Aug 2026 16:17:29 +0200 Subject: [PATCH 4/4] fix: remove runtime dependency from readbackSync --- packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp | 36 +++++++++----------- packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h | 6 ++-- 2 files changed, 18 insertions(+), 24 deletions(-) diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp index b5655316f..4e14973b8 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp +++ b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.cpp @@ -4,6 +4,7 @@ #include #include #include +#include #include #include "Convertors.h" @@ -39,7 +40,7 @@ GPUBuffer::getMappedRange(std::optional o, std::optional size) { } std::shared_ptr -GPUBuffer::readbackSync(jsi::Runtime &runtime, std::optional offsetIn, +GPUBuffer::readbackSync(std::optional offsetIn, std::optional sizeIn, std::optional timeoutMsIn) { // Synchronous small-buffer readback: blocks the calling thread until all @@ -52,14 +53,13 @@ GPUBuffer::readbackSync(jsi::Runtime &runtime, std::optional offsetIn, // Instance::WaitAny, which this library's instance enables via the // TimedWaitAny feature at creation; external (Skia-provided) instances // without it fail the wait and throw rather than hang. - auto toByteSize = [&runtime](const char *name, double value) -> size_t { + auto toByteSize = [](const char *name, double value) -> size_t { constexpr double kMaxSafeInteger = 9'007'199'254'740'991.0; if (!std::isfinite(value) || value < 0 || std::floor(value) != value || value > kMaxSafeInteger || value > static_cast(std::numeric_limits::max())) { - throw jsi::JSError(runtime, std::string("GPUBuffer.readbackSync ") + - name + - " must be a non-negative safe integer"); + throw std::runtime_error(std::string("GPUBuffer.readbackSync ") + name + + " must be a non-negative safe integer"); } return static_cast(value); }; @@ -68,21 +68,20 @@ GPUBuffer::readbackSync(jsi::Runtime &runtime, std::optional offsetIn, offsetIn.has_value() ? toByteSize("offset", *offsetIn) : 0; const uint64_t bufferSize = _instance.GetSize(); if (offset > bufferSize) { - throw jsi::JSError(runtime, - "GPUBuffer.readbackSync offset exceeds the buffer size"); + throw std::runtime_error( + "GPUBuffer.readbackSync offset exceeds the buffer size"); } const size_t size = sizeIn.has_value() ? toByteSize("size", *sizeIn) : static_cast(bufferSize - offset); if (size > bufferSize - offset) { - throw jsi::JSError(runtime, - "GPUBuffer.readbackSync range exceeds the buffer size"); + throw std::runtime_error( + "GPUBuffer.readbackSync range exceeds the buffer size"); } constexpr size_t kMaxReadbackSyncBytes = 1 << 20; if (size > kMaxReadbackSyncBytes) { - throw jsi::JSError( - runtime, + throw std::runtime_error( "GPUBuffer.readbackSync is limited to 1 MiB; use mapAsync for larger " "readbacks"); } @@ -94,8 +93,7 @@ GPUBuffer::readbackSync(jsi::Runtime &runtime, std::optional offsetIn, static_cast(std::numeric_limits::max()) / kNanosecondsPerMillisecond; if (!std::isfinite(timeoutMs) || timeoutMs < 0 || timeoutMs > maxTimeoutMs) { - throw jsi::JSError( - runtime, + throw std::runtime_error( "GPUBuffer.readbackSync timeoutMs must be a finite, non-negative " "number"); } @@ -118,21 +116,19 @@ GPUBuffer::readbackSync(jsi::Runtime &runtime, std::optional offsetIn, // Cancels the pending map request. The callback owns MapResult so a late // completion cannot access stack memory after this method returns. _instance.Unmap(); - throw jsi::JSError( - runtime, + throw std::runtime_error( "GPUBuffer.readbackSync did not complete before timeoutMs, or the Dawn " "instance does not support timed waits"); } if (mapResult->status != wgpu::MapAsyncStatus::Success) { - throw jsi::JSError(runtime, "GPUBuffer.readbackSync mapping failed: " + - mapResult->message); + throw std::runtime_error("GPUBuffer.readbackSync mapping failed: " + + mapResult->message); } const void *ptr = _instance.GetConstMappedRange(offset, size); if (ptr == nullptr) { _instance.Unmap(); - throw jsi::JSError(runtime, - "GPUBuffer.readbackSync could not access the mapped " - "range"); + throw std::runtime_error( + "GPUBuffer.readbackSync could not access the mapped range"); } // Allocate owned native storage. The JSI converter wraps this memory without // copying it again and reports its external memory pressure to the runtime. diff --git a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h index 41598e4b9..ede646201 100644 --- a/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h +++ b/packages/webgpu/cpp/rnwgpu/api/GPUBuffer.h @@ -38,8 +38,7 @@ class GPUBuffer : public NativeObject { std::optional size); std::shared_ptr getMappedRange(std::optional offset, std::optional size); - std::shared_ptr readbackSync(jsi::Runtime &runtime, - std::optional offset, + std::shared_ptr readbackSync(std::optional offset, std::optional size, std::optional timeoutMs); void unmap(); @@ -61,8 +60,7 @@ class GPUBuffer : public NativeObject { &GPUBuffer::mapAsync); installMethod(runtime, prototype, "getMappedRange", &GPUBuffer::getMappedRange); - installMethodWithRuntime(runtime, prototype, "readbackSync", - &GPUBuffer::readbackSync); + installMethod(runtime, prototype, "readbackSync", &GPUBuffer::readbackSync); installMethod(runtime, prototype, "unmap", &GPUBuffer::unmap); installMethod(runtime, prototype, "destroy", &GPUBuffer::destroy); installGetter(runtime, prototype, "size", &GPUBuffer::getSize);