ImagePreprocessorBuffer: Dedicated class to handle buffer used by spot finding.
Build Packages / build:rpm (ubuntu2404_nocuda) (push) Successful in 15m13s
Build Packages / build:rpm (rocky9_nocuda) (push) Successful in 15m49s
Build Packages / build:rpm (rocky8_sls9) (push) Successful in 16m45s
Build Packages / build:rpm (rocky8) (push) Successful in 17m36s
Build Packages / build:rpm (rocky9_sls9) (push) Successful in 18m27s
Build Packages / build:rpm (rocky9) (push) Successful in 18m28s
Build Packages / Generate python client (push) Successful in 1m31s
Build Packages / Unit tests (push) Successful in 1h0m56s
Build Packages / build:rpm (ubuntu2204) (push) Failing after 9m13s
Build Packages / Create release (push) Has been skipped
Build Packages / Build documentation (push) Successful in 1m52s
Build Packages / XDS test (durin plugin) (push) Successful in 10m18s
Build Packages / XDS test (neggia plugin) (push) Successful in 8m22s
Build Packages / XDS test (JFJoch plugin) (push) Successful in 9m24s
Build Packages / build:rpm (ubuntu2404) (push) Successful in 13m6s
Build Packages / DIALS test (push) Successful in 13m14s
Build Packages / build:rpm (rocky8_nocuda) (push) Successful in 11m15s
Build Packages / build:rpm (ubuntu2204_nocuda) (push) Failing after 14m5s

This commit is contained in:
2026-04-23 20:20:24 +02:00
parent e51a45a5ea
commit 17619bd19a
26 changed files with 223 additions and 117 deletions
@@ -94,7 +94,6 @@ ImagePreprocessorGPU::ImagePreprocessorGPU(const DiffractionExperiment &experime
stream(stream),
gpu_mask(npixels),
gpu_decompressed_image(npixels * sizeof(uint32_t)), // Overshoot - if input image is 1- or 2-byte, then it is still fine, while memory loss is minimal
gpu_image(npixels),
gpu_stats(1),
cpu_stats(1),
cpu_stats_reg(cpu_stats) {
@@ -113,7 +112,7 @@ ImagePreprocessorGPU::ImagePreprocessorGPU(const DiffractionExperiment &experime
blocks = 4 * prop.multiProcessorCount;
}
ImageStatistics ImagePreprocessorGPU::Analyze(std::vector<int32_t> &processed_image, const uint8_t *image_ptr,
ImageStatistics ImagePreprocessorGPU::Analyze(ImagePreprocessorBuffer &processed_image, const uint8_t *image_ptr,
CompressedImageMode image_mode) {
switch (image_mode) {
case CompressedImageMode::Int8:
@@ -134,7 +133,7 @@ ImageStatistics ImagePreprocessorGPU::Analyze(std::vector<int32_t> &processed_im
}
template<class T>
ImageStatistics ImagePreprocessorGPU::Analyze(std::vector<int32_t> &processed_image,
ImageStatistics ImagePreprocessorGPU::Analyze(ImagePreprocessorBuffer &processed_image,
const uint8_t *input,
T err_value,
T sat_value) {
@@ -148,23 +147,15 @@ ImageStatistics ImagePreprocessorGPU::Analyze(std::vector<int32_t> &processed_im
preprocess_kernel<T> <<< blocks, threads, 0, *stream >>>(
reinterpret_cast<const T *>(gpu_decompressed_image.get()),
gpu_mask,
gpu_image,
processed_image.getGPUBuffer(),
gpu_stats,
sat_value,
err_value,
npixels);
cudaMemcpyAsync(processed_image.data(), gpu_image, npixels * sizeof(int32_t), cudaMemcpyDeviceToHost, *stream);
cudaMemcpyAsync(processed_image.data(), processed_image.getGPUBuffer(), npixels * sizeof(int32_t), cudaMemcpyDeviceToHost, *stream);
cudaMemcpyAsync(cpu_stats.data(), gpu_stats, sizeof(ImageStatistics), cudaMemcpyDeviceToHost, *stream);
cudaStreamSynchronize(*stream);
return cpu_stats[0];
}
const int32_t *ImagePreprocessorGPU::GetImageDevicePtr() const {
return gpu_image;
}
std::shared_ptr<CudaStream> ImagePreprocessorGPU::GetStream() const {
return stream;
}