Skip to content

[Caspar] Fix FlushSumShared warp reduction and zero-initialize reduction accumulators - #472

Open
sarlinpe wants to merge 3 commits into
symforce-org:mainfrom
sarlinpe:caspar-fix-flush-sum-shared
Open

[Caspar] Fix FlushSumShared warp reduction and zero-initialize reduction accumulators#472
sarlinpe wants to merge 3 commits into
symforce-org:mainfrom
sarlinpe:caspar-fix-flush-sum-shared

Conversation

@sarlinpe

@sarlinpe sarlinpe commented Oct 2, 2026

Copy link
Copy Markdown

Summary

Follows up on #469 (a768308e) with two related fixes in symforce/caspar/source/runtime/memops.cuh when compiling Caspar kernels with Clang CUDA:

  1. Zero-initialize StorageT tot = 0; in FlushSumShared, FlushSumBlock, and FlushSumBlockAdd:
    When a thread is inactive (idx.argsort == 0xffff or !valid), tot was previously uninitialized across __syncthreads() barriers, triggering -Wsometimes-uninitialized under Clang and creating an undef/poison phi node in LLVM IR.
  2. Replace cg::labeled_partition + cg::reduce in FlushSumShared with an explicit warp-shuffle reduction:
    cg::labeled_partition(cg::coalesced_threads(), unique) was called outside the if (idx.argsort != 0xffff) guard while cg::reduce(group, ...) was called inside it. For float and double, cg::reduce over a coalesced_group falls back to cg::details::inclusive_scan_non_contiguous, which runs a warp-shuffle loop (__shfl_sync(mask, ...)) whose trip count (groupSize) varies per partition within the same warp. Under Clang CUDA, this divergent shuffle loop produces corrupted reductions when multiple factors in a warp share a target node. Replacing it with a uniform warp-shuffle reduction over __activemask() ensures all active threads execute __shfl_sync in lockstep.
  3. Add a test case in caspar_solver_cuda_test.cc:
    Existing solver tests used kNumNodes = 4096 (an exact multiple of the 1024-thread block size, leaving no inactive lanes) with a 1-to-1 factor-to-node mapping (indices[i] = i, so all partitions in FlushSumShared had size 1). Added a test case with shared nodes (kNumSharedNodes = 64) and a partial block (kNumFactors = 1500).

@sarlinpe
[Caspar] Fix FlushSumShared warp reduction and zero-initialize reduct…
ab2089b 
…ion accumulators

In FlushSumShared, cg::labeled_partition(cg::coalesced_threads(), unique)
followed by cg::reduce inside the divergent (idx.argsort != 0xffff) branch
can produce incorrect warp reductions when compiled with Clang CUDA
(where floating-point coalesced_group reductions fall back to
inclusive_scan_non_contiguous with per-group variable shuffle loop trip
counts, and tot is uninitialized across __syncthreads() in inactive lanes).

Replace the labeled_partition + cg::reduce call in FlushSumShared with an
explicit uniform warp-shuffle reduction over __activemask(), and
zero-initialize tot in FlushSumShared, FlushSumBlock, and FlushSumBlockAdd
(following up on a768308).

Also add a test case in caspar_solver_cuda_test covering shared nodes and
partial blocks.

Signed-off-by: Paul-Edouard Sarlin <15985472+sarlinpe@users.noreply.github.com>
@sarlinpe sarlinpe mentioned this pull request Oct 2, 2026
Merged
ahojnnes pushed a commit to colmap/colmap that referenced this pull request Oct 7, 2026
@sarlinpe
Regenerate Caspar kernels with upstream SymForce fixes (#4799)
68b722b 
## Summary

Regenerates `src/thirdparty/Symforce-Caspar/generated/{f32,f64}` using
`caspar_generate.py` with upstream SymForce (`symforce-org/symforce`)
and the following upstream PRs:

- **Fix `FlushSumShared` warp reduction and zero-initialize `tot` in
`memops.cuh`** (symforce-org/symforce#472):
Replaces `cg::labeled_partition` + `cg::reduce` inside the divergent
`idx.argsort != 0xffff` branch with an explicit warp-shuffle reduction
over `__activemask()`, and zero-initializes `StorageT tot = 0` in
`FlushSumShared`, `FlushSumBlock`, and `FlushSumBlockAdd` (fixing
corrupted warp reductions and uninitialized-variable UB under Clang
CUDA).
- **Rematerialize numeric constants at point of use**
(symforce-org/symforce#473): Avoids caching
numeric literals globally in registers across the entire kernel.
- **Scope subexpression elimination in Dabseg and CPSE to output
parameter groups** (symforce-org/symforce#475):
Scopes CSE and partial CSE reductions to individual parameter groups
(`res`, `pose`, `calib`, `point`). This prevents intermediate variables
(e.g. from the pose Jacobian) from staying live across multiple
`FlushSumShared` reduction barriers and spilling/corrupting stack
registers in subsequent calibration Jacobian evaluations.
- **Zero-initialize kernel registers (`r0 = 0, r1 = 0, ...`)**
(`symforce-org/symforce@a768308e` / `symforce-org/symforce#469`):
Prevents uninitialized register values in inactive padding lanes of
partial blocks from poisoning warp-level reductions in `SumStore`.
- **Fix solver loop defects and populate `initial_score`**
(`symforce-org/symforce@eb3d04a7` / `symforce-org/symforce#468`,
`symforce-org/symforce@475065da`, `symforce-org/symforce@b3d30737`).
- **Adjust f64 tolerance in `MergedCalibMatchesCeres`**: Relaxed `kPPTol
= 2.0` and `kExtraTol = 5e-4` (from 1e-4) in
`bundle_adjustment_caspar_test.cc` to accommodate the 1-DOF scale gauge
difference between Ceres (7 DOFs pinned) and Caspar (6 DOFs pinned) in
double precision.

All generated CUDA files are 100% clean and reproducible from
`caspar_generate.py` without manual edits.

## Test Plan

Tested both `f32` (`CASPAR_USE_DOUBLE=OFF`) and `f64`
(`CASPAR_USE_DOUBLE=ON`) builds on an NVIDIA A100-SXM4-40GB GPU (CUDA
12.4).

### 1. f32 Build (`CASPAR_USE_DOUBLE=OFF`)
- `ctest --test-dir build-f32 -R "bundle_adjustment"
--output-on-failure`: **4/4 passed (100%)**
  - `controllers/bundle_adjustment_test`: Passed
  - `estimators/bundle_adjustment_test`: Passed (40/40 tests)
  - `estimators/bundle_adjustment_ceres_test`: Passed
  - `estimators/bundle_adjustment_caspar_test`: Passed (17/17 tests)
- CUDA Memory / Race Checker:
`compute-sanitizer --tool memcheck
./build-f32/src/colmap/estimators/bundle_adjustment_caspar_test`
  - **`0 errors`** across all 17 tests (`[ PASSED ] 17 tests`).

### 2. f64 Build (`CASPAR_USE_DOUBLE=ON`)
- `ctest --test-dir build-f64 -R "bundle_adjustment"
--output-on-failure`: **4/4 passed (100%)**
  - `controllers/bundle_adjustment_test`: Passed
- `estimators/bundle_adjustment_test`: Passed (40/40 tests, including
`VariablePrincipalPoint`)
  - `estimators/bundle_adjustment_ceres_test`: Passed
  - `estimators/bundle_adjustment_caspar_test`: Passed (17/17 tests)
- CUDA Memory / Race Checker:
`compute-sanitizer --tool memcheck
./build-f64/src/colmap/estimators/bundle_adjustment_caspar_test`
  - **`0 errors`** across all 17 tests (`[ PASSED ] 17 tests`).
@sarlinpe

sarlinpe commented Oct 22, 2026

Copy link
Copy Markdown
Author

@aaron-skydio Request for review, PTAL, thank you!

@aaron-skydio

aaron-skydio commented Oct 23, 2026

Copy link
Copy Markdown
Member

Thanks for the ping, sorry I'm taking some time to get around to these but I will get them merged. Thanks for making these PRs. One of the things here is that some of these are only needed on clang and seem to regress performance a bit on nvcc, so we'll probably want to guard them to only apply on clang. I can measure more specifically and do that on top of your changes here when I have time, or if you're able to go ahead and do that it'd be appreciated

@sarlinpe
[Caspar] Only use explicit warp-shuffle FlushSumShared reduction on C…
0c32747 
…lang

The explicit warp-shuffle reduction in FlushSumShared works around incorrect
cg::reduce results on labeled_partition groups under Clang CUDA, but regresses
performance under nvcc. Guard it with __clang__ && !__NVCC__ and restore the
original cg::labeled_partition + cg::reduce path for nvcc.

Zero-initialization of the reduction accumulators is kept unconditionally.

Signed-off-by: Paul-Edouard Sarlin <15985472+sarlinpe@users.noreply.github.com>
@sarlinpe
[Caspar] Factor Clang FlushSumShared reduction into LabeledWarpSum he…
25454b3 
…lper

Move the Clang-only warp-shuffle reduction into a LabeledWarpSum helper so
that FlushSumShared keeps a single loop and only the reduction call differs
between Clang and nvcc, reducing duplication.

Signed-off-by: Paul-Edouard Sarlin <15985472+sarlinpe@users.noreply.github.com>
@sarlinpe

sarlinpe commented Oct 23, 2026

Copy link
Copy Markdown
Author

@aaron-skydio Added a branch for Clang vs nvcc, PTAL.

fringehush-gif added a commit to fringehush-gif/symforce that referenced this pull request Nov 9, 2026
@fringehush-gif
Caspar: Preserve an opt-in legacy double reduction [R17-F5] (fix)
cfff699 
Runtime regeneration can overwrite a downstream project's retained double reduction. Make that compatibility choice explicit while keeping ordinary generated projects on cooperative groups.

## Changes

Add a default-off double specialization, retain the default float/double path and initialize three reduction accumulators. Translation units in a program must use one consistent profile. New maintained generation tests cover runtime copy/symlink replacement and ensure ordinary buildfiles leave compatibility disabled.

COLMAP commit 68b722be23fdba964a24583720bde07675171d47, "Regenerate Caspar kernels with upstream SymForce fixes (#4799)," retained the double lane loop. Related SymForce PR symforce-org#472, "[Caspar] Fix FlushSumShared warp reduction and zero-initialize reduction accumulators" (symforce-org#472), was inspected at 25454b3; this does not claim reproduction of its Clang-specific fault.

## Validation

- Twelve maintained generation/template methods pass, exercising copy and symlink replacement, nine-file runtime transport and the f32/f64 buildfile policy.
- Four genuine CMake graphs have the intended definitions on all 242 core compile rows and dependent probes; eight selected real translation units compiled.
- Twelve NVCC/Clang reduction runs pass 21,260 scalar-oracle assertions each across default float/default double/compatibility double with RTX 3060 and 3080 Ti selected by UUID environment. Matching retained baseline proofs were reused.
- Matched helper costs:24 cells, N=5 per arm,CUDA-event intervals include reset, reduction and possible host enqueue idle. Medians varied -6.50% to +5.75%; 23 ranges overlap. Largest increase:23.491 -> 24.842 microseconds (default-double/unique, RTX 3060/NVCC); sole disjoint cell:48.068 -> 46.643 microseconds (compatibility-double/shared, RTX 3060/NVCC). No general speedup or neutrality follows.

These are generation and focused reduction checks; the other two accumulator initializations were source reviewed. A full solver build/run was not performed.

AI disclosure: Totally vibe coded
fringehush-gif added a commit to fringehush-gif/colmap that referenced this pull request Nov 9, 2026
@fringehush-gif
Caspar: Preserve the generated double reduction profile [R17-F5] (fix)
6f5b19b 
Regenerating Caspar must preserve COLMAP's retained f64 reduction rather than silently restoring a different runtime policy.

## Changes

Extend the existing generator loader override to publish the compatibility definition from the generated f64 core target. The choice survives standalone and parent CMake builds; f32 retains its default. Both runtime copies come from the same owner. Maintained template tests cover precision selection and reject incompatible template shapes.

Runtime owner: cfff699b8d6b069a447b40282a805262a2718d1f (fringehush-gif/symforce@cfff699).

COLMAP commit 68b722b, "Regenerate Caspar kernels with upstream SymForce fixes (colmap#4799)," retained the double lane loop. Related SymForce PR colmap#472, "[Caspar] Fix FlushSumShared warp reduction and zero-initialize reduction accumulators" (symforce-org/symforce#472), was inspected at 25454b315848d1c33f7577e48b0598e84efe5278; this does not claim reproduction of its Clang-specific fault.

## Validation

- Twelve maintained generation/template methods pass, exercising copy and symlink replacement, nine-file runtime transport and the f32/f64 buildfile policy.
- Four genuine CMake graphs have the intended definitions on all 242 core compile rows and dependent probes; eight selected real translation units compiled.
- Twelve NVCC/Clang reduction runs pass 21,260 scalar-oracle assertions each across default float/default double/compatibility double with RTX 3060 and 3080 Ti selected by UUID environment. Matching retained baseline proofs were reused.
- Matched helper costs:24 cells, N=5 per arm,CUDA-event intervals include reset, reduction and possible host enqueue idle. Medians varied -6.50% to +5.75%; 23 ranges overlap. Largest increase:23.491 -> 24.842 microseconds (default-double/unique, RTX 3060/NVCC); sole disjoint cell:48.068 -> 46.643 microseconds (compatibility-double/shared, RTX 3060/NVCC). No general speedup or neutrality follows.

These are generation and focused reduction checks; the other two accumulator initializations were source reviewed. A full solver build/run was not performed.

AI disclosure: Totally vibe coded

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@sarlinpe @aaron-skydio