Repository navigation
Merge upstream main, build the Caspar tests with HIP, fix the first PCG residual update - #2
Merged
Merged
Conversation
GitOrigin-RevId: 30ba631fedfccf6e2efd588ed4b69f369b7773f7
- Drop Python 3.9 support from symforce; minimum is now Python 3.10 - Python 3.9 reached EOL and security-patched packages (pip>=26.1, urllib3>=2.7.0, requests>=2.33.0) require Python >=3.10, making it impossible to resolve dependabot alerts while supporting py39 - No code changes to use new features, so you could still build SymForce from source on py3.9. This is unsupported though and will not continue to be checked. GitOrigin-RevId: 5886f72eadd864069fd6d87e160c5b57aab6dc53
…x CUDA architecture list in Caspar template [Caspar] Minor fix: Export CASPAR_MIN_ARCH and fix CUDA architecture list in Caspar template CASPAR_MIN_ARCH was previously a local CMake variable, invisible to parent projects consuming Caspar via FetchContent. This change makes it a CACHE INTERNAL variable so downstream consumers (e.g. COLMAP) can read it after FetchContent_MakeAvailable(caspar) without hardcoding the value. Also fixes the CMAKE_CUDA_ARCHITECTURES list: the previous 75 80 86 89 75-virtual generated PTX for all four architectures and emitted compute_75 twice. The corrected list 75 80-real 86-real 89-real generates SASS for all and PTX only from sm_75, matching the stated intent. Closes symforce-org#463 GitOrigin-RevId: bc4a9a1f02b41ec4b804bfdd5021e88d2bad1ee9
Move the recursive _get_hash_recursive / _computeHash algorithm out of
each language template and onto Struct.compute_recursive_hash in
syntax_tree.py. Templates now render the precomputed integer literal
returned by the AST. Cross-package types are resolved via JSON files
that thread {full_name -> hash_hex} through skymarshal via
--hash-info-out / --hash-deps-info.
GitOrigin-RevId: 4bff1b98a159fd7f5e9c28a36159690aba8ef244
GitOrigin-RevId: cf7b840faa103ec2f49dd47884a9329456cd5d0e
GitOrigin-RevId: d878d26c0406581d45dab59260906d0be49ec319
GitOrigin-RevId: 2c848376c2a1f50b4976214ff9085c0d81c9c5fb
The recursive-hash precompute added in the skymarshal codegen change
resolves cross-package types (e.g. eigen_lcm.Vector2d) via the eigen_lcm
.lcm sources, located through path_util.eigen_lcm_lcmtypes_dir(). That
lookup checks third_party/eigen_lcm/lcmtypes, which exists in source
checkouts but is not shipped in wheels, so wheel codegen failed with:
KeyError: compute_recursive_hash: cannot resolve eigen_lcm.Vector2d
Copy the .lcm sources into the same third_party/eigen_lcm/lcmtypes path
(relative to the install root) during the wheel build so the existing
lookup finds them. Editable installs already have them in the source
tree.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Topic: sf-fix-lcm-hash-wheels
Reviewers: jerry,nathan,ryan-br
Assignees: jerry
GitOrigin-RevId: ba89098baa1ae1013af6d0002e979fec4661419f
GitOrigin-RevId: a2b0fe32d8ed52611e69bf9ab153e225dce0c087
GitOrigin-RevId: 1117a1a1729d393c11ac74870522900252cce1f8
On gcc 13, building some lcm types is failing due to uninitialized hash. It should be pretty cheap to resolve the warning even if it is a false positive. Signed-off-by: Brian Kubisiak <brian.kubisiak@skydio.com> GitOrigin-RevId: 47a7f3ce300aa151d8833d918f647fd08747c286
GitOrigin-RevId: 12130b4e2f6130e794de7713ec0afefa5e719088
Topic: red-cast Reviewers: william-a GitOrigin-RevId: aa123dbcf035d1df8af9b381a3a70bc4a73ccd16
Build manylinux aarch64 wheels This adds Linux aarch64 wheels to the wheel build, next to the existing x86_64 Linux and arm64 macOS ones. The `build-linux` job now runs the same cibuildwheel configuration natively on GitHub's free `ubuntu-24.04-arm` hosted runner (no QEMU), so build times should be in line with the x86_64 jobs. The `os` matrix entry becomes a runner/arch object like the macOS job already uses, `CIBW_BUILD` selects `manylinux_<arch>`, and the per-job artifact names include the arch so the two Linux builds don't collide before `merge-wheel-artifacts` picks them up. Nothing changes for the existing x86_64 build beyond the artifact name. I built and tested `cp312-manylinux_aarch64` locally with cibuildwheel 3.4.1 against the same `manylinux_2_28_aarch64` image (arm64 Docker on Apple Silicon), including `run_wheel_tests.py` under both `symengine` and `sympy`. auditwheel tags the result `manylinux_2_27_aarch64.manylinux_2_28_aarch64`, matching the x86_64 wheels. Partially addresses symforce-org#146. Assisted-by: Claude:claude-fable-5 Signed-off-by: Ramon Roche <mrpollo@gmail.com> Closes symforce-org#466 GitOrigin-RevId: ea69df6cc28834a06ed916c7b3919f576256f019
Upstream moved on since this branch was last synchronized, leaving the pull request conflicting. Merge current main and resolve the one conflict so it applies cleanly again. Nothing is rebased: the previously reviewed commits stay exactly where they were. symforce/caspar/source/templates/buildfiles/CMakeLists.txt.jinja: upstream introduced the CASPAR_MIN_ARCH cache variable and revised the CUDA architecture list in the same block this branch had moved into the else() of the USE_HIP switch. The resolution keeps the USE_HIP/CUDA structure and adopts upstream's new CUDA block verbatim inside the else() branch, CASPAR_MIN_ARCH included. The HIP branch of the template is untouched, and CASPAR_MIN_ARCH has no other consumer, so scoping it to the CUDA branch changes nothing for either vendor. Authored with assistance from Claude (Anthropic). Test Plan: Linux gfx1100 (Radeon Pro W7800, ROCm 7.2.1) for the HIP path, and a CUDA 12.8 compile+link check with nvcc on the same host (no NVIDIA GPU present). The runtime CMakeLists.txt is rendered from templates/buildfiles/CMakeLists.txt.jinja with python_bindings=False. ``` cmake symforce/caspar/source/runtime -B build_hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 cmake --build build_hip -j16 # libcaspar_runtime_core.a links CUDACXX=nvcc cmake symforce/caspar/source/runtime -B build_cuda -DUSE_HIP=OFF cmake --build build_cuda -j16 # CUDA path with the new arch list links ``` The gfx1100 device ISA of all three kernel objects (shared_indices, solver_tools, sort_indices) is identical before and after the merge, so GPU behavior is unchanged from the reviewed branch.
SumStore's second stage has the first warp tile read inout_shared and reduce it into shared_tmp, and the function returns without a trailing __syncthreads(). The buffer is therefore not safe to reuse until the next block-wide barrier, and today that safety rests on non-local invariants of the call sites: SumStore and SumFlushFinal each open with a block-wide barrier, and every other store into the buffer is indexed by threadIdx.x scaled by the storing accessor's element stride, so nothing below element 32 is written by a thread outside the reading tile, whose own reads the shuffle collectives order ahead of any lane leaving the function. The trailing barrier replaces those implicit guarantees with a local one, at the cost of one barrier per call. Suggested by a user running this branch on gfx1151 / ROCm 7.2, where solver runs showed intermittent NaNs and run-to-run nondeterminism on the same binary, and adding this barrier made their determinism protocol pass 22/22 across fresh seeds and a larger synthetic problem. The precise interleaving behind their failure was not pinned down (their own revert-and-retest was inconclusive), so this change is offered on the strength of that empirical evidence and of the contract argument above, not a root-caused race. The barrier is not gated to HIP: the reuse contract is not platform-specific. Authored with assistance from Claude (Anthropic). Test Plan: Linux gfx1100 (Radeon Pro W7800, ROCm 7.2.1). A standalone harness compiles memops.cuh with hipcc and launches a 1024-thread kernel issuing four back-to-back SumStore calls against the same inout_shared buffer followed by SumFlushFinal, over a problem size that exercises the valid mask: ``` hipcc -x hip sumstore_test.cpp -I symforce/caspar/source/runtime --offload-arch=gfx1100 -o sumstore_test ./sumstore_test # all 4 sums exact, 20 reruns bit-identical ``` The pre-change code passes the same harness on this card (that call pattern is protected by SumStore's own leading barrier, and the gfx1151 failure does not reproduce here), so this is verified as a no-regression hardening. The changed header is not part of the prebuilt runtime library; it is compiled by the generated kernels, so the CUDA side is covered by compiling a representative translation unit with nvcc 12.8: ``` nvcc -arch=sm_75 -c nvcc_memops.cu -I symforce/caspar/source/runtime -o nvcc_memops.o # nvcc_memops.cu: includes memops.cuh and instantiates SumStore, SumFlushFinal, FlushSumBlock ``` Co-authored-by: bjoernellens1 <64093272+bjoernellens1@users.noreply.github.com>
Collaborator
Author
|
To approve this fix round, leave a comment containing this line by itself: To send it back to the porter instead: Approving covers the commits on this branch and the section under '## Upstream reply' in the body, which is posted verbatim as a comment on the upstream pull request after the merge, behind a standing line disclosing that it was drafted by an AI assistant. |
GitOrigin-RevId: 1db1ee491af31c1251ee3393f7693cfe6e651cc5
Fix Rust codegen matrix return types ## Problem Rust codegen formats every matrix returned directly from a generated function as `nalgebra::SVector<scalar, rows * cols>`. This loses the original two-dimensional shape. For example, a function returning `Matrix23` is currently generated with an `SVector<f64, 6>` signature and constructor. The multi-output path already handles this correctly through `format_typename`, which selects `SVector` only when either matrix dimension is one and otherwise selects `SMatrix`. ## Approach Reuse the same shape-aware formatter in both places used by the single-return path: - the generated function return type; - the constructor used for the returned value. This keeps the single-return and multi-output paths consistent without duplicating their vector-versus-matrix condition. Scalar returns and vector-shaped matrices retain their existing behavior. ## Regression coverage The Rust codegen test now includes a non-square `2×3` matrix return. A non-square shape verifies both dimensions explicitly instead of checking only the flattened element count. Before this change, the generated code used: ```rust nalgebra::SVector<f64, 6> ``` It now uses: ```rust nalgebra::SMatrix<f64, 2, 3> ``` The generated crate still includes the existing `3×1` return case, which remains an `SVector<f64, 3>`. ## Validation - Rust codegen test with the SymEngine backend - Rust codegen test with the SymPy backend - Cargo build of both generated test-data crates - Ruff 0.9.10 lint and formatting checks - `git diff --check` Fixes symforce-org#439 Closes symforce-org#467 GitOrigin-RevId: cb6eedbe1d8f7e42c779e33d27056b77338c61bd
Existing import cycle:
```
symforce.cam.__init__
→ symforce.cam.atan_camera_cal
→ symforce.cam.camera_cal
→ symforce.util
→ symforce.codegen
→ symforce.codegen.codegen
→ symforce.symbolic
→ symforce.util.lambdify ← FAILS
```
Topic: sf-cycle
Reviewers: ryan-br,ashley-m
GitOrigin-RevId: fa2ab220246008a3b8157f631127fe3757f25318
Assigning a fixed-size vector into a dynamically-sized one hides the destination size from Eigen, which then selects a linear vectorized assignment loop. When the result is shorter than one packet that loop body can emit a read past the end of the array that ends up discarded. Work around this by copying out through a fixed-size Map instead, following the idiom already used in JacobianFixedMap. This makes the destination size a compile-time constant, so Eigen never emits the vectorized path and the load disappears. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Signed-off-by: Brian Kubisiak <brian.kubisiak@skydio.com> GitOrigin-RevId: afc3b23c93acd3f1d63cb52a204d96e7aa69df51
Notations (the `#name{prop = value}` syntax) could only be attached to a struct
or an enum. Add a `notations` prefix to every member production so they can also
be attached to an individual member, which lets a backend take per-member
direction rather than only per-type direction.
A notation declares which scopes it is valid in via `NotationSpec.allowed`, so
add a "member" scope alongside "struct" and "enum". Notations written ahead of a
multi-member declaration (`int64_t first, second;`) apply to every member the
declaration introduces.
Member notations are included in `__repr__` so that a type round-trips through
the .lcm source that `--lcmout` emits. They are not part of the type hash, the
same as struct and enum notations.
Also initialize `Notation.properties` unconditionally. It was only assigned for
known notations, so reading it on an unknown one (which is reachable via
`--allow-unknown-notations`) raised AttributeError.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
GitOrigin-RevId: 8ac048972991ad2c9ca6dd00125f4e0ae21431c3
Add a member notation that stores a dynamically-sized array inline, with room for a fixed number of elements, instead of in a std::vector. Embedded targets must not allocate, and some go further and overlay an LCM type directly on a hardware payload buffer, which needs the array to be a plain inline one at a fixed offset. The encoded form is a normal dynamically-sized array, and the capacity is not part of the type hash, so it is a bound on local storage only and two builds may disagree about it without disagreeing about the type. We use a std::array of that capacity plus a bounds check on both encode and decode, a message longer than the capacity fails to decode. The length is cast to size_t first, so a negative length in a signed length member is also rejected. Only supports C++ for now. The other backends ignore the notation and keep their existing containers, which is fine because nothing about the wire format changes. Rust is the other one that would benefit from this. It is only valid on a one-dimensional array whose length lives in a member of the same struct (`items[num_items]`), because that member stores the length. An automatic length (`items[uint8_t]`) is derived from the container's own size, which doesn't work for a fixed-size container. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> GitOrigin-RevId: 626a0d9ef8550be6f24c3ef3202c482242f74d55
GitOrigin-RevId: 517ff7b1231a64857a9dd3a734c1076f47014a8c
GitOrigin-RevId: c7f5580bec849f41719b786a0d4ced1bef951d54
Upgrade fmt to 12.2.0 (the latest tagged release) Signed-off-by: Brian Kubisiak <brian.kubisiak@skydio.com> GitOrigin-RevId: dede9296e9d4619352aa54d940d52aa7bb7c58e5
The CMake build still needs the fstring fix because its include directories are not SYSTEM. Take it from upstream 1685e694 and d5af52d9, which landed after the 1.17.0 tag and are in no release. patch_spdlog.sh stamps a hash of the whole series instead of reverse-dry-running each patch, since 0002 rewrites lines 0001 adds and the per-patch check would read an already-patched tree as unpatched. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Signed-off-by: Brian Kubisiak <brian.kubisiak@skydio.com> GitOrigin-RevId: 86edce82319bd11f39f2af26856568862aaf97b8
GitOrigin-RevId: 6884434b3265638d98fadbec3d61eafdec06ea28
Fixes symforce-org#469 Threads past `problem_size` still reach the accessor calls emitted outside the valid-lane guard, and pass them registers nothing assigned in those threads. For `AddSum` that register is the summand `SumStore` reduces. `SumStore` masks invalid lanes with `valid ? data : 0`, but on a large enough kernel LLVM speculates the guarded body and drops the mask, so the padding lanes' garbage lands in the sum. It was reported from a COLMAP OpenCV bundle adjustment on HIP (AMD) scoring NaN; the same kernel built with clang for sm_75 (Nvidia) scores NaN on any problem size that isn't a multiple of the block size. nvcc keeps the mask and emits identical code with or without this change. Zeroing the registers where they are declared makes the read defined for every accessor, including the pair accessors, which hit the same path. `caspar_padding_lane_cuda_test` runs a checked-in generated score kernel over problem sizes that leave padding lanes. It is the first test here that needs a GPU: `*_cuda_test.cc` tests are built in CMake only when a CUDA compiler is available. Its codegen test also imports caspar under the sympy API, where the symengine in SymForce's build directory isn't on `sys.path` until `symforce.find_symengine()` puts it there. caspar now calls it, and it loses its leading underscore since caspar isn't the only caller. Running this test built with clang on a Quadro RTX 4000 fails 4 of 5 assertions without the fix and passes with it. Assisted-by: Claude:claude-opus-5 Signed-off-by: Aaron Miller <aaron@skydio.com> GitOrigin-RevId: 4c2a6d0744b03cd0ac0659b819ef20349bceb45a
Generated solvers and kernels were declaring unused variables, which trips nvcc's "declared but never referenced" and "set but never used" diagnostics under warnings-as-errors. `DoNormalize` and `DoAlphaFirst` declared dead pointer locals, now dropped. A kernel whose body comes out empty still emitted its thread index, shared scratch, and accessor index setup; those are now emitted only when there is a body. No kernel computes anything different. Assisted-by: Claude:claude-opus-5 Signed-off-by: Aaron Miller <aaron@skydio.com> GitOrigin-RevId: 91611ee3c284b3545711afe226fe638ce5b6f72f
This adds a custom block tri-diagonal solver for use with trajectory optimization problems (where the hessian structure is block tri-diagonal). This is often dramatically faster than using the existing sparse cholesky solver `sym::SparseCholeskySolver` when the diagonal and off-diagonal blocks are relatively dense. GitOrigin-RevId: a676e2285da81ee977fa30c567f951476d5f24e3
Two independent, opt-in flags on Skymarshal's proto emitter, mirroring the existing `--python-package-prefix` pattern - both default to today's behavior, so no existing caller is affected: - `--proto-package-prefix`: override the generated proto's package namespace. - `--proto-flat-output`: skip the `<package>/pbtypes/` subdirectory convention and write the generated `.proto` directly under `--proto-path`. ## Test plan - [x] Verified both flags standalone against a real `.lcm` file, with and without each flag, confirming unchanged output when omitted. - [x] Verified the two flags combine correctly (package prefix applied, output flattened). GitOrigin-RevId: ee928a4f1a3c1c08bc55b1a2e88c0fabc48702a9
include-cleaner only understands macro-based header guards if they're at the top of the file, otherwise it treats the file as not self-contained which interferes with include-cleaner correctness for files that include headers like this. Fix this ordering, and temporarily add `IWYU pragma: export` where necessary for compatibility. GitOrigin-RevId: 7b0d3e67575520d52b48149b43c735a6248c67a2
Generated kernels declared their scratch buffer and one SharedIndex array per shared argument as static `__shared__` arrays. Static shared memory is capped at 48 KB per block, and fp64 factors with several shared (Tunable) node arguments can exceed it at the fixed block size of 1024. The scratch buffer and the shared index arrays now live in one dynamic shared memory buffer, at offsets computed at generation time, and each wrapper opts its kernel in with `cudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemorySize, ...)` once before the first launch, throwing if the GPU refuses. Kernels under 48 KB behave as before. Generation also checks the size: it raises, naming the kernel and its shared arguments, above 227 KB (the largest opt-in limit of any GPU), and warns above 48 KB, since such kernels only launch on GPUs with a larger opt-in limit (64 KB on sm_75, 99 KB on sm_86, 163 KB on sm_80). The checked-in outputs of `symforce_caspar_codegen_test` are regenerated with dynamic shared memory, and pass with their CUDA tests. GitOrigin-RevId: 68c16e606cbca80545288d9941c881f40c0c3e19
…nd short chains Caspar gives every node type used by a `Pair` factor a block-tridiagonal preconditioner, solved by cyclic reduction in `DoNormalize()`. The generated host code for it had three bugs: - The reduction loop was emitted inline for each node type, with `i`, `offset` and `remaining` declared at function scope. With two or more such node types (say an IMU factor over consecutive poses and velocities, plus bias random walks) solver.cc didn't compile. - For a chain of one node, or none, the final two-node solve still wrote a second node's solution past the end of the buffer. For a 3D node type it landed on the first node's third component, which then never moved. - For chains of one or two nodes that solve used a stride of `2 << -1`, which is undefined. Chains shorter than two nodes now use the block-diagonal `Normalize` kernel, which is exact when there are no off-diagonal blocks. The stride is `1 << (i + 1)`, and the reduction is scoped per node type. `symforce_caspar_codegen_test` now also generates a library with a 2D and a 3D chain, each linked by its own `Pair` factor, priors on every node, and a factor coupling the two. `caspar_tridiagonal_cuda_test` solves it for chains of 1, 2, 3, 7 and 8 nodes and compares against a dense Eigen solve. Without coupling the preconditioner is exact, so one LM iteration with one CG iteration lands on the dense solution; with coupling the solver converges to it. The single-node 3D cases fail without the fix. Rather than add a third copy of the CMake block each Caspar CUDA test had, the CMake build now defines them in one loop over `*_cuda_test.cc`, as it does the C++ tests. Each test compiles the code generated into `symforce_caspar_<name>_test` and links one library of the Caspar runtime sources, and Eigen. `caspar_padding_lane_cuda_test` is the one exception: its kernel is generated into `symforce_caspar_codegen_test`. Built the Caspar CUDA tests from `test/CMakeLists.txt`, configured alone with CUDA 12.6 and stub targets for the SymForce libraries, and they passed under `ctest -R caspar` on a Quadro RTX 4000, both here and with the next commit's backend coverage test. GitOrigin-RevId: 16c04302de4d5b61e28410a60d39a077121d5723
Caspar only turned `Number` atoms into constants, so an expression with pi or e, which SymEngine represents as number symbols, failed to generate with a `KeyError`. It also had no mapping for floor, so a factor couldn't use `sf.wrap_angle` (for a yaw error, say), and none for ceil, log, the hyperbolic functions and their inverses, erf, erfc, lgamma or tgamma. Any atom that isn't a symbol is now a constant, each of those functions maps to its CUDA math function, and `exp(x)` is generated as `exp` rather than `pow(e, x)`. To cover every expression type at once, `symforce_caspar_codegen_test` now generates a kernel from `backend_test_function`, the expression list the CUDA, Rust and PyTorch backends generate, with the same extra functions as the CUDA backend. Next to it the test checks in the values SymPy gives at five (x, y) points, NaN where the result isn't real. `caspar_backend_coverage_cuda_test` runs the kernel at those points and checks all 60 outputs to 1e-12 relative. Nothing had to be excluded. The SymForce CMake build picks the new test up with the other Caspar CUDA tests. GitOrigin-RevId: 1645d1e8cb66dd8ff819128b01a860cfa1048881
The generated `.pyi` stub declared the `ExitReason` members as `int`, but `SolveResult.exit_reason` is an `ExitReason`, so comparing a result's exit reason to a member needed a cast to pass mypy. The stub now declares the members as class variables of type `ExitReason`, which is what the `py::enum_` binding defines. GitOrigin-RevId: dc5416986a415d82db913c5a10863e253a39d56c
The manifold IMU preintegration (`ImuManifoldPreintegrationUpdate`, used by `sym::ImuPreintegrator`) propagated the covariance of (DR, Dv, Dp) with a discrete noise model: each step's accelerometer noise entered as one vector n with Dv += R n dt and Dp += R n dt^2 / 2, so one step's (Dv, Dp) covariance per axis was sigma^2 dt [1, dt/2; dt/2, dt^2/4], which is rank 1. A window consisting of a single sample therefore had a rank 6 covariance, so the Cholesky factorization to convert to information form fails. Also a window of n equal steps had a per-axis (Dv, Dp) determinant of sigma^4 T^4 (1 - 1/n^2) / 12 instead of the continuous-time sigma^4 T^4 / 12. This change models the measurement noise as continuous-time white noise with the given spectral densities acting over the whole step (the meaning and units of `accel_cov` and `gyro_cov` are unchanged), and adds its covariance integrated exactly over the step. With the measurement and the rotation held over the step, the linearized error dynamics are d/dt (dR, dv, dp) = F (dR, dv, dp) + G (n_gyro, n_accel) with a nilpotent F (F^3 = 0), so the transition matrix is the exact polynomial I + F s + F^2 s^2 / 2 and the step noise covariance is a closed-form sum of nine terms. Per axis the accelerometer part of the (Dv, Dp) block becomes sigma^2 [dt, dt^2/2; dt^2/2, dt^3/3], which is positive definite for any dt > 0. The same integral also yields the cross terms by which the rotation noise accumulated during the step leaks into Dv and Dp within that step (A Sg dt^2 / 2, A Sg A^T dt^3 / 3, and so on, with A = -R [accel]x); without them a window of n steps carried only a fraction (1 - 3/(2n) + ...) of the continuous-time rotation-to-velocity covariance. With the exact step covariance the covariance of a window no longer depends on how the window is divided into steps, and it is full rank even for a single step. The derivation is documented in `step_noise_covariance` in `manifold_symbolic.py`. The generated update function grows from 1253 to 1301 ops. On an i7-14700K, with -O2 -march=haswell, one `IntegrateMeasurement` call goes from about 123 ns to about 133 ns. This is for the preintegration step, not inside the optimization, so this seems fine. This adds tests that one step's covariance equals the analytic closed form and is positive definite, and that a window's covariance does not depend on its subdivision into steps (to machine precision without rotation, and to 1e-6 with rotation against a 100x finer subdivision). The existing Monte Carlo covariance test still passes; it simulates the noise as a perturbation held over each step, which agrees with the white-noise model to the order of 1/n in the sub-dominant terms and is well within its 5% tolerance. GitOrigin-RevId: fd1b2967f93460587ebf3a6ec22470783099c4b8
GitOrigin-RevId: baaeb152430a873c3ac1f70c94315ed4179eab7b
GitOrigin-RevId: c2e7a8e5054391e66caca727bb173a1d404d1808
Generated LCM enums compared through plain inline operators and generated structs had a plain inline member constructor, so neither could take part in a constant expression. The enum comparison operators are now constexpr. The member constructor of a struct is constexpr when every member is a number, boolean, byte, enum or fixed-size array of those, including arrays stored inline with a capacity. Structs holding strings, dynamic arrays or other structs keep the inline constructor, because clang rejects a constexpr constructor that can never produce a constant expression. GitOrigin-RevId: 467f5c6531e64150914b3871eaef93beaa341390
GitOrigin-RevId: cbbac47dbd097780d8660038c9690fb5ef504230
GitOrigin-RevId: f486c59314156fc43e8886a74f6d3e5a80182b5c
The expected values for the caspar backend coverage test were computed by substituting Python floats into SymPy expressions and calling `evalf(30)`. SymPy 1.14 rounds float arguments differently from 1.13, so the generated header differed in the last digit for several outputs, e.g. `(x + y)**2` at (1.6, 2.5) came out as 16.81 instead of 16.810000000000002. The pip wheel CI installs an unpinned SymPy and picked up 1.14, failing `symforce_caspar_codegen_test`. Substitute the exact values of the input doubles as rationals instead, so the result is the correctly rounded value at the double inputs regardless of SymPy version. This shifts eight expected values by about one ulp, well inside the 1e-12 relative tolerance of `caspar_backend_coverage_cuda_test`. This was tested by evaluating the new expressions under SymPy 1.13.3 and 1.14.0 and confirming identical output, and running `caspar_backend_coverage_cuda_test` (on a local GPU) and `symforce_caspar_codegen_test`, which passed. GitOrigin-RevId: e3e5ed95c5b84d5fc21dd53a640c4b1fca3ac34c
Brings the branch up to date with upstream main (acdedd0). Conflicts: symforce/caspar/source/runtime/sort_indices.h Upstream added `#include <sys/types.h>` where this branch selects `hip/hip_runtime.h` or `cuda_runtime.h`; the resolution keeps both, with the new include first. Authored with Claude.
Upstream now launches generated kernels with dynamic shared memory, opting in with cudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemorySize, ...) when a kernel needs more than 48 KB, and the generated solver checks cudaGetDeviceCount. The HIP compatibility header did not map these, so generated code no longer compiled for ROCm. hipFuncSetAttribute only takes the kernel as const void*, while the CUDA runtime also accepts the kernel function itself, so the header adds a small template overload that casts it. Authored with Claude. Test Plan: Build the Caspar GPU tests (generated kernels and solvers checked into test/) for gfx90a and run them on an MI250X: ``` cmake -B build -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_PREFIX_PATH=<ROCm prefix> cmake --build build --target bin.tests.caspar_padding_lane_cuda_test bin.tests.caspar_solver_cuda_test bin.tests.caspar_backend_coverage_cuda_test ./build/bin/tests/caspar_padding_lane_cuda_test ./build/bin/tests/caspar_solver_cuda_test ./build/bin/tests/caspar_backend_coverage_cuda_test ```
This reverts the trailing __syncthreads() added to SumStore in "[ROCm] Harden SumStore's shared scratch buffer reuse contract". That barrier was added for a gfx1151 report of NaN scores and run-to-run nondeterminism in solves, without a root-caused race. Upstream has since fixed the mechanism behind that symptom family in "[Caspar] Zero initialize caspar kernel registers" (symforce-org#469): lanes past problem_size passed uninitialized registers to SumStore, and LLVM can drop the valid-lane mask. Uninitialized registers explain both the NaN and the run-to-run variation. With register zeroing in place, the barrier is not needed: - SumStore and SumFlushFinal open with a block-wide barrier, and every other store into the scratch buffer below element 32 comes from a lane of the tile that reads it. - On an MI250X (gfx90a), the padding-lane score kernel and the solver's alpha numerator/denominator kernel (two SumStore calls back to back into one buffer) give bit-identical per-block partial sums over 200 runs at sizes from 1000 to 1000003, with and without the barrier. With integer inputs, the alpha kernel's per-block sums equal an exact host reference in every run. - Without register zeroing, the padding-lane score is NaN at every size that leaves padding lanes, whether or not the barrier is present. So the barrier does not fix that defect either. Authored with Claude. Test Plan: Build upstream's checked-in kernels against memops.cuh with and without the barrier, with SumFlushFinal's atomicAdd replaced by a per-block store so cross-block float atomic ordering does not mask anything, and run each 200 times: ``` hipcc -x hip -std=c++17 -O3 -DUSE_HIP --offload-arch=gfx90a -I. -Iruntime main.cc kernel_caspar_padding_lane_score.cu runtime/shared_indices.cu -o padding_lane hipcc -x hip -std=c++17 -O3 -DUSE_HIP --offload-arch=gfx90a -I. -Iruntime main_alpha.cc kernel_Matrix21_alpha_numerator_denominator.cu -o alpha ./padding_lane 200 ./alpha 200 ```
The generated kernels and solvers include cuda_to_hip.h instead of the CUDA runtime and cooperative groups headers, so the code symforce_caspar_codegen_test checks in no longer matched what the templates generate. This regenerates it with the test's own --update; nothing was edited by hand. With USE_HIP unset, cuda_to_hip.h includes the same CUDA headers as before, so the CUDA build of this code is unchanged. Authored with Claude. Test Plan: ``` python test/symforce_caspar_codegen_test.py --update python test/symforce_caspar_codegen_test.py ```
DoUpdateRFirst() zeroes r_0_norm2_tot and then accumulates into both r_0_norm2_tot and r_kp1_norm2_tot, but only DoUpdateR() zeroes r_kp1_norm2_tot. The first PCG iteration therefore adds the new residual norm onto whatever the solver's allocation held. The bad value feeds the PCG exit test and beta, so the result depends on uninitialized device memory. Freshly allocated device memory is usually zero, which hides this. On ROCm, a second solver in the same process gets memory freed by the first, and caspar_tridiagonal_cuda_test failed a varying 5 to 7 of its 20 uncoupled cases on an MI250X: step rejected, 199 PCG iterations. Filling the allocation with 0xFF makes the failures deterministic, and zero-filling it makes them disappear. That narrowed the read to the solver's scalar accumulators. The same defect is reachable on CUDA, so the fix is not gated on HIP. Authored with Claude. Test Plan: ``` python test/symforce_caspar_codegen_test.py --update python test/symforce_caspar_codegen_test.py cmake --build build --target bin.tests.caspar_tridiagonal_cuda_test bin.tests.caspar_solver_cuda_test ./build/bin/tests/caspar_tridiagonal_cuda_test # 5/5 runs pass, also with the allocation filled with 0xFF ./build/bin/tests/caspar_solver_cuda_test ```
The *_cuda_test.cc tests run code checked in by symforce_caspar_codegen_test and are built only when CMake finds a CUDA compiler. When it finds no CUDA compiler but finds a HIP compiler, this builds the same tests for AMD GPUs instead. The runtime, the generated .cu/.cc files and the test sources are compiled as HIP, with USE_HIP defined and hip::host and hip::hipcub linked. The tests include cuda_to_hip.h, which includes the CUDA runtime when USE_HIP is unset, so the CUDA build is unchanged. The build itself needs no AMD GPU: with -DCMAKE_HIP_ARCHITECTURES set, CMake does not query the device, so a ROCm-only machine or container compiles every checked-in kernel and solver with the HIP toolchain. That catches HIP compile breaks from template or runtime changes. Running the tests needs an AMD GPU, as the CUDA tests need an NVIDIA one. HIP language support needs CMake 3.21, so the HIP check is skipped on older CMake. Authored with Claude. Test Plan: On an MI250X (gfx90a), ROCm 7.x: ``` cmake -S . -B build -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_PREFIX_PATH=<ROCm prefix> cmake --build build --target bin.tests.caspar_padding_lane_cuda_test bin.tests.caspar_solver_cuda_test bin.tests.caspar_tridiagonal_cuda_test bin.tests.caspar_backend_coverage_cuda_test cd build && ctest -R caspar_ --output-on-failure ```
symforce_caspar_codegen_test checks in the kernels and solvers caspar generates, but not the CMakeLists.txt it generates next to them. That file holds the only generated code that differs between the CUDA and HIP builds: the USE_HIP branch, which selects the HIP language, finds hip and hipcub, compiles the .cu sources (and, with Python bindings, the pybind sources) as HIP, and sets the HIP math flags. A change to that branch, or to the CUDA one, went unnoticed until someone built a generated library on that backend. test_generate_build_files generates the build files for the padding-lane library with and without Python bindings and compares them with checked-in copies, like the other codegen tests. The copies were written with the test's own --update. Changing the HIP compile flags in CMakeLists.txt.jinja makes the test fail. Authored with Claude. Test Plan: ``` python test/symforce_caspar_codegen_test.py --update python test/symforce_caspar_codegen_test.py SYMFORCE_SYMBOLIC_API=sympy python test/symforce_caspar_codegen_test.py ```
The generated kernel and solver headers include cuda_to_hip.h, and so does host C++ that uses them: the generated solver.cc, the Python bindings, and the Caspar GPU tests, all compiled by the host compiler in a CUDA build. On the CUDA path the header also included cooperative_groups and CUB, which only compile under nvcc. A g++ build of caspar_padding_lane_cuda_test failed in cooperative_groups/details/sync.h ('threadIdx' was not declared).
Those includes, the cg alias and CG_REDUCE_SUM are only used by device code, so they are now included only when __CUDACC__ is defined. Host-only translation units get just the CUDA runtime API, which is what they included before the HIP port. The HIP path is unchanged: every file that includes the header is compiled as HIP there.
Authored with Claude.
Test Plan:
CUDA 12.8 (no NVIDIA GPU on this machine, so compile and link only):
```
cmake -S . -B build_cuda -DCMAKE_CUDA_COMPILER=/path/to/cuda-12.8/bin/nvcc
cmake --build build_cuda --target bin.tests.caspar_padding_lane_cuda_test bin.tests.caspar_solver_cuda_test bin.tests.caspar_tridiagonal_cuda_test bin.tests.caspar_backend_coverage_cuda_test
```
HIP, on an MI250X (gfx90a):
```
cmake --build build --target bin.tests.caspar_padding_lane_cuda_test bin.tests.caspar_solver_cuda_test bin.tests.caspar_tridiagonal_cuda_test bin.tests.caspar_backend_coverage_cuda_test
ctest --test-dir build -R caspar_
```
make lint runs clang-format --dry-run -Werror over the tree, and the HIP compatibility header was not formatted with the repository's .clang-format. Formatting only; no code change. Authored with Claude. Test Plan: ``` make lint ```
The ROCm branch of cuda_to_hip.h defined caspar_hip::reduce_sum, reduce_max, match_any_mask and labeled_reduce_sum, plus the CG_REDUCE_SUM and CG_LABELED_REDUCE_SUM macros, and the CUDA branch defined its own CG_REDUCE_SUM. Nothing in the tree calls any of them: memops.cuh and shared_indices.cu do their HIP reductions and copies inline. labeled_reduce_sum was also wrong when lanes with the same label are not contiguous, because its xor butterfly only adds a partner with the same label, so lanes 0 and 3 sharing a label never combine. Remove them so the header carries only the runtime API mapping. The cg namespace alias stays on the CUDA path. Templates and macros that are never instantiated emit no code, so the device code objects of the Caspar test binaries are unchanged. Authored with Claude. Test Plan: ``` grep -rn 'caspar_hip\|CG_REDUCE_SUM\|CG_LABELED_REDUCE_SUM' symforce test cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DCMAKE_HIP_ARCHITECTURES=gfx90a cmake --build build ctest --test-dir build -R caspar cmake -S . -B build-cuda -DCMAKE_CUDA_COMPILER=nvcc cmake --build build-cuda --target bin.tests.caspar_padding_lane_cuda_test bin.tests.caspar_solver_cuda_test bin.tests.caspar_tridiagonal_cuda_test bin.tests.caspar_backend_coverage_cuda_test ```
The README's AMD GPU section called a compile_caspar_library() function that does not exist. The API is the CasparLibrary.compile() static method, so the example now calls caslib.compile(out_dir, use_hip=True, hip_arch=...) after caslib.generate(out_dir), as the examples do. The same section also described the compatibility header wrongly. It does not map __syncthreads or any cooperative groups spelling; it maps the CUDA runtime calls Caspar emits and the cub namespace. It also said hipcc works as the compiler, but the generated build uses CMake's HIP language support, which rejects the hipcc wrapper and needs the ROCm Clang compiler. The text now says so, and notes that without hip_arch CMake builds for the GPU it detects. The comment above the HIP branch of GetDeviceId() in pybind_array_tools.cc still mentioned __device__ helpers in cuda_to_hip.h, which were removed, and local aliases the code never defines. It now says what the code does. No code change outside the comment. Authored with Claude. Test Plan: ``` cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DCMAKE_HIP_ARCHITECTURES=gfx90a cmake --build build ctest --test-dir build -R caspar --output-on-failure ctest --test-dir build python symforce/caspar/examples/kernel_example/gen_and_run.py # with caslib.compile(out_dir, use_hip=True, hip_arch="gfx90a") ```
Collaborator
Author
|
/moat approve |
Collaborator
Author
|
/moat approve |
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.
Fix round for symforce-org#465, refreshed on 2026-11-04. The staging branch is now at e7e0b00 (it was at 1a9e977).
What the maintainer asked
aaron-skydio (2026-09-14) is willing to merge and wants tests that cover at least the codegen part and preferably the build, and AMD GPU testing where we can manage it. The PR also conflicted with main (
sort_indices.h), and the 2026-08-19 reply to bjoernellens1 promised results on their SumStore report.What changed since the published PR head
ce4f2532merges upstream main at acdedd0. The one conflict wasruntime/sort_indices.h: upstream added#include <sys/types.h>where the branch selects the HIP or CUDA runtime header, and the resolution keeps both.38f16ea3mapscudaFuncSetAttributeandcudaFuncAttributeMaxDynamicSharedMemorySizeincuda_to_hip.hfor upstream's new dynamic-shared-memory launch code (78e380a), andcudaGetDeviceCount, which the generated solver has called since 2ee2978 (GPU device selection). That mapping was already missing at the published head e994ef0, so a HIP build of any generated solver failed there. The new HIP build of the tests (9ef7a6a3) is what caught it. The 38f16ea commit body says upstream "now" usescudaGetDeviceCount; that is wrong, and the commit is pushed, so it stays.1a984fc0drops the trailingSumStorebarrier added in 1a9e977. See the reconciliation below.5256d2adregenerates the checked-in codegen test data withsymforce_caspar_codegen_test.py --update. The templates emitcuda_to_hip.h, so the test failed 4/4 right after the merge.daa1a7bfzeroesr_kp1_norm2_totin the solver template'sDoUpdateRFirst(). OnlyDoUpdateR()zeroed it, so the first PCG step of every solver iteration added the new residual norm to whatever the accumulator held: the previous iteration's last value, or uninitialized memory on the first iteration. The value feeds only the PCG exit test (DoBeta()readsbeta_numeratorandalpha_numerator, not this). The new tridiagonal test exposed it on ROCm, and it is not specific to ROCm. The daa1a7b commit body also names beta and "freshly allocated memory"; the commit is pushed, so it stays.9ef7a6a3is the build test.test/CMakeLists.txtbuilds the*_cuda_test.ccCaspar GPU tests with HIP when CMake finds no CUDA compiler but finds a HIP compiler. With-DCMAKE_HIP_ARCHITECTURESset, the tests build with no AMD GPU present.92ec5462is the codegen test.test_generate_build_fileschecks in the generatedCMakeLists.txt(with and without Python bindings), which is where the HIP-specific generated output lives.1ee93b2cincludes cooperative_groups and CUB fromcuda_to_hip.honly when compiling CUDA (__CUDACC__). A CUDA 12.8 build of the test targets failed under g++ without this.d5eb335dapplies clang-format tocuda_to_hip.h, whichmake lintrequires.f82e4abfremoves the unused reduction helpers (caspar_hip::reduce_sum,reduce_max,match_any_mask,labeled_reduce_sum) and theCG_*macros fromcuda_to_hip.h. Nothing called them, and the labeled reduction gave wrong sums for lanes whose labels are not contiguous. The HIP code objects of all 78 compiled.cufiles are byte-identical before and after, for gfx90a and for gfx1100.e7e0b00cfixes the README's ROCm example, which called acompile_caspar_libraryfunction that does not exist (the API iscaslib.compile(out_dir, use_hip=True, hip_arch=...)). It also corrects what the README sayscuda_to_hip.hmaps, replaces its claim thathipccworks as the compiler (CMake's HIP language rejects the hipcc wrapper and needs the ROCm Clang compiler), and updates a comment inpybind_array_tools.ccthat still referred to the removed helpers. No code change outside that comment.SumStore reconciliation
Upstream a768308 (symforce-org#469) and our barrier aimed at the same symptoms (NaN and run-to-run variation in
SumStoresums on AMD), but they are different changes.SumStoreandSumFlushFinalopen with a block-wide barrier, and every other store into the scratch buffer below element 32 comes from a lane of the tile that reads it, so back-to-back calls are ordered without a trailing barrier on wave32 and wave64 alike. Using upstream's checked-in kernels:SumStorecalls back to back into one buffer), whose sums with integer inputs equal an exact host reference in every run.SumStorefails in 159 and 200 of 200 runs at the two largest sizes without the barrier and in none with it, so this check can see a race when there is one.SumFlushFinal. It is the same with and without the barrier.So no race remains for the barrier to fix, and it is removed. bjoernellens1's commit stays in the branch history.
Evidence
Validator records at e7e0b00 come first, then the porter and reviewer evidence.
rustfmtmissing for 2 Rust codegen tests, andnbsphinxmissingjsonschemafor the docs test), Caspar 6/6 with tridiagonal 140/140 in 5 of 5 runs, the 200-run SumStore reproduction exact at every size, the README HIP example runs.rustfmtmissing (2 Rust codegen tests) andnbsphinxmissingjsonschema(docs). At f82e4ab the gfx90a code objects of all 78 HIP object files were byte-identical to d5eb335. The README example (caslib.generate(out_dir), thencaslib.compile(out_dir, use_hip=True, hip_arch="gfx90a"), run with the kernel example's inputs) builds and its checks pass.caspar_tridiagonal_cuda_testpasses 5/5 runs. Without the daa1a7b line, 15 of 140 tridiagonal assertions fail in 5 of 5 runs. At f82e4ab, the gfx1100 code objects and whole object files of the same 78.cufiles are byte-identical to d5eb335.make lintis clean, andjargon.py --port symforceis clean.Reviewing our changes apart from upstream's
This round contains two merges of upstream main, so the PR diff mixes upstream commits with ours. To see our part alone:
git show --cc ce4f2532locally)Upstream reply
Merged current main into the branch. The only conflict was an include in
sort_indices.h.On SumStore: after a768308 the trailing barrier from bjoernellens1's report isn't needed, so I took it back out. The NaN and run-to-run variation they saw fit symforce-org#469. On an MI250X the padding-lane kernel without the register zeroing gives NaN at every padded size, with or without the barrier. With the zeroing, per-block sums are bit-identical over 200 runs whether or not the barrier is there, including the back-to-back
SumStorekernel they pointed at. On a W7800 (wave32) that kernel matches an exact reference in every run, with or without the barrier. Thanks @bjoernellens1 for the report and the reproduction.Tests: the
*_cuda_test.cctests now also build with HIP when CMake finds no CUDA compiler. With-DCMAKE_HIP_ARCHITECTURESset they build without an AMD GPU, so a ROCm container can compile every checked-in Caspar kernel and solver.symforce_caspar_codegen_testnow also checks in the generatedCMakeLists.txt, which is where the HIP build differs.At the head I pushed before, a HIP build of a generated solver failed. The solver calls
cudaGetDeviceCount(since the GPU device selection change, 2ee2978), andcuda_to_hip.hnever mapped it. The new HIP build of the tests is what caught it, and it is mapped now. Running the tests on ROCm turned up one more bug:DoUpdateRFirst()never zeroedr_kp1_norm2_tot, so the first PCG step of each solver iteration added onto the previous iteration's value (or uninitialized memory on the first) before the PCG exit test read it. That's fixed in the template, and it affects CUDA too. A CUDA build of the tests also showed that host.ccfiles including the generated headers failed under g++, socuda_to_hip.hnow includes cooperative_groups and CUB only when compiling CUDA. I also removed some reduction helpers fromcuda_to_hip.hthat nothing used, and fixed the README's ROCm example, which called a function that doesn't exist.Tested on an AMD Instinct MI250X and an AMD Radeon PRO W7800 on Linux (the Caspar GPU tests, the codegen tests and the rest of the test suite, apart from three tests that need tools those machines lack: rustfmt and a docs dependency) and an AMD Radeon 8060S on Windows (the Caspar GPU tests).