Repository navigation
Conversation
…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>
ahojnnes
pushed a commit
to colmap/colmap
that referenced
this pull request
Oct 7, 2026
## 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`).
Author
|
@aaron-skydio Request for review, PTAL, thank you! |
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 |
…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>
…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>
Author
|
@aaron-skydio Added a branch for Clang vs nvcc, PTAL. |
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Follows up on #469 (
a768308e) with two related fixes insymforce/caspar/source/runtime/memops.cuhwhen compiling Caspar kernels with Clang CUDA:StorageT tot = 0;inFlushSumShared,FlushSumBlock, andFlushSumBlockAdd:When a thread is inactive (
idx.argsort == 0xffffor!valid),totwas previously uninitialized across__syncthreads()barriers, triggering-Wsometimes-uninitializedunder Clang and creating anundef/poisonphi node in LLVM IR.cg::labeled_partition+cg::reduceinFlushSumSharedwith an explicit warp-shuffle reduction:cg::labeled_partition(cg::coalesced_threads(), unique)was called outside theif (idx.argsort != 0xffff)guard whilecg::reduce(group, ...)was called inside it. Forfloatanddouble,cg::reduceover acoalesced_groupfalls back tocg::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_syncin lockstep.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 inFlushSumSharedhad size 1). Added a test case with shared nodes (kNumSharedNodes = 64) and a partial block (kNumFactors = 1500).