diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index c81d8c7c..0b1c1984 100644 --- a/src/CMakeLists.txt +++ b/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) diff --git a/src/mcs/CMakeLists.txt b/src/mcs/CMakeLists.txt new file mode 100644 index 00000000..1ca1cb56 --- /dev/null +++ b/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) diff --git a/src/mcs/fmcs_cuda/fmcs_match_tables.cu b/src/mcs/fmcs_cuda/fmcs_match_tables.cu new file mode 100644 index 00000000..1f17f94c --- /dev/null +++ b/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& 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(totalWords, stream); + out.tables.resize(host.size()); + + std::vector 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 diff --git a/src/mcs/fmcs_cuda/fmcs_match_tables.cuh b/src/mcs/fmcs_cuda/fmcs_match_tables.cuh new file mode 100644 index 00000000..af18adfc --- /dev/null +++ b/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 + +#include +#include +#include + +#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 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(rows) * wordsPerRow, 0u); + } + + void setBit(int row, int col) { data[static_cast(row) * wordsPerRow + col / 32] |= (1u << (col % 32)); } + + bool testBit(int row, int col) const { + return (data[static_cast(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(row) * wordsPerRow + col / 32] >> (col % 32)) & 1u; + } +}; + +struct PairMatchTablesHost { + MatchTableHost atoms; + MatchTableHost bonds; +}; + +struct PairMatchTablesDevice { + MatchTableDevice atoms; + MatchTableDevice bonds; +}; + +struct UploadedPairMatchTables { + nvMolKit::AsyncDeviceVector storage; + std::vector 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& host, cudaStream_t stream); + +} // namespace fmcs +} // namespace mcs + +#endif // FMCS_CUDA_FMCS_MATCH_TABLES_CUH diff --git a/src/mcs/mcs_common/mcs_types.cuh b/src/mcs/mcs_common/mcs_types.cuh index 0be638e4..6c63d3d2 100644 --- a/src/mcs/mcs_common/mcs_types.cuh +++ b/src/mcs/mcs_common/mcs_types.cuh @@ -35,6 +35,11 @@ struct Graph { std::vector 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>& edges); + /** * @brief Result of a maximum common substructure computation. */ diff --git a/src/mcs/mcs_graph.cpp b/src/mcs/mcs_graph.cpp new file mode 100644 index 00000000..0de2763f --- /dev/null +++ b/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 +#include +#include +#include +#include + +#include "src/mcs/mcs_common/mcs_types.cuh" + +namespace mcs { + +Graph buildGraphFromEdges(size_t numVertices, const std::vector>& edges) { + Graph graph; + graph.numVertices = static_cast(numVertices); + graph.numEdges = static_cast(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 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(graph.rowOffsets[atomIdx]); + const auto end = graph.colIndices.begin() + static_cast(graph.rowOffsets[atomIdx + 1]); + std::sort(begin, end); + } + + return graph; +} + +} // namespace mcs diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 2bb4fd26..28c432e5 100644 --- a/tests/CMakeLists.txt +++ b/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) diff --git a/tests/test_fmcs_foundations.cu b/tests/test_fmcs_foundations.cu new file mode 100644 index 00000000..ae3e6c88 --- /dev/null +++ b/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 +#include + +#include +#include +#include + +#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{0, 1, 2, 5, 6})); + EXPECT_EQ(graph.colIndices, (std::vector{2, 2, 0, 1, 3, 2})); +} + +TEST(FMCSGraph, BuildsEmptyAndDisconnectedGraphs) { + const auto empty = mcs::buildGraphFromEdges(0, {}); + EXPECT_EQ(empty.rowOffsets, (std::vector{0})); + EXPECT_TRUE(empty.colIndices.empty()); + + const auto disconnected = mcs::buildGraphFromEdges(3, {}); + EXPECT_EQ(disconnected.rowOffsets, (std::vector{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 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 copied(4); + upload.storage.copyToHost(copied); + cudaCheckError(cudaStreamSynchronize(upload.storage.stream())); + EXPECT_EQ(copied, (std::vector{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