GPU throughput: measured bottlenecks and the plan to remove them

Status: accepted, 2026-10-21, as the order of work (owner request: maximum parallel throughput). Each item is verified by its own pull request. No kernel changes in this document's pull request. It adds the measurement script scripts/profile-kernels.mjs and this record. Each item in "The plan" is its own pull request with its own evidence.

Request source: the owner, 2026-10-21 (translated): more memory only lets a stage use more within its limit. What matters more is to accelerate, to process with the maximum possible parallelism, to split work finer where it can be split, and to put data in workgroup memory and the other faster tiers. Moving memory up and down is a bottleneck. Recent papers, techniques from games, other libraries and compression are all open.

Context: pull request #161 (dense MVS residency, docs/dense-mvs.md Amendment 3) and #167 (BA and triangulation streaming) made every stage run within any memory budget. They did not make it faster. This record asks where the time goes on the owner's GPU and what changes it.

Baseline: main at 4b16cd5, compiler pin 3c6e964. TypeShade documents are those of the pin (node_modules/typeshade/AUTHORING.md, docs/use-typeshade-surface.md, changes/).

1. Summary

Facts (section 3):

1. On the owner's RTX 5060 Ti, five stages hold 80 to 90 percent of a run: guided matching, pair verification, bundle adjustment, dense photometric propagation and dense geometric propagation. Triangulation and PnP hold another 11 to 14 percent.
2. The GPU is mostly idle in the sparse stages. They run 6 to 40 ms per dispatch and 0.5 ms per dispatch in pair verification, with 8-lane and 64-lane workgroups of tiny grids. This is latency and launch bound, not arithmetic or bandwidth bound.
3. The BA's PCG iteration costs about 7 ms (newest pin) to 17 ms (older pins). The bytes it moves would take about 25 microseconds at the card's memory bandwidth. The iteration is over 100 times above its bandwidth floor.
4. Chunk timing measures the fence, not the kernel. Many timed chunks end at 3 to 8 ms whatever their size. The learned rate then makes chunks too small, and the GPU never fills.
5. Dense propagation is f32 and holds 97 percent of the dense GPU time on SwiftShader. It rereads the reference window and recomputes a weight and a relative pose for every hypothesis and every source.

Plan (section 7), by measured gain per effort. Estimates are inferences, and each pull request measures its own gain.

Open decisions for the owner are in section 9.

2. Method and configurations

Run records. The owner's collector (Cloudflare KV namespace 1c8702ece02640b7bfdf7dca1af2e6e2, keys record/<ts>-<8hex>) was read on 2026-10-21 with wrangler kv key get. A record holds timing.stages (wall time, calls, dispatches and submissions per stage), timing.chunks, timing.kernels (the kernel names in first-use order, no times), gpu, mainThread, transfers and dense. A record holds no time for a single kernel. Section 4 says what that costs.

Local profile. scripts/profile-kernels.mjs wraps the TypeShade device's fuse() and sums, for each kernel name, the gpuTime that timestamp-query reports. It ran on 2026-10-21 in headless Chromium 1194 with SwiftShader, on a shared 4-core container, at 4b16cd5:

npm run build
node scripts/profile-kernels.mjs --case mvs    --scene south-building-24 --ids 0,1,2,3,4,5 --scale 2
node scripts/profile-kernels.mjs --case bundle --scene south-building-24 --scale 1
node scripts/profile-kernels.mjs --case sift   --scene south-building-24 --ids 0,1,2,3
node scripts/profile-kernels.mjs --case pair   --scene south-building-24 --scale 1

Warning about SwiftShader. It is a CPU emulation. It ranks kernels by arithmetic and by barrier cost, and it has no launch latency, no occupancy and no memory hierarchy worth the name. A SwiftShader share is evidence of where arithmetic goes. It is not evidence of where a physical GPU waits. The earlier comparison in docs/lab-run-records.md ("SwiftShader against the RTX 2080") shows the gap: features ran 111 times faster on the card, pair verification 1.4 times.

Adapter facts. SwiftShader in this container offers the features subgroups, timestamp-query, float32-filterable, indirect-first-instance and the texture-compression set. It does not offer shader-f16. Its limits: 32768 bytes of workgroup storage, 256 invocations per workgroup, 10 storage buffers per stage. The run record keeps neither the features nor the limits of the owner's adapter, so this record cannot say whether the 5060 Ti under Chrome 153 offers shader-f16 and subgroups. Chrome ships subgroups since version 134 (new in WebGPU 134), so the owner's browser is expected to offer it. Item 1 records it.

3. Measured baseline

3.1 Time per stage on the owner's GPU

Seconds, with the share of the run's total. Source: timing.stages and dense.stageMs of the records. Blank means the record has no such stage. The two multi-video records are in 3.4.

The sum of the rows is below the total by 5 to 10 percent: frame decode, upload, mapper control on the host and the viewer.

Observation (fact). The bundle row falls from 36 to 39 s in four records to 11 s in f1edec5c (27 views and 33458 observations, more than 28dcfe16's 24 views and 27062). Between them the pin moved to 3c6e964 (#164, the removal of the typeshade/typeshade#717 stopgap) and #152 cut the host cost of pair verification. Inference: the older BA was bound by the host cost of checking bound arrays before each dispatch, and #164 removed most of it. The records do not isolate the cause. The same inference fits the multi-video pair (3.4).

3.2 Dispatches, submissions and the time of one dispatch

Source: timing.stages. "ms per dispatch" is the stage's wall time divided by its dispatch count. It includes everything: the host, the queue, the fence and the GPU.

Readings (inferences from these counts):

• Guided matching runs 6 dispatches and one submission per pair, yet costs 180 ms. The plain matcher runs the same 6 dispatches for 10 ms. docs/lab-run-records.md shows the scalar guided kernel runs about 3 percent of the plain matcher's sums. So the 170 ms difference is not the match kernels. The stage also runs the guide: a full verifyImagePairGpu (50 hypotheses) on the same correspondences (guidedMatches in website/src/workers/sfm.worker.ts). Pair verification then runs the full 500-hypothesis search on the same pair again.
• Pair verification runs about 195 dispatches and 35 submissions per pair. Its kernels are 8-lane workgroups (preparePairModels, appendPairModelRows, solvePairModelNull, finishPairModels) with one model per lane. A batch of 32 models is 4 workgroups. The batch size is the "32 trials up" floor of the time budget (docs/gpu-geometry.md).
• Triangulation and observations run 9 to 42 ms per dispatch. These are single dispatches that wait for the GPU (submissions equals dispatches). Their cost is the fence plus a per-lane serial solve, 16 lanes to a workgroup.

3.3 The chunk rates show a floor of 3 to 8 ms

timing.chunks of ed932d77, 28dcfe16, 0de918f8, 4a8aa2ff, f1edec5c:

Fact: the longest chunk of nearly every key is 3 to 8 ms, whatever its units. The one exception class is bundleCg (57 to 182 ms), the only key whose chunks do real work per unit.

Inference: a chunk is timed with a queue fence (device.wait()), and a fence on this stack costs about 3 to 5 ms. The learned rate adds that latency to the first timed chunks. The rate is then too high, the planner picks small chunks, and each chunk fills a small part of the GPU. The time budget is 250 ms (timing.budgetMs), but the chunks run at 3 to 8 ms. An RTX 5060 Ti has about 36 streaming multiprocessors and holds about 55000 resident threads. A 118-track chunk of 16-lane workgroups does not fill one tenth of it. For a latency-bound kernel the cost of 118 and of 6341 tracks is nearly the same. Item 2 tests this. The test is a SwiftShader-independent one: the timestamp of the chunk (item 1) against its fence time.

3.4 Multi-video (host-bound)

Fact: the same input ran 3.7 times faster between two site builds that changed no kernel (#152 host-cost cuts, #164 pin move). Pair verification is 79 ms per pair after the change. Inference: for this input the host, not the GPU kernels, held most of the time. The record's mainThread shows frame gaps of 41 s (19f18274), 15 s (fad31ad6) and 33 s with 164 long tasks (d40f8950, a failed 839-frame run). The page stalls even when the worker does the work. This record leaves the page's own stalls to the multi-video record (docs/multi-video.md).

3.5 Transfers

transfers of the records: application uploads 204 to 584 MB per run, readbacks 7 to 11 MB. applicationUploadBytes counts the values given to allocate*() and also the zeros of every new f64 array (ShadeDevice.statistics). Much of the upload is therefore zero fill. At 10 GB/s effective, 584 MB is 60 ms in a 75 s run. Fact: the volume is small in time. Transfer is not a bottleneck in these runs. A GPU-side clear would still save the host copy and the allocation (item 9).

3.6 Per-kernel profile on SwiftShader

Warning: SwiftShader. Read as the distribution of emulated arithmetic.

Dense propagation (--case mvs, 6 views of south-building-24 at 320x240, default iterations, gpuTime sum 315 s):

Bundle adjustment (--case bundle, 24 views, refineSharedFocal, 4 LM iterations, gpuTime sum 9522 ms; 401 PCG iterations):

Pair verification (--case pair, 23 consecutive pairs, gpuTime sum 6120 ms): scorePairModels 55.5%, appendPairModelRows 16.6% (1229 groups, 53 per pair), triangulatePoseCandidates 10.1%, solvePairModelNull 6.6%, the rest 11%.

SIFT (--case sift, 4 frames at 640x480, first octave -1, gpuTime sum 25 s): scanFlags 49.1%, siftBlurV and siftBlurH 21.3%, siftScatterExtrema and siftFlagExtrema 18.6%, siftDescribe 6.8%. scanFlags is a 64-lane Hillis-Steele scan with 12 barriers. SwiftShader pays dearly for barriers, a card does not. Do not act on this row.

3.7 What the BA numbers say about the card

bundleCg learned 0.0007 ms per unit on ed932d77 and 0.0002 on f1edec5c. A unit is one row of a PCG iteration (M + 6C, docs/gpu-sfm.md). For f1edec5c that is 0.0002 x (33458 + 162) = 6.7 ms per PCG iteration, and for ed932d77 0.0007 x 24312 = 17 ms.

The data one iteration reads, assuming 8 bytes per emulated f64 and 21 stored cross-block words per observation (crossAt in src/gpu/joint-focal.shade.ts): 33458 x 168 B = 5.6 MB per pass, two passes, about 12 MB with the vectors. At about 450 GB/s that is 25 microseconds. The arithmetic is about 33458 x 21 MACs x 2 passes x 20 f32 operations per emulated MAC = 28 MFLOP, below 2 microseconds at 20 TFLOP/s. Inference: the iteration is more than 100 times above its roofline, and the cause is the critical path, not the volume. In bundleSchurCameraRow, one invocation walks all observations of its camera, about 1200 for f1edec5c, with four dependent loads per step (cameraObservation, index, blocks, vectors), and only 6C + 1 invocations exist.

4. What the records cannot tell, and how to close it

1. No kernel time. timing.kernels lists names only. Item 1 adds gpuMs, groups and dispatches per kernel, summed from FuseReport.groups[].gpuTime, and the wall time of the plans that held them. The difference of a stage's wall time and its summed GPU time is its host and fence time. This one table decides the order of items 2 to 8 on the owner's card.
2. No adapter features. gpu holds two limits. Item 1 adds features and the compute limits (maxComputeWorkgroupStorageSize, maxComputeInvocationsPerWorkgroup, maxStorageBuffersPerShaderStage), and shader-f16 and subgroups in particular.
3. mvsPropagate does not separate the photometric and geometric pass in dense.stageMs's kernels. The stage times separate them (3.1). Item 1 keeps the kernel's pass in the key.
4. Guided matching has no split between the guide and the match kernels. Item 1 times them.

5. Research

Each row: the method, the source, what it removes here and the fit to WebGPU. The WebGPU limits that decide fit: no native f64; subgroups is optional and absent from many adapters; 16 KiB of workgroup memory is the guaranteed minimum (maxComputeWorkgroupStorageSize), 32 KiB is common; 8 storage buffers per stage is the guaranteed minimum; no device-wide barrier inside a dispatch; timestamp-query is optional; a forward-progress guarantee between workgroups does not exist.

5.1 Bundle adjustment

5.2 Features and matching

5.3 Dense MVS

5.4 Game-engine and library techniques

6. TypeShade capabilities for each technique

Source: node_modules/typeshade at the pin (AUTHORING.md, docs/use-typeshade-surface.md, changes/). Status "available" means the surface documents it. Nothing here was run on a card.

Draft issues for typeshade/typeshade (not filed)

G1. Title: An f16 scalar type for "use typeshade" kernels, so that half-precision storage and arithmetic are authorable. Body: AUTHORING.md lists f16 as a capability (enable f16;, host feature shader-f16) and says no f16 scalar type exists. TypeCV's dense MVS and BA store gray, depth, cost and Jacobian blocks in f32. It can pack two halves with pack2x16float but cannot compute on them. Expected: f16 scalar and vec2h, vec4h types, with the CPU oracle rounding to binary16 after each operation (as quantizeToF16 does) and the WebGL2 backend refusing them. Tests owed: a kernel with f16 storage and f32 accumulation against the oracle; the SD0030 refusal on GLSL. Pinned commit: 3c6e964.

G2. Title: Subgroup intrinsics (subgroupAdd, subgroupMax, subgroupShuffle, subgroupBallot) with a CPU oracle meaning. Body: The pin accepts @builtin("subgroup_invocation_id") and @builtin("subgroup_size") and derives the subgroups capability, but no intrinsic. TypeCV's reductions (bundleMetricPartials, the PCG dot products, the top-2 selection of matching) use a workgroup-memory tree with several barriers. A subgroup reduction needs none. Expected: the intrinsics behind the subgroups capability, with the oracle running a declared subgroup size of 1 or a stated simulated size, and a refusal on WebGL2. Tests owed: a reduction against the tree reduction on SwiftShader (which offers subgroups).

G3. Title: A sampled texture with a filtering sampler read from a compute entry, with a stated tolerance class for the CPU oracle. Body: The surface table lists textureSampleLevel as valid in any stage. TypeCV would store the gray image of each view as an r16float or r8unorm texture and read bilinear samples in mvsPropagate. Hardware bilinear weights have fewer than 32 bits of fraction, so the CPU oracle cannot match bit for bit. Expected: documented behavior of textureSampleLevel in a @compute entry on WebGPU and the CPU tier, and a tolerance class in the surface's accuracy table for the filtered result (a bound in units of the texel's last place). Tests owed: a compute entry that samples at fractional coordinates, compared with a four-texel oracle within the bound.

G4. Title: The CPU tier runs a tiled entry (workgroup memory and barriers) as a whole workgroup. Body: Surface section 25 says the oracle runs one invocation per call, so a barrier has no partner. TypeCV keeps a sequential ...Oracle sibling for each tiled entry (scanPairRowsOracle, bundleMetricPartialsOracle, scorePairModelsOracle). Each pair is a second kernel to write and to hold to the first. Expected: a CPU tier that runs the lanes of a workgroup between barriers (a generator or a thread pool), so that one source runs on every tier, as change 0069 allows for a tiled entry in a fuse scope. Tests owed: the existing tile examples on the CPU tier.

G5. Title: The chunk learner of a device fits time = latency + rate x units from GPU timestamps when the adapter has timestamp-query. Body: Change 0074 times a chunk from its plan start to the queue fence. On the owner's RTX 5060 Ti the longest chunk of nearly every key is 3 to 8 ms whatever its size (TypeCV record ed932d77: scorePairModels 1 unit in 3.0 ms, scorePnP 500 units in 4.1 ms, triangulateTracks 118 units in 5.4 ms). The fence latency enters the rate, the rate sizes small chunks, and small chunks leave the GPU empty. Expected: when FuseReport.groups[].gpuTime exists, learn the rate from it and keep the fence time only as a watchdog check. Without timestamps, fit two chunk sizes to separate the latency. Also expose a per-entry gpuTime for the entries of one plan. Tests owed: a synthetic kernel with a known latency floor, and a check that the chosen chunk grows to the budget.

G6. Title: A host-readable workgroup-memory limit and a way to size a tile from it. Body: A tile kernel must fit maxComputeWorkgroupStorageSize (16 KiB guaranteed). The device owner exposes limits (change 0051). Expected: a documented read of that limit in the program runtime, and an override constant that a kernel can size its tile from (change 0057 covers the workgroup size). Tests owed: one kernel built at two tile sizes.

7. The plan

Order: measured gain per effort, on the owner's card. Gains are inferences until the pull request measures them. Each pull request states its before and after numbers: SwiftShader with scripts/profile-kernels.mjs for the work done (dispatches, groups, GPU ms), and, by request, the owner's GPU through a run record (item 1 makes that table).

Verification classes:

• Exact. The pull request shows equal output on the CPU tier and on SwiftShader (the tests named in each item). A pull request in this class changes order of loop iterations only where it says so.
• Tolerance. A kernel reorders a floating-point sum or reads a filtered texel. The pull request states the bound, shows the CPU oracle held to it, and shows the end-to-end figures (landmarks, mean reprojection error, fused points) within the run-to-run spread of the baseline.

Item 1. Per-kernel GPU time and adapter features in the run record

• Stage: all. Technique: measurement. Cost: small (website/src/lib/run-record.ts, website/src/workers/sfm.worker.ts, docs/lab-run-records.md, record test).
• Change: record format v6. timing.gpu is a list of { name, groups, dispatches, gpuMs } from the summed FuseReport.groups (a wrapper on the TypeShade device's report, as in scripts/profile-kernels.mjs), plus timing.stages[*].gpuMs. gpu.features holds the adapter's feature list and gpu.limits the compute limits. The cap is 120 kernel rows.
• Gain: none directly. It turns sections 3.2 and 3.3 from inference into fact on the owner's card and decides items 2 to 8.
• Risk: timestamp-query rounds timestamps coarsely in some browsers; a group may carry no gpuTime. A kernel row keeps gpuMs: null then. Timestamp queries cost a little per plan.
• Verification: record validator test (website/src/lib/run-record.ts tests), a record sent from SwiftShader carries the rows, v5 records still validate. Exact class (no kernel changes).
• TypeShade gaps: none (G5 later improves it).

Item 2. Fill the device before splitting a chunk

• Stage: triangulation (5.4 s), PnP (5.0 s), pair verification (16 s, scoring and model batches), and the same rule for observations. Technique: occupancy-aware chunks. A chunk key gets a minimum number of units that fills the device (for example 2048 workgroups), and the rate learned from GPU time (item 1) replaces the fence time as soon as the pin offers it (G5). Until G5, TypeCV computes the rate itself from gpuTime of the chunk's report.
• Cost: medium-small. Files: src/gpu/device.ts (planSteps), src/gpu/triangulation.ts, src/gpu/pnp.ts, src/gpu/pair-estimation.ts (the trials expression, 32 floor), src/gpu/observations.ts.
• Expected gain: triangulation 5.4 s to about 1 s (a call of 6341 tracks as one 400-workgroup dispatch, because a latency-bound serial solve costs the same for 118 and 6341 tracks until the card fills); PnP 5 s to about 2 s; pair verification 16 s to about 9 s (a batch of 256 to 512 models instead of 32 cuts the 195 dispatches per pair to about 60). Sum 8 to 12 s of 75 s.
• Risk: the Windows watchdog (TDR) and the earlier Blackwell hang in triangulateTracks (docs/hardware-observations.md, A1). Mitigation: a ceiling in measured GPU milliseconds (for example 100 ms) per dispatch, taken from gpuTime, and an upper bound of the first dispatch of each key at the old chunk size times four. The idle-lane rule stays.
• Verification: exact. tests/typeshade/pnp-batch.test.mjs, pair-batch.test.mjs, tests/typeshade/streaming.test.mjs, plan-steps.test.mjs pass unchanged; the RANSAC search stops "where a sequential RANSAC would, whatever the batch" (docs/gpu-sfm.md), so the models and inliers are equal. The pull request adds a test that a call of N tracks plans at most ceil(N / floor) dispatches. Benchmark: dispatch and submission counts per stage on SwiftShader (must fall), GPU ms from the profile script, and the owner's run record (timing.stages).
• TypeShade gaps: G5 (rate from GPU time). Not blocking.

Item 3. The guide of guided matching estimates E only

• Stage: guided matching (17.6 s, 23 percent). Technique: the guide needs a band, not a pose. Run the essential-matrix RANSAC alone (estimateEssentialRansacGpu, src/gpu/two-view.ts) in one batch of 64 to 128 hypotheses and build the band from E directly. Skip pose recovery, triangulation and refinement (triangulatePoseCandidates is 10 percent of pair verification GPU time on SwiftShader) and the essentialFromPose round trip.
• Cost: small to medium. Files: website/src/workers/sfm.worker.ts (guidedMatches), website/src/lib/guided-matching.ts, maybe src/gpu/two-view.ts.
• Expected gain: the guide drops from about 170 ms to about 40 to 60 ms per pair: 17.6 s to 6 to 8 s. A further step, reuse of the mapper's E (verify first, then rematch), would remove the guide entirely but reorders the worker. Record it as a follow-up.
• Risk: the band differs slightly (E unscaled against [t]x R). The mapper still verifies the chosen matches, so a wrong band costs matches, not correctness.
• Verification: tolerance. tests/typeshade/guided-matching-webgpu.test.mjs and the end-to-end figures. The pull request shows the guided match count per pair on the pyramid and south-building samples within 2 percent of the baseline, and equal landmarks within the baseline spread.
• TypeShade gaps: none.

Item 4. Dense propagation: hoist the loop-invariant work, then pack the gray

• Stage: dense photometric (7 to 26 s) and geometric (3 to 5 s). On SwiftShader mvsPropagate is 97 percent of the dense GPU time. Technique, in two exact steps and one tolerance step:
  1. Reference window cache (exact). referenceWindow and sourceCost read the reference window again for each of up to 10 hypotheses and 6 sources per pixel and compute an exp bilateral weight for each of 36 samples each time. Read the 36 reference values and weights once into array<f32, 36> registers per pixel and iteration. This removes about 2000 global loads and 2000 exp per pixel and iteration.
  2. Relative pose cache (exact). sourceCost rebuilds R_rs and t_rs (about 100 flops and 12 loads) for each hypothesis and source. They depend on the reference and source only. Compute them once per pass into a small storage array (mvsRelative, 6 x 16 floats), the way mvsCameras is held.
  3. Gray as unorm8 (tolerance). The gray comes from 8-bit RGB. Store the undistorted gray with pack4x8unorm (4 samples per word). Source bytes fall 4 times. Quantization error is at most 0.5/255 of the range per sample and enters NCC as noise below the cost's own noise.
• Cost: step 1 and 2 small (src/gpu/mvs.shade.ts, src/gpu/mvs.ts); step 3 medium (mvsUndistort, mvsGray layout in docs/dense-mvs.md 4.1 and the residency plan of Amendment 3).
• Expected gain: steps 1 and 2: 20 to 30 percent of the propagate time (a count of the loads and transcendental calls, to be measured on SwiftShader first); step 3: a further 10 to 20 percent where sampling is bandwidth bound. Together 25 to 40 percent of the dense time (4 to 12 s per run).
• Risk: more registers per lane (36 values and 36 weights), which can lower occupancy. The pull request measures register-bound effects with a second variant (weights recomputed from 36 offsets, values cached).
• Verification: steps 1 and 2 exact: tests/typeshade/mvs.test.mjs, mvs-geometric.test.mjs and mvs-webgpu.test.mjs compare depth and normal bit for bit; the fused point count equals the baseline. Step 3 tolerance: the fused point count and the 99th percentile depth change within 1 percent on south-building-24 and the pyramid sample. Benchmark: mvsPropagate GPU ms before and after with the profile script, same views and iterations.
• TypeShade gaps: none for steps 1 and 2. unorm8 uses pack4x8unorm (available).

Item 5. BA: one workgroup per camera in the Schur camera product

• Stage: bundle (11 to 39 s). Technique: parallelize over observations, as MegBA does, and reduce in workgroup memory. Today bundleSchurCamera runs 6C + 1 invocations, each walking its camera's whole observation list with four dependent loads per step. New: one workgroup of 64 lanes per camera, lanes stride over that camera's observations, each lane keeps six f64 partial sums, and a tree reduction in workgroup<array<f64, 64>> finishes the six rows (512 B per row, 3 KiB for six rows, well inside 16 KiB). Sort the camera's observation index so lane loads are contiguous. Keep the existing kernel as the CPU tier's Oracle.
• Cost: medium. Files: src/gpu/joint-focal.shade.ts (bundleSchurCamera, the cached variants), src/gpu/bundle.shade.ts (fixed-focal route), src/gpu/joint-focal.ts, src/gpu/bundle.ts, the oracle siblings, tests. The camera count per workgroup stays small, so grid size is C workgroups: at 27 cameras still fewer than a card's SMs. Item 6 is the answer to that; this item cuts the critical path from about 1200 serial steps to about 20.
• Expected gain: the critical path of the camera product falls about 20 to 50 times. The PCG iteration is the camera product plus the point product plus launches, so the iteration falls 1.5 to 3 times (from 6.7 ms toward 2 to 4 ms). BA stage: 11 s to 4 to 7 s on the newest pin. SwiftShader cannot show this gain (it has no latency); the owner's record decides.
• Risk: reduction order changes the f64 sum, so results differ in the last bits. LM acceptance is a threshold test and tolerates this; the CPU oracle keeps the sequential order. Barriers need uniform control flow (TS8052): every lane runs every step, with masked loads.
• Verification: tolerance, relative 1e-9 on the PCG vectors per iteration against the oracle and equal LM iteration counts on the fixtures of tests/typeshade/bundle-checkpoint.test.mjs, solver-checkpoint.test.mjs, schur-cache.test.mjs, bundle-camera-webgpu.test.mjs; final cost and mean reprojection error within 1e-6 relative. Benchmark: profile-kernels.mjs --case bundle (dispatch counts equal; GPU ms for the kernel may not fall on SwiftShader, state it) and the owner's bundleCg ms per unit (timing.chunks), which must fall.
• TypeShade gaps: G4 (a CPU twin per tiled kernel). Subgroup intrinsics G2 would shorten the tree.

Item 6. BA: explicit reduced camera system and tiled Cholesky

• Stage: bundle. Technique: assemble the 6C x 6C reduced camera system once per LM iteration (one workgroup per camera pair that shares landmarks, accumulating over their common landmarks in a fixed order, no atomics) and factor it with a blocked Cholesky in workgroup memory (Ceres DENSE_SCHUR on a GPU). At C = 27 the matrix is 162 x 162, 26244 entries, about 210 KB in f64 pairs. This is the schurCache assembly (cacheBundleSchurCross, cacheBundleSchurRows), lifted from its 16-camera guard and followed by a direct solve in place of PCG.
• Cost: large. Files: src/gpu/joint-focal.shade.ts, src/gpu/joint-focal.ts, a new src/gpu/cholesky.shade.ts (blocked, 16 x 16 tiles, 2 KiB each in f64), tests, and a record in docs/gpu-geometry.md. A direct solve also needs a damping retry loop for LM.
• Expected gain: 20 to 100 PCG iterations (60 to 300 dispatches) per LM step become about 10 dispatches. A Cholesky of 162 costs 1.4 M emulated MACs, tens of microseconds of arithmetic on a card with a latency-bound chain of 162 column steps, so about 1 to 3 ms. The recorded schurCache result (RTX 2080, 2.4 times on the final BA, docs/gpu-sfm.md) is the lower evidence. Estimate 3 to 8 times on the BA stage. The solver also stops depending on the preconditioner and on maxCgIterations.
• Risk: memory grows as (6C)^2 (a 300-camera problem is 25 MB in f64 pairs, fine; 1000 cameras is 288 MB, so keep PCG above a camera limit chosen from device.workingBytes). Cholesky in emulated f64 needs a pivot guard for gauge-free directions.
• Verification: tolerance. The step equals the PCG step to the PCG tolerance (cgRelativeTolerance) on every BA fixture; LM trajectories compared per iteration on the CPU tier; end-to-end figures as item 5. Benchmark: dispatch count per BA call (must fall 5 to 10 times), GPU ms on SwiftShader, and the owner's bundle stage.
• TypeShade gaps: G4 (CPU twin), G6 (tile size from the limit). Subgroups (G2) are optional.

Item 7. Hardware bilinear lookups for the source images

• Stage: dense propagation. Technique: upload each undistorted gray as an r16float or r8unorm texture; sourceAt becomes one textureSampleLevel. This keeps the 2D locality of the warped window in the texture cache, and the four loads and three mixes become one instruction. Needs the device owner to make textures (change 0077 for a kernel-written texture; an uploaded texture is available).
• Cost: medium. Files: src/gpu/mvs.shade.ts, src/gpu/mvs.ts, src/gpu/mvs-residency.ts (bytes per view change), docs/dense-mvs.md.
• Expected gain: 1.5 to 2.5 times on propagation where it is texture-bound. Unproven: it depends on the card (item 1 tells whether propagation is bandwidth or ALU bound).
• Risk: the CPU oracle's bilinear differs from the filter hardware (weights are 8-bit fractions). The "no data" sentinel (-1) cannot ride through filtering: the validity check moves to a mask texture or to textureLoad at the four corners.
• Verification: tolerance (G3 sets the class). Fused points within 1 percent, depth median within 0.1 percent on the two scenes.
• TypeShade gaps: G3.

Item 8. Multi-resolution propagation (ACMM style)

• Stage: dense photometric. Technique: run 2 to 3 iterations at 1/4 and 1/2 resolution first, upsample the depth and normal as the full-resolution initialization, and run 2 to 3 iterations at full resolution instead of 8. Estimated 2 to 3 times on the photometric stage.
• Cost: large (src/gpu/mvs.ts plan, mvs-residency.ts, a design amendment to docs/dense-mvs.md 3.5 and 4.x). Result changes; needs a quality gate (ETH3D-style accuracy or the repository's TUM and south-building comparisons).
• Verification: tolerance with a quality gate against scripts/compare-mvs-tum.mjs. Own design amendment first.
• TypeShade gaps: none.

Items later or not worth it now

Item 9. Small items for the same pull request series

• Clear new f64 arrays on the GPU instead of uploading zeros (ShadeDevice.allocate*).
• Add the pass (photometric or geometric) to the kernel key of item 1.

8. The first three pull requests

PR 1. Run record v6: GPU time per kernel, adapter features, the profile script

1. Move scripts/profile-kernels.mjs into the pull request (it exists in the branch of this record). Keep its wrapper of device.shade.fuse as the model.
2. In src/gpu/device.ts, add a small collector in ShadeDevice.plan: for each report, add groups and gpuTime to a map keyed by calls.join('+'). Expose it as statistics.kernels. Do not keep more than 160 keys.
3. In website/src/workers/sfm.worker.ts, add the collector's rows to the record's timing.gpu and a per-stage gpuMs. website/src/lib/run-record.ts validates them (v6; v5 stays valid).
4. In website/src/lib/send-record.ts and the worker, add adapter.features and the three compute limits to gpu.
5. Update docs/lab-run-records.md ("What a record holds").
6. Tests: the validator accepts a v6 record and a v5 record, and rejects a negative gpuMs; website/worker.test.mjs sees the rows.
7. Ask the owner for one run of the pyramid sample with Dense MVS on the 5060 Ti. The table of section 3.2 then gets a GPU-ms column.

Checks: npm run format:check, npm run typecheck, npm test, npm run test:website, npm run test:shade.

PR 2. Fill the device before splitting a chunk (item 2)

1. Add a minUnits option to GpuChunk in src/gpu/device.ts and use it in ShadeDevice.chunkSize so that a key never returns less than minUnits, unless a measured dispatch exceeded the ceiling.
2. Set minUnits from workgroups: triangulation 16 lanes x 1024 workgroups, scorePnP 512 trials, scorePairModels 256 trials, observations all units. Replace the Math.min(batchSize, 32) floor in src/gpu/pair-estimation.ts and the equal floor in src/gpu/pnp.ts.
3. Teach ShadeDevice to read the chunk's gpuTime and to learn the rate from it when it exists (device.shade.chunkRates stays the fence-based fallback). Add a ceiling of 100 GPU ms per dispatch that halves the key's minUnits after one overrun.
4. Tests: a planner test that N tracks give ceil(N / floor) dispatches; the existing batch tests unchanged.
5. Measure: dispatches and submissions per stage on SwiftShader through scripts/benchmark-sfm-pyramid-e2e.mjs; the owner's record against f1edec5c.

PR 3. The guide estimates E only (item 3)

1. In website/src/workers/sfm.worker.ts guidedMatches, replace verifyImagePairGpu with estimateEssentialRansacGpu at 64 to 128 hypotheses (read src/gpu/two-view.ts for the options), and build E from its result.
2. Remove the pose round trip (essentialFromPose) and guideVerificationOptions's unused options; keep GUIDE_RANSAC_ITERATIONS as the hypothesis count.
3. Tests: website/guided-matching.test.mjs for the new helpers; tests/typeshade/guided-matching-webgpu.test.mjs unchanged.
4. Measure: guided stage time on SwiftShader through the pyramid benchmark, the guided match count per pair against the baseline (2 percent), landmarks and mean error against the baseline spread, and the owner's guidedMatches stage against f1edec5c.

After these, order: item 4 (steps 1 and 2 are exact and small), item 5, item 6, then 7 and 8. If the owner's kernel table (PR 1) shows dense propagation or BA holds the GPU time, move items 4 to 6 before PR 2 and PR 3.

9. Open items

10. References

• Wu, Agarwal, Curless, Seitz. Multicore Bundle Adjustment. CVPR 2011. https://grail.cs.washington.edu/projects/mcba/
• Ren, Liang, Yan, Mai, Liu, Liu. MegBA: A GPU-Based Distributed Library for Large-Scale Bundle Adjustment. ECCV 2022. https://arxiv.org/abs/2112.01349
• Demmel, Sommer, Cremers, Usenko. Square Root Bundle Adjustment for Large-Scale Reconstruction. CVPR 2021. https://arxiv.org/abs/2103.01843
• Weber, Demmel, Chan, Cremers. Power Bundle Adjustment for Large-Scale 3D Reconstruction. arXiv 2022 (v4 2023). https://arxiv.org/abs/2204.12834
• Agarwal, Snavely, Seitz, Szeliski. Bundle Adjustment in the Large. ECCV 2010. https://grail.cs.washington.edu/projects/bal/
• Haidar, Tomov, Dongarra, Higham. Harnessing GPU Tensor Cores for Fast FP16 Arithmetic to Speed up Mixed-Precision Iterative Refinement Solvers. SC 2018.
• Bleyer, Rhemann, Rother. PatchMatch Stereo. BMVC 2011.
• Galliani, Lasinger, Schindler. Massively Parallel Multiview Stereopsis by Surface Normal Diffusion. ICCV 2015.
• Xu, Tao. Multi-Scale Geometric Consistency Guided Multi-View Stereo. CVPR 2019. https://arxiv.org/abs/1904.08103
• Xu, Tao. Planar Prior Assisted PatchMatch Multi-View Stereo. AAAI 2020. https://arxiv.org/abs/1912.11744
• Schonberger, Zheng, Pollefeys, Frahm. Pixelwise View Selection for Unstructured Multi-View Stereo. ECCV 2016.
• Schonberger, Frahm. Structure-from-Motion Revisited. CVPR 2016.
• Johnson, Douze, Jegou. Billion-scale similarity search with GPUs. IEEE Transactions on Big Data 2021. https://arxiv.org/abs/1702.08734
• Arandjelovic, Zisserman. Three things everyone should know to improve object retrieval. CVPR 2012.
• AliceVision PopSift. https://github.com/alicevision/popsift
• Cigolle, Donow, Evangelakos, Mara, McGuire, Meyer. A Survey of Efficient Representations for Independent Unit Vectors. JCGT 2014. https://jcgt.org/published/0003/02/01/
• Kerbl, Kopanas, Leimkuehler, Drettakis. 3D Gaussian Splatting for Real-Time Radiance Field Rendering. SIGGRAPH 2023. https://arxiv.org/abs/2308.04079
• Ye et al. gsplat: An Open-Source Library for Gaussian Splatting. https://arxiv.org/abs/2409.06765
• Adinets, Merrill. Onesweep: A Faster Least Significant Digit Radix Sort for GPUs. https://arxiv.org/abs/2206.01784
• Merrill, Garland. Single-pass Parallel Prefix Scan with Decoupled Look-back. NVIDIA TR NVR-2016-002.
• Haar, Aaltonen. GPU-Driven Rendering Pipelines. SIGGRAPH 2015 Advances in Real-Time Rendering.
• Barrett. Sparse Virtual Textures. GDC 2008.
• Harris. Optimizing Parallel Reduction in CUDA. NVIDIA 2007.
• Lewis. Fast Normalized Cross-Correlation. Vision Interface 1995.
• Chrome for Developers. What's new in WebGPU (Chrome 134), subgroups. https://developer.chrome.com/blog/new-in-webgpu-134

Citations were checked for title, authors and venue only where a link is given. The others are cited from the authors' publications and need a check before a pull request relies on them.