Skip to content

MCS Fast match pathway #238

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
scal444 merged 4 commits into from
Jul 30, 2026
Merged

MCS Fast match pathway #238

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
27ac684
Add fMCS fast incremental matching
scal444 Jul 24, 2026
d982a2b
Use project-rooted includes in fMCS matching
scal444 Jul 24, 2026
2c14740
Drop the topology-generality concept from fMCS matching
scal444 Jul 28, 2026
b12721c
Clarify greedy fMCS match contract
scal444 Jul 28, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
22 changes: 11 additions & 11 deletions src/mcs/fmcs_cuda/fmcs_grow.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,14 +19,12 @@
#include <cstdint>

#include "src/mcs/fmcs_cuda/fmcs_seed.cuh"
#include "src/mcs/fmcs_cuda/fmcs_topology.cuh"
#include "src/mcs/mcs_common/mcs_cooperative_copy.cuh"

namespace mcs {
namespace fmcs {

constexpr int kGrowBondEndpointShift = 16;
constexpr std::uint32_t kGrowBondEndpointMask = 0xFFFFu;

template <int maxAtoms, int maxBonds>
__device__ __forceinline__ void seedAddNewBondWithinThread(Seed<maxAtoms, maxBonds>& seed, const NewBond& bond) {
seedAddBondWithinThread(seed, bond.bondIdx);
Expand Down Expand Up @@ -68,10 +66,10 @@ __device__ __forceinline__ void seedAddNewBondWithinThread(Seed<maxAtoms, maxBon
/// and clamps @p outCount to @p maxNewBonds (slots beyond that bound
/// are not written). No I/O ordering on @p outBonds is guaranteed --
/// the order depends on lane race outcomes.
template <int maxAtoms, int maxBonds, class QueryTopology, class GroupT>
template <int maxAtoms, int maxBonds, class GroupT>
__device__ __forceinline__ bool fillNewBondsCooperative(const GroupT& group,
const Seed<maxAtoms, maxBonds>& seed,
const QueryTopology& queryTopology,
const DeviceCsrView& queryTopology,
NewBond* outBonds,
int* outCount,
int maxNewBonds) {
Expand All @@ -96,8 +94,8 @@ __device__ __forceinline__ bool fillNewBondsCooperative(const GroupT&

// Decode this query bond's endpoints.
const std::uint32_t queryEndpoints = queryTopology.bondEndpoints[q];
const int queryEndpointU = static_cast<int>(queryEndpoints >> kGrowBondEndpointShift);
const int queryEndpointV = static_cast<int>(queryEndpoints & kGrowBondEndpointMask);
const int queryEndpointU = static_cast<int>(queryEndpoints >> kBondEndpointShift);
const int queryEndpointV = static_cast<int>(queryEndpoints & kBondEndpointMask);

// Bond is a "new boundary" candidate iff at least one endpoint
// was added in the most recent grow step.
Expand Down Expand Up @@ -153,10 +151,12 @@ __device__ __forceinline__ bool fillNewBondsCooperative(const GroupT&
/// across all lanes of @p group), and within each iteration the only
/// truly group-parallel step is the @c warpCopy of @c parent into
/// @p childWorkspace. Lane 0 then patches in the singleton bond /
/// atom; the per-bond match call is whatever @p matchFn does (when
/// the kernel passes @ref matchIncrementalFastCooperative the inner
/// target-bond scan IS lane-parallel); @p childSink is typically lane
/// 0 only (cache.insert + queue.push).
/// atom; the per-bond match call is whatever @p matchFn does. The
/// kernel's callback may try @ref tryMatchIncrementalGreedyCooperative,
/// but it must own the exact fallback too: a greedy miss alone must not
/// prune the bond. The greedy helper's inner target-bond scan is
/// lane-parallel; @p childSink is typically lane 0 only (cache.insert +
/// queue.push).
///
/// In other words: this function is cooperative only for the
/// parent->child copy and for delegating to its callbacks; it does
Expand Down
Loading
Loading