fMCS seed grow utilities

[scal444]

Utils for getting new candidate bonds to add and pruning. fMCS has as funky extension mechanism, where it first tries to add all new bonds. Then tries each bond singly. Then, other combinations.

PR for #221

[greptile-apps]

Greptile Summary

This PR adds fmcs_grow.cuh, which provides the two cooperative CUDA helpers needed by the fMCS seed-grow kernel: fillNewBondsCooperative (warp-parallel boundary-bond enumeration with ring-closing vs. atom-adding classification) and pruneIndividualBondsCooperative (the serial-outer Stage 1 per-bond match loop). A matching unit-test file covers five fillNewBonds scenarios and one Stage 1 prune scenario, with a correct CMake target including --extended-lambda.

• fillNewBondsCooperative correctly uses atomicAdd for slot allocation, clamps outCount on overflow, and returns a uniform success/overflow signal to all lanes.
• pruneIndividualBondsCooperative serializes the outer bond loop (intentionally) but lacks documentation that matchFn must return a uniform boolean across all warp threads.
• The test's uint64_t backing storage for QueuedT shared-memory slots has no explicit alignas(16), so alignment correctness for warpCopy's int4 path depends on declaration order.

Confidence Score: 4/5

Safe to merge; all findings are robustness and documentation gaps rather than current failures, and the core grow/prune logic is sound.

The cooperative logic in fmcs_grow.cuh is well-structured. The test's uint64_t backing storage currently lands at a 16-byte boundary by declaration order so no misaligned copy occurs today, but the arrangement is fragile. The undocumented uniform-return requirement on matchFn is similarly not broken today but is a latent trap for future callers.

tests/test_fmcs_grow.cu (pruneStage1Driver shared-memory layout and missing alive=false test case) and src/mcs/fmcs_cuda/fmcs_grow.cuh (matchFn uniform-return contract in pruneIndividualBondsCooperative).

Important Files Changed

Filename	Overview
src/mcs/fmcs_cuda/fmcs_grow.cuh	New header implementing two cooperative CUDA helpers. Logic and synchronization patterns look correct; the undocumented contract that matchFn must return a uniform value across all warp threads is a latent hazard if the interface is reused with a divergent functor.
tests/test_fmcs_grow.cu	Comprehensive unit tests for both new helpers. Two issues: the uint64_t backing arrays for QueuedT shared memory in pruneStage1Driver lack alignas(16), silently relying on declaration order for correct warpCopy alignment; the prune test has no coverage for the alive = false skip path.
tests/CMakeLists.txt	Adds test_fmcs_grow build target with correct includes, linkage, and the required --extended-lambda compile option; also registers it in TEST_LIST.

_{Reviews (1): Last reviewed commit: "Use project-rooted includes in fMCS grow..." | Re-trigger Greptile}

[greptile-apps]

Shared-memory alignment of QueuedT backing arrays is not explicit

warpCopy uses int4 (16-byte) loads/stores and its own comment requires 16-byte alignment. uint64_t arrays in shared memory are 8-byte aligned. The layout currently works because parentStorage is the first __shared__ variable (offset 0, always 16-byte aligned) and sizeof(QueuedT) is a multiple of 16 (due to alignas(16) on QueuedSeed). However, inserting any new __shared__ declaration before these arrays, or reordering them, would silently break the alignment and produce incorrect copies in shared memory. Consider adding alignas(16) (or alignas(alignof(QueuedT))) to both storage arrays to make the requirement self-documenting and robust to future edits.

Note: If this suggestion doesn't match your team's coding style, reply to this and let me know. I'll remember it for next time!

[greptile-apps]

matchFn must return a uniform value across all warp threads — not documented or enforced

The function uses if (matchFn(childWorkspace)) { childSink(childWorkspace); } inside a cooperative context. If matchFn returns different values to different lanes (warp divergence), some threads enter childSink while others do not. Any group.sync() inside childSink (which the test's own childSink has, and which a realistic childSink doing queue.push would likely have) would deadlock in that scenario. The outer group.sync() after the if would also deadlock. The existing test only exercises a uniform-returning stub, so this hazard is not caught today. Adding a note to the doc-comment that matchFn must return the same boolean on every thread in group would make the contract clear for future callers.

[greptile-apps]

No test coverage for the alive = false skip path in pruneIndividualBondsCooperative

pruneIndividualBondsCooperative has an early-continue guard if (!bonds[i].alive) continue; that Stage 1 uses to exclude bonds that failed singleton matching. The single test case for this function uses four bonds, all with alive = true, so the continue path is never exercised. A test that mixes alive = false bonds with surviving ones would verify the skip logic and also confirm that numMatchAttempts is not incremented for dead bonds.