From b94d7b9bbd91e2c98b2ba3da1b641be8e8bc171c Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Mon, 28 Sep 2026 19:04:02 +0200 Subject: [PATCH] ModelMaskGPUTest: upload on the mask's stream cudaMemcpy from pageable memory can return before the DMA lands, and the legacy stream it runs on does not order a non-blocking stream, so RemoveIslands could read a partly uploaded mask. Both uploads are now cudaMemcpyAsync on the test's stream. Co-Authored-By: Claude Opus 5.5 (1M context) Claude-Session: https://claude.ai/code/session_01D1G8gJVAy6gp1K5Dz3NE5C --- tests/ModelMaskGPUTest.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/ModelMaskGPUTest.cpp b/tests/ModelMaskGPUTest.cpp index 793187cd6..e5717c555 100644 --- a/tests/ModelMaskGPUTest.cpp +++ b/tests/ModelMaskGPUTest.cpp @@ -143,7 +143,7 @@ TEST_CASE("ModelMaskGPU_MatchesGemmi", "[ModelValidation][gpu]") { const std::vector atoms = MaskAtoms(st); const std::vector ops = MaskOps(*sg); CudaDevicePtr atoms_d(atoms.size()); - REQUIRE(cudaMemcpy(atoms_d, atoms.data(), atoms.size() * sizeof(ModelMaskAtom), cudaMemcpyHostToDevice) == cudaSuccess); + REQUIRE(cudaMemcpyAsync(atoms_d, atoms.data(), atoms.size() * sizeof(ModelMaskAtom), cudaMemcpyHostToDevice, stream) == cudaSuccess); for (double d_min : {6.0, 3.5}) { const gemmi::SolventMasker masker(gemmi::AtomicRadiiSet::Refmac); @@ -164,7 +164,7 @@ TEST_CASE("ModelMaskGPU_MatchesGemmi", "[ModelValidation][gpu]") { std::vector out(n), first; // The island step alone, on gemmi's pre-island mask: exact. - REQUIRE(cudaMemcpy(mask_d, pre.data.data(), n * sizeof(float), cudaMemcpyHostToDevice) == cudaSuccess); + REQUIRE(cudaMemcpyAsync(mask_d, pre.data.data(), n * sizeof(float), cudaMemcpyHostToDevice, stream) == cudaSuccess); mask.RemoveIslands(mask_d); REQUIRE(cudaMemcpyAsync(out.data(), mask_d, n * sizeof(float), cudaMemcpyDeviceToHost, stream) == cudaSuccess); REQUIRE(cudaStreamSynchronize(stream) == cudaSuccess);