From 4fba5ccea7067be698b743a440ef9cb8bb368c23 Mon Sep 17 00:00:00 2001 From: Kevin Boyd Date: Fri, 24 Jul 2026 10:07:03 -0400 Subject: [PATCH 1/3] Add fMCS graph and match-table foundations --- src/CMakeLists.txt | 1 + src/mcs/CMakeLists.txt | 19 ++++ src/mcs/fmcs_cuda/fmcs_match_tables.cu | 100 +++++++++++++++++++++ src/mcs/fmcs_cuda/fmcs_match_tables.cuh | 90 +++++++++++++++++++ src/mcs/mcs_common/mcs_types.cuh | 5 ++ src/mcs/mcs_graph.cpp | 63 +++++++++++++ tests/CMakeLists.txt | 4 + tests/test_fmcs_foundations.cu | 112 ++++++++++++++++++++++++ 8 files changed, 394 insertions(+) create mode 100644 src/mcs/CMakeLists.txt create mode 100644 src/mcs/fmcs_cuda/fmcs_match_tables.cu create mode 100644 src/mcs/fmcs_cuda/fmcs_match_tables.cuh create mode 100644 src/mcs/mcs_graph.cpp create mode 100644 tests/test_fmcs_foundations.cu 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..c96a93c5 --- /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 CUDA::cudart) 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..ea92850a --- /dev/null +++ b/src/mcs/fmcs_cuda/fmcs_match_tables.cu @@ -0,0 +1,100 @@ +// 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 "fmcs_cuda/fmcs_match_tables.cuh" + +namespace mcs { +namespace fmcs { + +namespace { + +void checkCuda(cudaError_t err, const char* context) { + if (err != cudaSuccess) { + throw std::runtime_error(std::string("fMCS match-tables CUDA error at ") + context + ": " + + cudaGetErrorString(err)); + } +} + +} // namespace + +std::vector uploadPairMatchTables(const std::vector& host, + cudaStream_t stream, + void** outDeviceBuffer, + size_t* outDeviceBufferBytes) { + std::vector out(host.size()); + + size_t totalWords = 0; + for (const auto& p : host) { + totalWords += p.atoms.data.size(); + totalWords += p.bonds.data.size(); + } + + void* buf = nullptr; + if (totalWords > 0) { + checkCuda(cudaMallocAsync(&buf, totalWords * sizeof(uint32_t), stream), "cudaMallocAsync (match tables)"); + } + + size_t cursor = 0; + uint32_t* base = reinterpret_cast(buf); + for (size_t i = 0; i < host.size(); ++i) { + const auto& ph = host[i]; + auto& pd = out[i]; + + if (!ph.atoms.data.empty()) { + uint32_t* dst = base + cursor; + checkCuda(cudaMemcpyAsync(dst, + ph.atoms.data.data(), + ph.atoms.data.size() * sizeof(uint32_t), + cudaMemcpyHostToDevice, + stream), + "cudaMemcpyAsync (atoms)"); + pd.atoms = {dst, ph.atoms.nRows, ph.atoms.nCols, ph.atoms.wordsPerRow}; + cursor += ph.atoms.data.size(); + } + + if (!ph.bonds.data.empty()) { + uint32_t* dst = base + cursor; + checkCuda(cudaMemcpyAsync(dst, + ph.bonds.data.data(), + ph.bonds.data.size() * sizeof(uint32_t), + cudaMemcpyHostToDevice, + stream), + "cudaMemcpyAsync (bonds)"); + pd.bonds = {dst, ph.bonds.nRows, ph.bonds.nCols, ph.bonds.wordsPerRow}; + cursor += ph.bonds.data.size(); + } + } + + if (outDeviceBuffer) + *outDeviceBuffer = buf; + if (outDeviceBufferBytes) + *outDeviceBufferBytes = totalWords * sizeof(uint32_t); + return out; +} + +void freePairMatchTablesBuffer(void* deviceBuffer, cudaStream_t stream) { + if (deviceBuffer) { + cudaFreeAsync(deviceBuffer, stream); + } +} + +} // 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..635fcf44 --- /dev/null +++ b/src/mcs/fmcs_cuda/fmcs_match_tables.cuh @@ -0,0 +1,90 @@ +// 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 "mcs_common/mcs_types.cuh" + +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; +}; + +/// Upload a batch of per-pair host match tables into one contiguous device +/// buffer. The returned buffer is owned by the caller and must be released +/// with @ref freePairMatchTablesBuffer on the same stream. +std::vector uploadPairMatchTables(const std::vector& host, + cudaStream_t stream, + void** outDeviceBuffer, + std::size_t* outDeviceBufferBytes); + +void freePairMatchTablesBuffer(void* deviceBuffer, 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..753ab468 --- /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 "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..95162512 --- /dev/null +++ b/tests/test_fmcs_foundations.cu @@ -0,0 +1,112 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include +#include + +#include +#include +#include + +#include "fmcs_cuda/fmcs_match_tables.cuh" +#include "mcs_common/mcs_types.cuh" + +namespace { + +using mcs::fmcs::MatchTableHost; +using mcs::fmcs::PairMatchTablesHost; + +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); + + void* buffer = nullptr; + std::size_t bytes = 0; + const auto device = mcs::fmcs::uploadPairMatchTables(host, nullptr, &buffer, &bytes); + + ASSERT_NE(buffer, nullptr); + ASSERT_EQ(device.size(), 2); + EXPECT_EQ(bytes, 4 * sizeof(std::uint32_t)); + EXPECT_EQ(device[0].atoms.data, static_cast(buffer)); + EXPECT_EQ(device[0].bonds.data, static_cast(buffer) + 1); + EXPECT_EQ(device[1].atoms.data, static_cast(buffer) + 2); + EXPECT_EQ(device[1].atoms.wordsPerRow, 2); + + std::vector copied(4); + ASSERT_EQ(cudaMemcpy(copied.data(), buffer, bytes, cudaMemcpyDeviceToHost), cudaSuccess); + EXPECT_EQ(copied, (std::vector{2u, 1u, 0u, 1u})); + + mcs::fmcs::freePairMatchTablesBuffer(buffer, nullptr); + ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess); +} + +TEST(FMCSMatchTable, EmptyUploadReturnsNoAllocation) { + void* buffer = reinterpret_cast(1); + std::size_t bytes = 1; + const auto device = mcs::fmcs::uploadPairMatchTables({}, nullptr, &buffer, &bytes); + + EXPECT_TRUE(device.empty()); + EXPECT_EQ(buffer, nullptr); + EXPECT_EQ(bytes, 0); + mcs::fmcs::freePairMatchTablesBuffer(buffer, nullptr); +} + +} // namespace From 6d92ec621f8072e5335afa59c547b0f929760d99 Mon Sep 17 00:00:00 2001 From: Kevin Boyd Date: Fri, 24 Jul 2026 13:33:13 -0400 Subject: [PATCH 2/3] Use one RAII upload for fMCS match tables --- src/mcs/CMakeLists.txt | 2 +- src/mcs/fmcs_cuda/fmcs_match_tables.cu | 75 ++++++------------------- src/mcs/fmcs_cuda/fmcs_match_tables.cuh | 17 +++--- tests/test_fmcs_foundations.cu | 36 ++++++------ 4 files changed, 43 insertions(+), 87 deletions(-) diff --git a/src/mcs/CMakeLists.txt b/src/mcs/CMakeLists.txt index c96a93c5..1ca1cb56 100644 --- a/src/mcs/CMakeLists.txt +++ b/src/mcs/CMakeLists.txt @@ -16,4 +16,4 @@ 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 CUDA::cudart) +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 index ea92850a..eb08eebd 100644 --- a/src/mcs/fmcs_cuda/fmcs_match_tables.cu +++ b/src/mcs/fmcs_cuda/fmcs_match_tables.cu @@ -13,87 +13,46 @@ // See the License for the specific language governing permissions and // limitations under the License. -#include - -#include -#include -#include - #include "fmcs_cuda/fmcs_match_tables.cuh" namespace mcs { namespace fmcs { -namespace { - -void checkCuda(cudaError_t err, const char* context) { - if (err != cudaSuccess) { - throw std::runtime_error(std::string("fMCS match-tables CUDA error at ") + context + ": " + - cudaGetErrorString(err)); - } -} - -} // namespace - -std::vector uploadPairMatchTables(const std::vector& host, - cudaStream_t stream, - void** outDeviceBuffer, - size_t* outDeviceBufferBytes) { - std::vector out(host.size()); - +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(); } - void* buf = nullptr; - if (totalWords > 0) { - checkCuda(cudaMallocAsync(&buf, totalWords * sizeof(uint32_t), stream), "cudaMallocAsync (match tables)"); - } + UploadedPairMatchTables out; + out.storage = nvMolKit::AsyncDeviceVector(totalWords, stream); + out.tables.resize(host.size()); - size_t cursor = 0; - uint32_t* base = reinterpret_cast(buf); + 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[i]; + auto& pd = out.tables[i]; if (!ph.atoms.data.empty()) { - uint32_t* dst = base + cursor; - checkCuda(cudaMemcpyAsync(dst, - ph.atoms.data.data(), - ph.atoms.data.size() * sizeof(uint32_t), - cudaMemcpyHostToDevice, - stream), - "cudaMemcpyAsync (atoms)"); - pd.atoms = {dst, ph.atoms.nRows, ph.atoms.nCols, ph.atoms.wordsPerRow}; - cursor += ph.atoms.data.size(); + 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()) { - uint32_t* dst = base + cursor; - checkCuda(cudaMemcpyAsync(dst, - ph.bonds.data.data(), - ph.bonds.data.size() * sizeof(uint32_t), - cudaMemcpyHostToDevice, - stream), - "cudaMemcpyAsync (bonds)"); - pd.bonds = {dst, ph.bonds.nRows, ph.bonds.nCols, ph.bonds.wordsPerRow}; - cursor += ph.bonds.data.size(); + 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 (outDeviceBuffer) - *outDeviceBuffer = buf; - if (outDeviceBufferBytes) - *outDeviceBufferBytes = totalWords * sizeof(uint32_t); - return out; -} - -void freePairMatchTablesBuffer(void* deviceBuffer, cudaStream_t stream) { - if (deviceBuffer) { - cudaFreeAsync(deviceBuffer, stream); + if (!packed.empty()) { + out.storage.copyFromHost(packed); } + return out; } } // namespace fmcs diff --git a/src/mcs/fmcs_cuda/fmcs_match_tables.cuh b/src/mcs/fmcs_cuda/fmcs_match_tables.cuh index 635fcf44..46da4712 100644 --- a/src/mcs/fmcs_cuda/fmcs_match_tables.cuh +++ b/src/mcs/fmcs_cuda/fmcs_match_tables.cuh @@ -23,6 +23,7 @@ #include #include "mcs_common/mcs_types.cuh" +#include "src/utils/device_vector.h" namespace mcs { namespace fmcs { @@ -74,15 +75,15 @@ struct PairMatchTablesDevice { MatchTableDevice bonds; }; +struct UploadedPairMatchTables { + nvMolKit::AsyncDeviceVector storage; + std::vector tables; +}; + /// Upload a batch of per-pair host match tables into one contiguous device -/// buffer. The returned buffer is owned by the caller and must be released -/// with @ref freePairMatchTablesBuffer on the same stream. -std::vector uploadPairMatchTables(const std::vector& host, - cudaStream_t stream, - void** outDeviceBuffer, - std::size_t* outDeviceBufferBytes); - -void freePairMatchTablesBuffer(void* deviceBuffer, cudaStream_t stream); +/// 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 diff --git a/tests/test_fmcs_foundations.cu b/tests/test_fmcs_foundations.cu index 95162512..888ecc0a 100644 --- a/tests/test_fmcs_foundations.cu +++ b/tests/test_fmcs_foundations.cu @@ -10,11 +10,13 @@ #include "fmcs_cuda/fmcs_match_tables.cuh" #include "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, @@ -78,35 +80,29 @@ TEST(FMCSMatchTable, UploadsPairsIntoOneContiguousBuffer) { host[1].atoms.resize(1, 33); host[1].atoms.setBit(0, 32); - void* buffer = nullptr; - std::size_t bytes = 0; - const auto device = mcs::fmcs::uploadPairMatchTables(host, nullptr, &buffer, &bytes); + const auto upload = mcs::fmcs::uploadPairMatchTables(host, nullptr); + const auto& device = upload.tables; - ASSERT_NE(buffer, nullptr); + ASSERT_NE(upload.storage.data(), nullptr); ASSERT_EQ(device.size(), 2); - EXPECT_EQ(bytes, 4 * sizeof(std::uint32_t)); - EXPECT_EQ(device[0].atoms.data, static_cast(buffer)); - EXPECT_EQ(device[0].bonds.data, static_cast(buffer) + 1); - EXPECT_EQ(device[1].atoms.data, static_cast(buffer) + 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); - ASSERT_EQ(cudaMemcpy(copied.data(), buffer, bytes, cudaMemcpyDeviceToHost), cudaSuccess); + upload.storage.copyToHost(copied); + cudaCheckError(cudaStreamSynchronize(upload.storage.stream())); EXPECT_EQ(copied, (std::vector{2u, 1u, 0u, 1u})); - - mcs::fmcs::freePairMatchTablesBuffer(buffer, nullptr); - ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess); } TEST(FMCSMatchTable, EmptyUploadReturnsNoAllocation) { - void* buffer = reinterpret_cast(1); - std::size_t bytes = 1; - const auto device = mcs::fmcs::uploadPairMatchTables({}, nullptr, &buffer, &bytes); - - EXPECT_TRUE(device.empty()); - EXPECT_EQ(buffer, nullptr); - EXPECT_EQ(bytes, 0); - mcs::fmcs::freePairMatchTablesBuffer(buffer, nullptr); + 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 From 52a0e9b81ef2d431a6b4fad8fb79d3642d8fcec9 Mon Sep 17 00:00:00 2001 From: Kevin Boyd Date: Fri, 24 Jul 2026 13:44:47 -0400 Subject: [PATCH 3/3] Use project-rooted includes in fMCS foundations --- src/mcs/fmcs_cuda/fmcs_match_tables.cu | 2 +- src/mcs/fmcs_cuda/fmcs_match_tables.cuh | 2 +- src/mcs/mcs_graph.cpp | 2 +- tests/test_fmcs_foundations.cu | 4 ++-- 4 files changed, 5 insertions(+), 5 deletions(-) diff --git a/src/mcs/fmcs_cuda/fmcs_match_tables.cu b/src/mcs/fmcs_cuda/fmcs_match_tables.cu index eb08eebd..1f17f94c 100644 --- a/src/mcs/fmcs_cuda/fmcs_match_tables.cu +++ b/src/mcs/fmcs_cuda/fmcs_match_tables.cu @@ -13,7 +13,7 @@ // See the License for the specific language governing permissions and // limitations under the License. -#include "fmcs_cuda/fmcs_match_tables.cuh" +#include "src/mcs/fmcs_cuda/fmcs_match_tables.cuh" namespace mcs { namespace fmcs { diff --git a/src/mcs/fmcs_cuda/fmcs_match_tables.cuh b/src/mcs/fmcs_cuda/fmcs_match_tables.cuh index 46da4712..af18adfc 100644 --- a/src/mcs/fmcs_cuda/fmcs_match_tables.cuh +++ b/src/mcs/fmcs_cuda/fmcs_match_tables.cuh @@ -22,7 +22,7 @@ #include #include -#include "mcs_common/mcs_types.cuh" +#include "src/mcs/mcs_common/mcs_types.cuh" #include "src/utils/device_vector.h" namespace mcs { diff --git a/src/mcs/mcs_graph.cpp b/src/mcs/mcs_graph.cpp index 753ab468..0de2763f 100644 --- a/src/mcs/mcs_graph.cpp +++ b/src/mcs/mcs_graph.cpp @@ -19,7 +19,7 @@ #include #include -#include "mcs_common/mcs_types.cuh" +#include "src/mcs/mcs_common/mcs_types.cuh" namespace mcs { diff --git a/tests/test_fmcs_foundations.cu b/tests/test_fmcs_foundations.cu index 888ecc0a..ae3e6c88 100644 --- a/tests/test_fmcs_foundations.cu +++ b/tests/test_fmcs_foundations.cu @@ -8,8 +8,8 @@ #include #include -#include "fmcs_cuda/fmcs_match_tables.cuh" -#include "mcs_common/mcs_types.cuh" +#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 {