MCS graph construction utils

src/CMakeLists.txt

@@ -17,6 +17,7 @@ add_subdirectory(conformer)
 add_subdirectory(data_structures)
 add_subdirectory(forcefields)
 add_subdirectory(minimizer)
+add_subdirectory(mcs)
 add_subdirectory(substruct)
 add_subdirectory(testutils)
 add_subdirectory(utils)

src/mcs/CMakeLists.txt

@@ -0,0 +1,19 @@
+# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+# All rights reserved. SPDX-License-Identifier: Apache-2.0
+#
+# Licensed under the Apache License, Version 2.0 (the "License"); you may not
+# use this file except in compliance with the License. You may obtain a copy of
+# the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
+# WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
+# License for the specific language governing permissions and limitations under
+# the License.
+
+add_library(mcs_fmcs_foundations fmcs_cuda/fmcs_match_tables.cu mcs_graph.cpp)
+target_include_directories(mcs_fmcs_foundations
+                           PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
+target_link_libraries(mcs_fmcs_foundations PUBLIC device_vector)

src/mcs/fmcs_cuda/fmcs_match_tables.cu

@@ -0,0 +1,59 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/mcs/fmcs_cuda/fmcs_match_tables.cuh"
+
+namespace mcs {
+namespace fmcs {
+
+UploadedPairMatchTables uploadPairMatchTables(const std::vector<PairMatchTablesHost>& host, cudaStream_t stream) {
+  size_t totalWords = 0;
+  for (const auto& p : host) {
+    totalWords += p.atoms.data.size();
+    totalWords += p.bonds.data.size();
+  }
+
+  UploadedPairMatchTables out;
+  out.storage = nvMolKit::AsyncDeviceVector<uint32_t>(totalWords, stream);
+  out.tables.resize(host.size());
+
+  std::vector<uint32_t> packed;
+  packed.reserve(totalWords);
+  uint32_t* base = out.storage.data();
+  for (size_t i = 0; i < host.size(); ++i) {
+    const auto& ph = host[i];
+    auto&       pd = out.tables[i];
+
+    if (!ph.atoms.data.empty()) {
+      const size_t offset = packed.size();
+      packed.insert(packed.end(), ph.atoms.data.begin(), ph.atoms.data.end());
+      pd.atoms = {base + offset, ph.atoms.nRows, ph.atoms.nCols, ph.atoms.wordsPerRow};
+    }
+
+    if (!ph.bonds.data.empty()) {
+      const size_t offset = packed.size();
+      packed.insert(packed.end(), ph.bonds.data.begin(), ph.bonds.data.end());
+      pd.bonds = {base + offset, ph.bonds.nRows, ph.bonds.nCols, ph.bonds.wordsPerRow};
+    }
+  }
+
+  if (!packed.empty()) {
+    out.storage.copyFromHost(packed);
+  }
+  return out;
+}
+
+}  // namespace fmcs
+}  // namespace mcs

src/mcs/fmcs_cuda/fmcs_match_tables.cuh

@@ -0,0 +1,91 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef FMCS_CUDA_FMCS_MATCH_TABLES_CUH
+#define FMCS_CUDA_FMCS_MATCH_TABLES_CUH
+
+#include <cuda_runtime.h>
+
+#include <cstddef>
+#include <cstdint>
+#include <vector>
+
+#include "src/mcs/mcs_common/mcs_types.cuh"
+#include "src/utils/device_vector.h"
+
+namespace mcs {
+namespace fmcs {
+
+/// Packed query-vs-target compatibility bitmap: row @c i covers query item
+/// @c i, column @c j covers target item @c j.  Row stride is @c wordsPerRow
+/// uint32 words.  Used for both atoms and bonds so every per-seed
+/// compatibility check on the device is a single word load.
+struct MatchTableHost {
+  int                   nRows       = 0;
+  int                   nCols       = 0;
+  int                   wordsPerRow = 0;
+  std::vector<uint32_t> data;
+
+  static int computeWordsPerRow(int nCols) { return (nCols + 31) / 32; }
+
+  void resize(int rows, int cols) {
+    nRows       = rows;
+    nCols       = cols;
+    wordsPerRow = computeWordsPerRow(cols);
+    data.assign(static_cast<size_t>(rows) * wordsPerRow, 0u);
+  }
+
+  void setBit(int row, int col) { data[static_cast<size_t>(row) * wordsPerRow + col / 32] |= (1u << (col % 32)); }
+
+  bool testBit(int row, int col) const {
+    return (data[static_cast<size_t>(row) * wordsPerRow + col / 32] >> (col % 32)) & 1u;
+  }
+};
+
+struct MatchTableDevice {
+  const uint32_t* data        = nullptr;
+  int             nRows       = 0;
+  int             nCols       = 0;
+  int             wordsPerRow = 0;
+
+  __host__ __device__ __forceinline__ bool testBit(int row, int col) const {
+    return (data[static_cast<size_t>(row) * wordsPerRow + col / 32] >> (col % 32)) & 1u;
+  }
+};
+
+struct PairMatchTablesHost {
+  MatchTableHost atoms;
+  MatchTableHost bonds;
+};
+
+struct PairMatchTablesDevice {
+  MatchTableDevice atoms;
+  MatchTableDevice bonds;
+};
+
+struct UploadedPairMatchTables {
+  nvMolKit::AsyncDeviceVector<uint32_t> storage;
+  std::vector<PairMatchTablesDevice>    tables;
+};
+
+/// Upload a batch of per-pair host match tables into one contiguous device
+/// allocation.  The returned object owns that allocation and keeps the
+/// per-pair device views valid for its lifetime.
+UploadedPairMatchTables uploadPairMatchTables(const std::vector<PairMatchTablesHost>& host, cudaStream_t stream);
+
+}  // namespace fmcs
+}  // namespace mcs
+
+#endif  // FMCS_CUDA_FMCS_MATCH_TABLES_CUH

src/mcs/mcs_common/mcs_types.cuh

@@ -35,6 +35,11 @@ struct Graph {
   std::vector<size_t> colIndices;       ///< Size = 2 * numEdges (symmetric).
 };
 
+/**
+ * @brief Build a CSR graph from a vertex count and undirected edge list.
+ */
+Graph buildGraphFromEdges(size_t numVertices, const std::vector<std::pair<size_t, size_t>>& edges);
+
 /**
  * @brief Result of a maximum common substructure computation.
  */

src/mcs/mcs_graph.cpp

@@ -0,0 +1,63 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <algorithm>
+#include <stdexcept>
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "src/mcs/mcs_common/mcs_types.cuh"
+
+namespace mcs {
+
+Graph buildGraphFromEdges(size_t numVertices, const std::vector<std::pair<size_t, size_t>>& edges) {
+  Graph graph;
+  graph.numVertices = static_cast<int>(numVertices);
+  graph.numEdges    = static_cast<int>(edges.size());
+  graph.rowOffsets.assign(numVertices + 1, 0);
+
+  for (const auto& [u, v] : edges) {
+    if (u >= numVertices || v >= numVertices) {
+      throw std::runtime_error("MCS graph edge endpoint out of range");
+    }
+    if (u == v) {
+      throw std::runtime_error("MCS graph self-loops are not supported");
+    }
+    ++graph.rowOffsets[u + 1];
+    ++graph.rowOffsets[v + 1];
+  }
+
+  for (size_t i = 1; i < graph.rowOffsets.size(); ++i) {
+    graph.rowOffsets[i] += graph.rowOffsets[i - 1];
+  }
+
+  graph.colIndices.assign(graph.rowOffsets.back(), 0);
+  std::vector<size_t> cursor = graph.rowOffsets;
+  for (const auto& [u, v] : edges) {
+    graph.colIndices[cursor[u]++] = v;
+    graph.colIndices[cursor[v]++] = u;
+  }
+
+  for (size_t atomIdx = 0; atomIdx < numVertices; ++atomIdx) {
+    const auto begin = graph.colIndices.begin() + static_cast<std::ptrdiff_t>(graph.rowOffsets[atomIdx]);
+    const auto end   = graph.colIndices.begin() + static_cast<std::ptrdiff_t>(graph.rowOffsets[atomIdx + 1]);
+    std::sort(begin, end);
+  }
+
+  return graph;
+}
+
+}  // namespace mcs

tests/CMakeLists.txt

@@ -72,6 +72,9 @@ target_include_directories(test_fmcs_primitives
                            PRIVATE ${CMAKE_SOURCE_DIR}/src/mcs)
 target_link_libraries(test_fmcs_primitives PRIVATE CUDA::cudart device_vector)
 
+add_executable(test_fmcs_foundations test_fmcs_foundations.cu)
+target_link_libraries(test_fmcs_foundations PRIVATE mcs_fmcs_foundations)
+
 add_executable(test_openmp_helpers test_openmp_helpers.cpp)
 target_link_libraries(test_openmp_helpers PRIVATE openmp_helpers
                                                   OpenMP::OpenMP_CXX)
@@ -417,6 +420,7 @@ set(TEST_LIST
     test_similarity
     test_thread_safe_queue
     test_fmcs_primitives
+    test_fmcs_foundations
     test_work_splitting
     test_fire_minimizer
     test_fire_minimizer_permol)

tests/test_fmcs_foundations.cu

@@ -0,0 +1,108 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+
+#include <cuda_runtime.h>
+#include <gtest/gtest.h>
+
+#include <cstdint>
+#include <utility>
+#include <vector>
+
+#include "src/mcs/fmcs_cuda/fmcs_match_tables.cuh"
+#include "src/mcs/mcs_common/mcs_types.cuh"
+#include "src/utils/cuda_error_check.h"
+
+namespace {
+
+using mcs::fmcs::MatchTableHost;
+using mcs::fmcs::PairMatchTablesHost;
+using nvMolKit::checkReturnCode;
+
+TEST(FMCSGraph, BuildsSortedSymmetricCsr) {
+  const auto graph = mcs::buildGraphFromEdges(4,
+                                              {
+                                                {2, 0},
+                                                {1, 2},
+                                                {3, 2}
+  });
+
+  EXPECT_EQ(graph.numVertices, 4);
+  EXPECT_EQ(graph.numEdges, 3);
+  EXPECT_EQ(graph.rowOffsets, (std::vector<std::size_t>{0, 1, 2, 5, 6}));
+  EXPECT_EQ(graph.colIndices, (std::vector<std::size_t>{2, 2, 0, 1, 3, 2}));
+}
+
+TEST(FMCSGraph, BuildsEmptyAndDisconnectedGraphs) {
+  const auto empty = mcs::buildGraphFromEdges(0, {});
+  EXPECT_EQ(empty.rowOffsets, (std::vector<std::size_t>{0}));
+  EXPECT_TRUE(empty.colIndices.empty());
+
+  const auto disconnected = mcs::buildGraphFromEdges(3, {});
+  EXPECT_EQ(disconnected.rowOffsets, (std::vector<std::size_t>{0, 0, 0, 0}));
+  EXPECT_TRUE(disconnected.colIndices.empty());
+}
+
+TEST(FMCSGraph, RejectsInvalidEdges) {
+  EXPECT_THROW((void)mcs::buildGraphFromEdges(2,
+                                              {
+                                                {0, 2}
+  }),
+               std::runtime_error);
+  EXPECT_THROW((void)mcs::buildGraphFromEdges(2,
+                                              {
+                                                {1, 1}
+  }),
+               std::runtime_error);
+}
+
+TEST(FMCSMatchTable, PacksAcrossWordBoundaries) {
+  MatchTableHost table;
+  table.resize(2, 65);
+
+  EXPECT_EQ(table.wordsPerRow, 3);
+  EXPECT_EQ(table.data.size(), 6);
+
+  table.setBit(0, 0);
+  table.setBit(0, 32);
+  table.setBit(1, 64);
+  EXPECT_TRUE(table.testBit(0, 0));
+  EXPECT_TRUE(table.testBit(0, 32));
+  EXPECT_TRUE(table.testBit(1, 64));
+  EXPECT_FALSE(table.testBit(1, 63));
+}
+
+TEST(FMCSMatchTable, UploadsPairsIntoOneContiguousBuffer) {
+  std::vector<PairMatchTablesHost> host(2);
+  host[0].atoms.resize(1, 2);
+  host[0].atoms.setBit(0, 1);
+  host[0].bonds.resize(1, 1);
+  host[0].bonds.setBit(0, 0);
+  host[1].atoms.resize(1, 33);
+  host[1].atoms.setBit(0, 32);
+
+  const auto  upload = mcs::fmcs::uploadPairMatchTables(host, nullptr);
+  const auto& device = upload.tables;
+
+  ASSERT_NE(upload.storage.data(), nullptr);
+  ASSERT_EQ(device.size(), 2);
+  EXPECT_EQ(upload.storage.size(), 4);
+  EXPECT_EQ(device[0].atoms.data, upload.storage.data());
+  EXPECT_EQ(device[0].bonds.data, upload.storage.data() + 1);
+  EXPECT_EQ(device[1].atoms.data, upload.storage.data() + 2);
+  EXPECT_EQ(device[1].atoms.wordsPerRow, 2);
+
+  std::vector<std::uint32_t> copied(4);
+  upload.storage.copyToHost(copied);
+  cudaCheckError(cudaStreamSynchronize(upload.storage.stream()));
+  EXPECT_EQ(copied, (std::vector<std::uint32_t>{2u, 1u, 0u, 1u}));
+}
+
+TEST(FMCSMatchTable, EmptyUploadReturnsNoAllocation) {
+  const auto upload = mcs::fmcs::uploadPairMatchTables({}, nullptr);
+
+  EXPECT_TRUE(upload.tables.empty());
+  EXPECT_EQ(upload.storage.data(), nullptr);
+  EXPECT_EQ(upload.storage.size(), 0);
+}
+
+}  // namespace