// SPDX-FileCopyrightText: 2025 Filip Leonarski, Paul Scherrer Institute // SPDX-License-Identifier: GPL-3.0-only #include #include "../common/CUDAWrapper.h" #ifdef JFJOCH_USE_CUDA #include "../image_analysis/spot_finding/ImageSpotFinderGPU.h" #include "../image_analysis/spot_finding/ImageSpotFinderCPU.h" #include "../image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.h" static void fill_test_image(ImagePreprocessorBuffer& buffer, size_t width, size_t height) { for (size_t i = 0; i < width * height; i++) buffer[i] = (i % 2) * 5 + 5; buffer[width * 50 + 50] = 20; buffer[width * 25 + 26] = 16; buffer[width * 75 + 25] = 12; } // Helper to run the GPU finder and collect its spot list static std::vector run_gpu_and_collect_spots(ImagePreprocessorBufferGPU &buffer, size_t width, size_t height, const SpotFindingSettings &settings, const std::vector &res_mask) { auto stream = std::make_shared(); ImageSpotFinderGPU gpu(static_cast(width), static_cast(height), stream); REQUIRE(get_gpu_count() > 0); gpu.SetResolutionMask(res_mask); REQUIRE(cudaMemcpyAsync(buffer.getGPUBuffer(), buffer.getBuffer().data(), width * height * sizeof(int32_t), cudaMemcpyHostToDevice, *stream) == cudaSuccess); return gpu.Run(buffer, settings); } // Mirror of ImageSpotFinder_SignalToNoise TEST_CASE("ImageSpotFinderGPU_SignalToNoise") { if (get_gpu_count() == 0) { WARN("No CUDA GPU present. Skipping ImageSpotFinderGPU_SignalToNoise"); return; } const size_t width = 100, height = 100; std::vector res_mask(width * height, false); std::vector mask(width * height, false); ImagePreprocessorBufferGPU buffer(width * height); fill_test_image(buffer, width, height); SpotFindingSettings settings{ .signal_to_noise_threshold = 3.0, .photon_count_threshold = 0, .min_pix_per_spot = 1, .max_pix_per_spot = 20, .high_resolution_limit = 0.5, .low_resolution_limit = 3.0, }; // GPU produces strong pixels; the resolution mask is handed to the finder separately and the // connected-component search then runs on the device (SpotExtractorGPU). The spot-level filter // matches CPU ImageSpotFinder test behavior for these synthetic inputs. auto spots = run_gpu_and_collect_spots(buffer, width, height, settings, res_mask); REQUIRE(spots.size() == 2); REQUIRE(spots[0].RawCoord().y == 25); REQUIRE(spots[1].RawCoord().y == 50); } TEST_CASE("ImageSpotFinderGPU_CountThreshold") { if (get_gpu_count() == 0) { WARN("No CUDA GPU present. Skipping ImageSpotFinderGPU_CountThreshold"); return; } const size_t width = 100, height = 100; std::vector res_mask(width * height, false); std::vector mask(width * height, false); ImagePreprocessorBufferGPU buffer(width * height); fill_test_image(buffer, width, height); SpotFindingSettings settings{ .signal_to_noise_threshold = 0.0, .photon_count_threshold = 11, .min_pix_per_spot = 1, .max_pix_per_spot = 20, .high_resolution_limit = 0.5, .low_resolution_limit = 3.0, }; // GPU produces strong pixels; the resolution mask is handed to the finder separately and the // connected-component search then runs on the device (SpotExtractorGPU). The spot-level filter // matches CPU ImageSpotFinder test behavior for these synthetic inputs. auto spots = run_gpu_and_collect_spots(buffer, width, height, settings, res_mask); REQUIRE(spots.size() == 3); REQUIRE(spots[0].RawCoord().y == 25); REQUIRE(spots[1].RawCoord().y == 50); REQUIRE(spots[2].RawCoord().y == 75); } TEST_CASE("ImageSpotFinderGPU_20M") { if (get_gpu_count() == 0) { WARN("No CUDA GPU present. Skipping ImageSpotFinderGPU_20M"); return; } const size_t width = 4500, height = 4500; std::vector res_mask(width * height, false); std::vector mask(width * height, false); ImagePreprocessorBufferGPU buffer(width * height); fill_test_image(buffer, width, height); SpotFindingSettings settings{ .signal_to_noise_threshold = 3.0, .photon_count_threshold = 0, .min_pix_per_spot = 1, .max_pix_per_spot = 20, .high_resolution_limit = 0.5, .low_resolution_limit = 3.0, }; // GPU produces strong pixels; the resolution mask is handed to the finder separately and the // connected-component search then runs on the device (SpotExtractorGPU). The spot-level filter // matches CPU ImageSpotFinder test behavior for these synthetic inputs. auto spots = run_gpu_and_collect_spots(buffer, width, height, settings, res_mask); REQUIRE(spots.size() == 2); REQUIRE(spots[0].RawCoord().y == 25); REQUIRE(spots[1].RawCoord().y == 50); } // The two finders must return the same spots for the same frame - a dataset processed on a machine // without a GPU has to give the same answer as one processed with it. // // The spots here are deliberately broad. Both finders measure a pixel against a 31x31 local // background, so a blob several pixels across sits inside its own background window and inflates // the mean and variance it is tested against. That is what the second pass exists to undo: it // recomputes the background with the pixels found strong by the first pass excluded. A single-pass // finder loses the outer pixels of every broad spot, so this comparison fails unless both sides run // the same two passes. TEST_CASE("ImageSpotFinder_CPU_GPU_Parity", "[ImageSpotFinder]") { if (get_gpu_count() == 0) SKIP("No CUDA GPU present"); const size_t width = 100, height = 100; ImagePreprocessorBufferGPU gpu_buffer(width * height); ImagePreprocessorBuffer cpu_buffer(width * height); // Background alternating 5/10 (mean 7.5, sd 2.5), plus two spots shaped to make the second pass // matter: a bright 5x5 core (300) with a thin one-pixel ring around it (25). The ring is well // above the clean background, but the core sitting inside the ring's own 31x31 window drags that // window's mean to ~16 and its sd to ~46, so on a single pass the ring fails the SNR test and the // spot comes out as the 25-pixel core. The second pass takes the core out of the background and // the ring passes, giving 49 pixels. The ring is kept thin on purpose: a wide halo would swamp // its own background and stay undetectable either way. auto fill = [&](ImagePreprocessorBuffer &b) { for (size_t i = 0; i < width * height; i++) b[i] = (i % 2) * 5 + 5; const struct { int cx, cy; } spots[] = {{50, 50}, {22, 74}}; for (const auto &s : spots) { for (int dy = -3; dy <= 3; dy++) { for (int dx = -3; dx <= 3; dx++) { const bool ring = std::abs(dx) == 3 || std::abs(dy) == 3; b[(s.cy + dy) * width + (s.cx + dx)] = ring ? 25 : 300; } } } }; fill(gpu_buffer); fill(cpu_buffer); SpotFindingSettings settings{ .signal_to_noise_threshold = 3.0, .photon_count_threshold = 0, .min_pix_per_spot = 1, .max_pix_per_spot = 1000, .high_resolution_limit = 0.5, .low_resolution_limit = 3.0, }; const std::vector res_mask(width * height, false); ImageSpotFinderCPU cpu(static_cast(width), static_cast(height)); cpu.SetResolutionMask(res_mask); const auto cpu_spots = cpu.Run(cpu_buffer, settings); const auto gpu_spots = run_gpu_and_collect_spots(gpu_buffer, width, height, settings, res_mask); REQUIRE(cpu_spots.size() == gpu_spots.size()); REQUIRE(cpu_spots.size() == 2); // guard against both finding nothing and "agreeing" for (size_t i = 0; i < cpu_spots.size(); i++) { CHECK(cpu_spots[i].RawCoord().x == Catch::Approx(gpu_spots[i].RawCoord().x)); CHECK(cpu_spots[i].RawCoord().y == Catch::Approx(gpu_spots[i].RawCoord().y)); // The pixel set is what the second pass changes, so compare it rather than the centroid, // which stays put for a symmetric spot whether or not the halo was picked up. CHECK(cpu_spots[i].PixelCount() == gpu_spots[i].PixelCount()); CHECK(cpu_spots[i].Count() == gpu_spots[i].Count()); } // Both finders must reach past the bright core into the ring: 49 pixels, not the 25 of the core // alone. Without this the comparison above would still pass if both ran a single pass. CHECK(cpu_spots[0].PixelCount() == 49); } // The same comparison on a TALL image, which is what makes it bite. // // The GPU kernel splits the image into ImageSpotFinderGPU::numberOfWaves (32) horizontal waves and // each wave walks its rows in three stages: a priming loop, a main loop, and a drain loop. The main // loop only runs while front < rmax, and front starts NBX+1 rows ahead of the wave's first row - so // for an image of 100 rows, where a wave owns ceil(100/32) = 4 rows, THE MAIN LOOP NEVER EXECUTES. // Every row goes through priming/drain. On a real detector frame (4362 rows -> 137 rows per wave) // the main loop instead carries ~121 of every 137 rows, i.e. almost the whole image. // // That is why a bug living only in the main loop survived: it was unreachable at the test's size. // This case is 1024 rows (32 per wave, ~16 through the main loop) so the path is covered, and the // spots are placed deep inside a wave rather than at its edges. Values above 46340 are included // deliberately: the window keeps a sum of squares, and squaring in 32 bits overflows past that, // while a real detector saturates around 1e6. TEST_CASE("ImageSpotFinder_CPU_GPU_Parity_TallImage", "[ImageSpotFinder]") { if (get_gpu_count() == 0) SKIP("No CUDA GPU present"); const size_t width = 128, height = 1024; ImagePreprocessorBufferGPU gpu_buffer(width * height); ImagePreprocessorBuffer cpu_buffer(width * height); // Spots at rows a wave reaches through its MAIN loop, not its priming or drain rows. const struct { int cx, cy; int32_t core; } spots[] = { {64, 80, 300}, // ordinary bright spot {64, 400, 100000}, // above 46340: 32-bit squaring wraps, 64-bit does not {64, 720, 300}, }; auto fill = [&](ImagePreprocessorBuffer &b) { for (size_t i = 0; i < width * height; i++) b[i] = (i % 2) * 5 + 5; for (const auto &s : spots) for (int dy = -3; dy <= 3; dy++) for (int dx = -3; dx <= 3; dx++) { const bool ring = std::abs(dx) == 3 || std::abs(dy) == 3; b[(s.cy + dy) * width + (s.cx + dx)] = ring ? 25 : s.core; } }; fill(gpu_buffer); fill(cpu_buffer); SpotFindingSettings settings{ .signal_to_noise_threshold = 3.0, .photon_count_threshold = 0, .min_pix_per_spot = 1, .max_pix_per_spot = 1000, .high_resolution_limit = 0.5, .low_resolution_limit = 3.0, }; const std::vector res_mask(width * height, false); ImageSpotFinderCPU cpu(static_cast(width), static_cast(height)); cpu.SetResolutionMask(res_mask); const auto cpu_spots = cpu.Run(cpu_buffer, settings); const auto gpu_spots = run_gpu_and_collect_spots(gpu_buffer, width, height, settings, res_mask); REQUIRE(cpu_spots.size() == 3); // all three found, so "agreeing on nothing" cannot pass REQUIRE(gpu_spots.size() == cpu_spots.size()); for (size_t i = 0; i < cpu_spots.size(); i++) { INFO("spot " << i); CHECK(cpu_spots[i].RawCoord().x == Catch::Approx(gpu_spots[i].RawCoord().x)); CHECK(cpu_spots[i].RawCoord().y == Catch::Approx(gpu_spots[i].RawCoord().y)); CHECK(cpu_spots[i].PixelCount() == gpu_spots[i].PixelCount()); CHECK(cpu_spots[i].Count() == gpu_spots[i].Count()); CHECK(cpu_spots[i].PixelCount() == 49); // core + ring, i.e. the second pass did its job } } // The local-box SNR test on the knife edge, where the GPU used to answer differently from the CPU. // Both compare (in_minus_mean)^2 against var * strong2 computed in float; the CPU compares the // integer square with the ceiling of that float, the GPU compared the square converted to float - // which, for squares past 2^24, can round onto the threshold and fail a pixel the CPU accepts. // // The centre of a checkerboard of 0 and 57 holding 171 is such a pixel at 5 sigma: its 960 // neighbours are 480 zeros and 480 of 57, so var = 748569600 and in_minus_mean = 136800, whose // square 18714240000 is var * 25 exactly. var * 25 rounds down in float to 18714238976, 1024 below // the square - half a float spacing - so the square rounds onto the threshold too. The CPU takes // the pixel (18714240000 > 18714238976), the GPU used not to. // // The second image is the count threshold: at photon_count_threshold = 0 the CPU still asks for a // positive value, the GPU took 0 to mean no threshold. A 0 above a negative background (the // checkerboard of -3 and -1, as a pedestal-subtracted detector gives) is 2 sigma above its mean. TEST_CASE("ImageSpotFinder_CPU_GPU_Parity_KnifeEdge", "[ImageSpotFinder]") { if (get_gpu_count() == 0) SKIP("No CUDA GPU present"); const size_t width = 64, height = 64; const std::vector res_mask(width * height, false); auto run_both = [&](int32_t even, int32_t odd, int32_t centre, float snr, size_t expected) { ImagePreprocessorBufferGPU gpu_buffer(width * height); ImagePreprocessorBuffer cpu_buffer(width * height); for (size_t row = 0; row < height; row++) for (size_t col = 0; col < width; col++) { const int32_t v = ((row + col) % 2 == 0) ? even : odd; gpu_buffer[row * width + col] = v; cpu_buffer[row * width + col] = v; } gpu_buffer[32 * width + 32] = centre; cpu_buffer[32 * width + 32] = centre; SpotFindingSettings settings{ .signal_to_noise_threshold = snr, .photon_count_threshold = 0, .min_pix_per_spot = 1, .max_pix_per_spot = 1000, .high_resolution_limit = 0.5, .low_resolution_limit = 3.0, }; ImageSpotFinderCPU cpu(static_cast(width), static_cast(height)); cpu.SetResolutionMask(res_mask); const auto cpu_spots = cpu.Run(cpu_buffer, settings); const auto gpu_spots = run_gpu_and_collect_spots(gpu_buffer, width, height, settings, res_mask); CHECK(cpu_spots.size() == expected); CHECK(gpu_spots.size() == cpu_spots.size()); }; SECTION("SNR test with the square past 2^24") { run_both(0, 57, 171, 5.0f, 1); } SECTION("Zero count threshold on a negative background") { run_both(-3, -1, 0, 1.5f, 0); } } #endif