From 228f0fc00f7251f7b245b683ddf8d951b4b9fca7 Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Wed, 14 Oct 2026 22:35:48 +0200 Subject: [PATCH] SpotTiles_MatchWholeFrame_GPU: wait for the upload before the squares are cut cudaMemcpy from pageable memory may return before the copy lands on the card, and ImagePreprocessorBufferGPU gathers its squares on a non-blocking stream that does not wait for the NULL stream, so on a loaded GPU the squares could be cut from the previous frame. Test-only race; the pre-scan's engines synchronise their own streams before the buffer is read. Co-Authored-By: Claude Opus 5.5 --- tests/SpotTilesTest.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/tests/SpotTilesTest.cpp b/tests/SpotTilesTest.cpp index 7da0047a0..d5c2f5d12 100644 --- a/tests/SpotTilesTest.cpp +++ b/tests/SpotTilesTest.cpp @@ -135,6 +135,9 @@ TEST_CASE("SpotTiles_MatchWholeFrame_GPU", "[SpotWidth][GPU]") { const auto upload = [&](Frame &f) -> ImagePreprocessorBuffer & { REQUIRE(cudaMemcpy(device.getGPUBuffer(), f.image.data(), f.image.size() * sizeof(int32_t), cudaMemcpyHostToDevice) == cudaSuccess); + // From pageable memory cudaMemcpy may return before the copy has landed, and the buffer cuts + // its squares on a non-blocking stream that does not wait for the NULL stream. + REQUIRE(cudaDeviceSynchronize() == cudaSuccess); return device; }; CheckFrames(frames, upload, SPOT_TILE_HALF, FOOTPRINT_TILE_HALF);