fMCS seed grow utilities - #231
Conversation
|
| Filename | Overview |
|---|---|
| src/mcs/fmcs_cuda/fmcs_grow.cuh | New header implementing fillNewBondsCooperative and pruneIndividualBondsCooperative; logic for bond scanning, ring-closing classification, overflow clamping, and warp-cooperative parent→child copy is correct |
| tests/test_fmcs_grow.cu | New test file covering 5 scenarios for fillNewBondsCooperative (atom-adding, excluded-bond skip, lastAdded filtering, ring-closing, overflow) and 1 for pruneIndividualBondsCooperative; correctness of shared-memory layout and sync structure verified |
| tests/CMakeLists.txt | Adds test_fmcs_grow build target with --extended-lambda CUDA flag and registers it in the test list; no issues |
| tests/test_fmcs_primitives.cu | Include paths updated from bare relative paths to src/mcs/… prefixes, matching the style used in test_fmcs_grow.cu; purely mechanical cleanup |
Reviews (2): Last reviewed commit: "Formatting" | Re-trigger Greptile
![greptile-apps[bot]](https://avatars.githubusercontent.com/in/867647?s=60&v=4){kind=small}
| __shared__ int survivorIdx; | ||
| QueuedT& parent = *reinterpret_cast<QueuedT*>(parentStorage); | ||
| QueuedT& workspace = *reinterpret_cast<QueuedT*>(workspaceStorage); | ||
|
|
||
| if (threadIdx.x == 0) { | ||
| mcs::fmcs::seedClearWithinThread(parent.seed); | ||
| mcs::fmcs::matchResultClearWithinThread(parent.match); | ||
| // Parent has bond 0 already in its bitset (so pruneIndividualBonds | ||
| // appears as "extending past bond 0" for tracking). |
There was a problem hiding this comment.
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!
| /// per-bond commits depend on shared visited-state from the previous | ||
| /// bond's match call, which is incompatible with running them | ||
| /// concurrently. | ||
| /// | ||
| /// @p childWorkspace must be a per-group shared slot the caller owns | ||
| /// -- the kernel reuses the same buffer it already allocated for the | ||
| /// Stage 0 biggest-child build. | ||
| template <int maxAtoms, int maxBonds, int maxTA, int maxTB, class GroupT, class MatchFn, class ChildSink> | ||
| __device__ __forceinline__ void pruneIndividualBondsCooperative( | ||
| const GroupT& group, | ||
| const QueuedSeed<maxAtoms, maxBonds, maxTA, maxTB>& parent, | ||
| QueuedSeed<maxAtoms, maxBonds, maxTA, maxTB>& childWorkspace, | ||
| const NewBond* bonds, | ||
| int nBonds, |
There was a problem hiding this comment.
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.
| // 4 candidate bonds with bondIdx 1, 2, 3, 4. Stub matchFn passes | ||
| // even bondIdx (-> 2, 4 succeed; 1, 3 fail). All 4 must be tried | ||
| // (matchAttempts == 4); only 2 survivors should reach childSink. | ||
| const std::vector<NewBond> hostBonds = { | ||
| NewBond{1, 5, NewBond::kNotInSeed, true}, | ||
| NewBond{2, 6, NewBond::kNotInSeed, true}, | ||
| NewBond{3, 7, NewBond::kNotInSeed, true}, | ||
| NewBond{4, 8, NewBond::kNotInSeed, true}, | ||
| }; | ||
| AsyncDeviceVector<NewBond> d_bonds(hostBonds.size()); | ||
| d_bonds.copyFromHost(hostBonds); | ||
|
|
||
| AsyncDevicePtr<PruneOut> d_out; | ||
| pruneStage1Driver<<<1, 32>>>(d_bonds.data(), 4, d_out.data()); | ||
| ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess); | ||
| PruneOut out{}; | ||
| ASSERT_EQ(cudaMemcpy(&out, d_out.data(), sizeof(out), cudaMemcpyDeviceToHost), cudaSuccess); | ||
|
|
||
| EXPECT_EQ(out.numMatchAttempts, 4); // all four bonds were tried | ||
| EXPECT_EQ(out.numSurvivors, 2); | ||
| std::set<int> survivors{out.survivorBondIdx[0], out.survivorBondIdx[1]}; | ||
| EXPECT_TRUE(survivors.count(2)); | ||
| EXPECT_TRUE(survivors.count(4)); | ||
| EXPECT_FALSE(survivors.count(1)); | ||
| EXPECT_FALSE(survivors.count(3)); | ||
| } | ||
|
|
||
| // --------------------------------------------------------------------------- |
There was a problem hiding this comment.
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.
cb5c6f7 to
1ce7d61
Compare
| // Try to extend the parent's recorded match by this single bond. | ||
| if (matchFn(childWorkspace)) { | ||
| childSink(childWorkspace); | ||
| } |
There was a problem hiding this comment.
Should we set bonds[i].alive to false if it could not match? This may help if bonds.alive is used again in a later stage. Not necessary if we only reconstruct from childSink from now.
evasnow1992
left a comment
evasnow1992
left a comment
There was a problem hiding this comment.
One minor comment. Changes look good to me. Thanks!
2 participants
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