Build Packages / build:viewer-tgz:cpu (push) Successful in 7m46s
Build Packages / build:viewer-tgz:cuda (push) Successful in 9m14s
Build Packages / build:rpm (ubuntu2404_nocuda) (push) Successful in 13m51s
Build Packages / build:rpm (rocky8_nocuda) (push) Successful in 14m17s
Build Packages / build:rpm (ubuntu2204_nocuda) (push) Successful in 14m14s
Build Packages / build:rpm (rocky9_nocuda) (push) Successful in 14m43s
Build Packages / build:rpm (rocky8_sls9) (push) Successful in 14m45s
Build Packages / build:rpm (rocky8) (push) Successful in 11m44s
Build Packages / build:rpm (rocky9_sls9) (push) Successful in 13m24s
Build Packages / XDS test (durin plugin) (push) Successful in 8m33s
Build Packages / Generate python client (push) Successful in 28s
Build Packages / Build documentation (push) Successful in 1m4s
Build Packages / Create release (push) Skipped
Build Packages / build:rpm (rocky9) (push) Successful in 12m45s
Build Packages / build:rpm (ubuntu2204) (push) Successful in 12m25s
Build Packages / build:rpm (ubuntu2404) (push) Successful in 13m1s
Build Packages / DIALS test (push) Successful in 14m29s
Build Packages / XDS test (neggia plugin) (push) Successful in 8m17s
Build Packages / XDS test (JFJoch plugin) (push) Successful in 9m5s
Build Packages / Unit tests (push) Successful in 1h16m19s
Build Packages / build:windows:nocuda (push) Failing after 2s
Build Packages / build:windows:cuda (push) Failing after 3s
The spot finder flagged strong pixels on the device and then labelled them on the host, so every frame sent the packed bitmask back - 2.26 MB on a large detector - and the host walked all of it to recover a few hundred pixels. Do the labelling on the device instead: compact the bitmask into a flat-index-sorted list, find each pixel's backward neighbours by binary search, union them lock-free with path halving, then label, accumulate and filter in one kernel. Only the spot list comes back, and only one stream synchronisation per frame. The gain in the ordinary case is modest - about a quarter off per-image spot finding - because the host algorithm is genuinely fast on a normal frame. What justifies it is the frame that is not ordinary. The host labels a sorted sparse list through a window spanning two detector lines, so its cost is quadratic in how many strong pixels share a line. A lit band of detector rows - a hot module, a panel edge - costs 33 ms at two rows and 377 ms at fifteen, all of it under the pixel cap that was supposed to bound this, and none of it maskable when the cause is a diffraction ring rather than a defect: a ring runs tangent to a row at its top and bottom, which is exactly the shape that hurts. The device version is flat at 0.05 to 0.64 ms across every geometry tried, so an online run no longer stalls a quarter of a second on an ice ring. Rejecting an over-cap frame is now free too, since the count is known before any pixel is written. Also label once and filter three times. The per-image minimum-pixel search runs the extraction at three settings, but that setting only decides which components are kept - it does not change the components - so the search itself need not be repeated. This helps the host path as much as the device one. The resolution mask moves to the device as a bit mask, uploaded when the limits change rather than per frame, since the compaction needs it there. Parity is asserted permanently rather than argued: five cases covering realistic frames, occupancy from a hundred pixels to past the cap, the pathological geometries including rings, the resolution mask, and a hundred-repeat determinism check - requiring the same partition, the same spot order, and identical counts. The centroid is a float sum and therefore order-dependent, so the device walks each component from its root in ascending order and fuses its multiply-add the way the host's does; note that whether the host fuses at all depends on the architecture flags, so exact centroid equality is asserted where the compiler fuses and a two-ulp bound otherwise. Making those accumulators integer would remove that dependence entirely and is worth doing separately. Regression set: all 37 crystals identical to the last printed digit. Unit suite passes with the new cases. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
214 lines
8.8 KiB
C++
214 lines
8.8 KiB
C++
// SPDX-FileCopyrightText: 2025 Filip Leonarski, Paul Scherrer Institute
|
|
// SPDX-License-Identifier: GPL-3.0-only
|
|
|
|
#include <catch2/catch_all.hpp>
|
|
#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<DiffractionSpot> run_gpu_and_collect_spots(ImagePreprocessorBufferGPU &buffer,
|
|
size_t width, size_t height,
|
|
const SpotFindingSettings &settings,
|
|
const std::vector<bool> &res_mask) {
|
|
auto stream = std::make_shared<CudaStream>();
|
|
ImageSpotFinderGPU gpu(static_cast<int32_t>(width), static_cast<int32_t>(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<bool> res_mask(width * height, false);
|
|
std::vector<bool> 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<bool> res_mask(width * height, false);
|
|
std::vector<bool> 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<bool> res_mask(width * height, false);
|
|
std::vector<bool> 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<bool> res_mask(width * height, false);
|
|
|
|
ImageSpotFinderCPU cpu(static_cast<int32_t>(width), static_cast<int32_t>(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);
|
|
}
|
|
|
|
#endif
|