Files
Jungfraujoch/image_analysis/spot_finding/ImageSpotFinderGPU.h
T
leonarski_fandClaude Opus 5.5 60ca872bc2 GPU adaptive spot finder: same spots, about half the kernel time per frame
Measured on an RTX 5080, one worker (nsys), 16M EIGER rotation set, per frame:
spot-finding kernels 757 -> ~383 us. Output is bit-identical: integer sums
and a deterministic candidate/spot order throughout.

- flag_strong reads the image only where a strong pixel can be: the plain
  ring pass notes each 32-pixel word's largest value, a geometry table gives
  the rings each word touches, and a word below the lowest of their
  thresholds is skipped. One thread per word writes the whole word, so the
  ring buffer needs no clearing or atomics. 128 -> 12 us.
- analyze_pixel with candidates uses 16-row tiles (only tiles near a
  candidate are worked, and each warp's serial chain of rows is short) and
  issues the entering/leaving row loads before the window scans; the window
  count is summed in 32 bits. 254 -> 72 us. The dense path keeps its tiling.
- analyze_candidates reads its column window with unconditional, unrolled
  loads; through value_at each load waited for its previous-pass bit and
  the row before. 74 -> 34 us.
- scatter_bits emits a block's range with the whole warp (prefix sum over
  the words' bit counts, same order). 32 -> 6 us.
- finish_components split into label (one block), sum (all blocks, integer
  atomics) and filter/compact (one block); a component's sums no longer walk
  the list from its first member to its last. 26 -> 12 us.

Tried and not kept (slower on sm_120): warp-aggregated ring sums and a
low-value histogram fast path in reduce_rings_shared, a precomputed
tile-needed map, unrolling analyze_pixel's row loop.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01EizBKhTqrAYago9KJG3dA3
2026-10-11 04:29:41 +02:00

53 lines
2.7 KiB
C++

// SPDX-FileCopyrightText: 2025 Filip Leonarski, Paul Scherrer Institute <filip.leonarski@psi.ch>
// SPDX-License-Identifier: GPL-3.0-only
#pragma once
#include <vector>
#include "SpotFindingSettings.h"
#include "ImageSpotFinder.h"
#include "SpotExtractorGPU.h"
#include "../indexing/CUDAMemHelpers.h"
class ImageSpotFinderGPU : public ImageSpotFinder {
protected:
// Protected rather than private because AdaptiveSpotFinderGPU derives from this engine: it is
// this same local-box detection with the fixed photon floor replaced by a per-resolution-ring
// one, so it reuses the stream, the bit buffers and the extractor rather than owning a second
// set of them.
std::shared_ptr<CudaStream> stream;
CudaDevicePtr<uint32_t> gpu_out_0;
CudaDevicePtr<uint32_t> gpu_out_1; // holds the strong-pixel bits after Detect()
SpotExtractorGPU extractor;
private:
const int numberOfCudaThreads = 128; // #threads per block of analyze_pixel (one warp per 32 columns)
const int numberOfWaves = 32; // #row bands of analyze_pixel
const int sparseRowsPerWave = 16; // rows per band of analyze_pixel when only near candidates
const int windowSizeLimit = 32; // limit on the window size (2nby+1, 2nbx+1): a warp holds 2 x 32 columns
int candidateBlocks = 0; // grid of analyze_candidates (256 threads per block)
void RunDetect(const ImagePreprocessorBuffer &image, const SpotFindingSettings &settings,
const uint32_t *gpu_candidates);
protected:
// Detect, but with the result wanted only at the candidate pixels (device bit buffer, the layout
// of the output): gpu_out_1 ends up holding (Detect's result & candidates), and the second pass is
// evaluated at the candidates alone. Asynchronous - it does not wait for the stream.
void DetectAt(const ImagePreprocessorBuffer &image, const SpotFindingSettings &settings,
const uint32_t *gpu_candidates);
public:
ImageSpotFinderGPU(int32_t width, int32_t height, std::shared_ptr<CudaStream> stream);
// Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr).
~ImageSpotFinderGPU() override { if (stream) cudaStreamSynchronize(*stream); }
void Detect(const ImagePreprocessorBuffer &image, const SpotFindingSettings &settings) override;
void SetResolutionMaskBits(const std::vector<uint32_t> &packed_mask) override;
[[nodiscard]] uint32_t StrongPixelCount() const override { return extractor.StrongPixelCount(); }
const std::vector<DiffractionSpot> &ExtractComponents(const ImagePreprocessorBuffer &image,
const SpotFindingSettings &settings) override;
};