// 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 GPU and get DiffractionSpot list via StrongPixelSet -> FindSpotsImage 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); REQUIRE(cudaMemcpyAsync(buffer.getGPUBuffer(), buffer.getBuffer().data(), width * height * sizeof(int32_t), cudaMemcpyHostToDevice, *stream) == cudaSuccess); return gpu.Run(buffer, settings, res_mask); } // 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; FindSpotsImage uses mask/resolution implicit in StrongPixelSet. // StrongPixelSet doesn't carry resolution/mask by itself, but FindSpotsImage(settings, vec) // 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; FindSpotsImage uses mask/resolution implicit in StrongPixelSet. // StrongPixelSet doesn't carry resolution/mask by itself, but FindSpotsImage(settings, vec) // 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; FindSpotsImage uses mask/resolution implicit in StrongPixelSet. // StrongPixelSet doesn't carry resolution/mask by itself, but FindSpotsImage(settings, vec) // 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)); const auto cpu_spots = cpu.Run(cpu_buffer, settings, res_mask); 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); } #endif