leonarski_fandClaude Opus 5.5
dcad4104ce
Every dense map that needs geometry now evaluates each image pixel at its place on the detector plane (DetectorModules), so it is right for a stacked panel layout and picks up refined module placements: - AzimuthalIntegrationMapping: bins, resolution, solid-angle/polarisation corrections built per pixel through the table; a pixel on no panel is left as a masked one. The geometry cache key now includes the panel table, so a mapping built after the stills module placement is rebuilt rather than served from the cache. ResolutionMaskBits follows (it is a lookup). - ROIMap / ExportROIMap: azimuthal ROIs on the plane (box and circle ROIs stay in image pixels). - ROIIntegration CPU+GPU: the weighted sums are taken per ROI and panel in image pixels and placed on the plane at the end (exact; the affine commutes with a weighted mean). - HotPixelFinder: ring/sector keys and the zeta of the persistence rule on the plane. - Background beam-centre walk CPU+GPU: pixels binned at their plane position; host and device share BackgroundPlanePoint so they stay bit-identical. New common/PanelRows.h: ForEachPixelInRows, ForEachPixel cut to a band of rows in index order, for the row-parallel builders. An empty table takes the old arithmetic in the old order: rotation sets and every single-panel input are unchanged. [PanelLayout] AzintMapping and ResolutionMask enabled (+ azimuthal ROI map and ROI centroid on the plane); new BeamCenterFromBackground_BinsOnThePlane and ROIIntegrationGPU_MatchesCPU_Panels. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_015CeEGgpKdfkCTJb3YNMJcu
43 lines
1.8 KiB
C++
43 lines
1.8 KiB
C++
// SPDX-FileCopyrightText: 2026 Filip Leonarski, Paul Scherrer Institute <filip.leonarski@psi.ch>
|
|
// SPDX-License-Identifier: GPL-3.0-only
|
|
|
|
#pragma once
|
|
|
|
#include <memory>
|
|
|
|
#include "ROIIntegration.h"
|
|
#include "../indexing/CUDAMemHelpers.h"
|
|
|
|
class ROIIntegrationGPU : public ROIIntegration {
|
|
std::shared_ptr<CudaStream> stream;
|
|
int threads;
|
|
int blocks;
|
|
size_t shared_needed;
|
|
|
|
CudaDevicePtr<uint16_t> gpu_roi_map;
|
|
// 64-bit sums are accumulated as unsigned long long (two's-complement bit
|
|
// pattern) because CUDA atomicAdd has no signed 64-bit overload.
|
|
CudaDevicePtr<unsigned long long> gpu_sum;
|
|
CudaDevicePtr<unsigned long long> gpu_sum2;
|
|
CudaDevicePtr<unsigned long long> gpu_pixels;
|
|
CudaDevicePtr<unsigned long long> gpu_panel_v; // per ROI and panel, as roi_panel_v
|
|
CudaDevicePtr<unsigned long long> gpu_panel_x;
|
|
CudaDevicePtr<unsigned long long> gpu_panel_y;
|
|
CudaDevicePtr<int> gpu_panel_box; // x0, y0, width, height per panel; unallocated without panels
|
|
CudaDevicePtr<int> gpu_max;
|
|
|
|
std::vector<unsigned long long> host_sum;
|
|
std::vector<unsigned long long> host_sum2;
|
|
std::vector<unsigned long long> host_pixels;
|
|
std::vector<unsigned long long> host_panel_v;
|
|
std::vector<unsigned long long> host_panel_x;
|
|
std::vector<unsigned long long> host_panel_y;
|
|
std::vector<int> host_max;
|
|
std::vector<int> max_init; // INT_MIN seed copied into gpu_max each frame
|
|
public:
|
|
ROIIntegrationGPU(const DiffractionExperiment &experiment, std::shared_ptr<CudaStream> stream);
|
|
// Waits for the stream: its pooled buffers are freed next (see CudaDevicePtr).
|
|
~ROIIntegrationGPU() override { if (stream) cudaStreamSynchronize(*stream); }
|
|
void Run(const ImagePreprocessorBuffer &image, std::map<std::string, ROIMessage> &out) override;
|
|
};
|