diff --git a/common/ROIAzimuthal.cpp b/common/ROIAzimuthal.cpp index 4c0d2103a..fd7a36689 100644 --- a/common/ROIAzimuthal.cpp +++ b/common/ROIAzimuthal.cpp @@ -18,9 +18,12 @@ ROIAzimuthal::ROIAzimuthal(const std::string &in_name, float in_d_min_A, float i : ROIElement(in_name), phi_min_deg(NormalizePhi_deg(in_phi_min_deg)), phi_max_deg(NormalizePhi_deg(in_phi_max_deg)) { - if ((in_d_min_A <= 0) || (in_d_max_A <= 0)) + // isfinite, not just > 0: NaN passes every <= test and inf passes > 0 (see ROICircle). + if (!std::isfinite(in_d_min_A) || !std::isfinite(in_d_max_A) + || (in_d_min_A <= 0) || (in_d_max_A <= 0) + || !std::isfinite(in_phi_min_deg) || !std::isfinite(in_phi_max_deg)) throw JFJochException(JFJochExceptionCategory::InputParameterInvalid, - "Resolution cannot be zero or negative"); + "Resolution and angles must be positive finite numbers"); if (in_d_min_A > in_d_max_A) { d_max_A = in_d_min_A; d_min_A = in_d_max_A; diff --git a/common/ROICircle.cpp b/common/ROICircle.cpp index 4fd72ec42..9be781020 100644 --- a/common/ROICircle.cpp +++ b/common/ROICircle.cpp @@ -1,14 +1,21 @@ // SPDX-FileCopyrightText: 2024 Filip Leonarski, Paul Scherrer Institute // SPDX-License-Identifier: GPL-3.0-only +#include + #include "ROICircle.h" #include "JFJochException.h" ROICircle::ROICircle(const std::string &name, float in_x, float in_y, float in_r_pxl) : ROIElement(name), center_x(in_x), center_y(in_y), r_pxl(in_r_pxl), r_pxl_2(in_r_pxl * in_r_pxl) { - if (r_pxl <= 0.0) + // isfinite, not just > 0: NaN passes every <= test, inf passes > 0, and both then reach the + // preview drawing, where a non-finite bound becomes an endless loop. + if (!std::isfinite(r_pxl) || r_pxl <= 0.0) throw JFJochException(JFJochExceptionCategory::InputParameterInvalid, - "ROICircle: radius must be positive number"); + "ROICircle: radius must be a positive finite number"); + if (!std::isfinite(center_x) || !std::isfinite(center_y)) + throw JFJochException(JFJochExceptionCategory::InputParameterInvalid, + "ROICircle: center must be finite"); } float ROICircle::GetX() const { diff --git a/image_analysis/azint/AzIntEngineGPU.cu b/image_analysis/azint/AzIntEngineGPU.cu index 9ea95342c..b35adaa95 100644 --- a/image_analysis/azint/AzIntEngineGPU.cu +++ b/image_analysis/azint/AzIntEngineGPU.cu @@ -130,11 +130,13 @@ void AzIntEngineGPU::Run(const ImagePreprocessorBuffer &image, AzimuthalIntegrat gpu_pixel_to_bin->get(),gpu_azint_correction->get(),image.getGPUBuffer(), gpu_sum, gpu_sum2, gpu_count, npixel, azint_bins ); + cuda_err(cudaGetLastError()); } else { gpu_azim<<>>( gpu_pixel_to_bin->get(),gpu_azint_correction->get(),image.getGPUBuffer(), gpu_sum, gpu_sum2, gpu_count, npixel, azint_bins ); + cuda_err(cudaGetLastError()); } cudaMemcpyAsync(azint_sum.data(), gpu_sum, sizeof(float) * azint_bins, cudaMemcpyDeviceToHost, *stream); diff --git a/image_analysis/beam_stop/ShadowAccumulatorGPU.cu b/image_analysis/beam_stop/ShadowAccumulatorGPU.cu index be7117b0b..240c83d86 100644 --- a/image_analysis/beam_stop/ShadowAccumulatorGPU.cu +++ b/image_analysis/beam_stop/ShadowAccumulatorGPU.cu @@ -61,6 +61,7 @@ void launch(const uint8_t *raw, int nframes, size_t frame_bytes, int64_t *max_va accumulate_kernel<<>>(reinterpret_cast(raw), nframes, frame_bytes / sizeof(T), max_value, sum_value, valid_count, npixels, masked); + cuda_err(cudaGetLastError()); } } // namespace diff --git a/image_analysis/beam_stop/ShadowFinder.cpp b/image_analysis/beam_stop/ShadowFinder.cpp index 28bb9f695..a0d3559a3 100644 --- a/image_analysis/beam_stop/ShadowFinder.cpp +++ b/image_analysis/beam_stop/ShadowFinder.cpp @@ -530,6 +530,15 @@ std::vector ShadowFinder::GetMask(size_t nthreads) const { if (frames == 0) return mask; + // The rings are meaningless about a centre that is not a real position - a NaN centre would + // even turn lround below into a negative array index - and a centre many detector sizes away + // makes max_radius, and every per-ring table sized by it, arbitrarily large. No centre, no + // shadow. + const float diag = std::hypot(static_cast(W), static_cast(H)); + if (!std::isfinite(beam_x) || !std::isfinite(beam_y) + || std::fabs(beam_x - W * 0.5f) > 4.0f * diag || std::fabs(beam_y - H * 0.5f) > 4.0f * diag) + return mask; + // mean projection, usable pixels and radius from the beam centre. The mean is divided by the // polarization factor, so that what is left varies around a ring only where something is in the // way; `pol` is kept because the counts the Poisson test is made of are the ones that were @@ -554,7 +563,11 @@ std::vector ShadowFinder::GetMask(size_t nthreads) const { pol[i] = pol_geometry.CalcAzIntPolarizationCorr(static_cast(x), static_cast(y), polarization.value()); - if (valid_count[i] > 0 && pixel_mask[i] == 0) { + // pol > 0 also rejects a NaN factor. The correction is exactly zero only where + // the polarization fully suppresses the scattering, so such a pixel carries no + // usable ring information anyway - and divided by zero it would put an inf into + // the pooled means, which the running box sums then turn into NaN for a whole row. + if (valid_count[i] > 0 && pixel_mask[i] == 0 && pol[i] > 0.0f) { mean[i] = static_cast(static_cast(sum_value[i]) / valid_count[i] / pol[i]); valid[i] = 1; diff --git a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu index c224c2870..ae320ec35 100644 --- a/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu +++ b/image_analysis/bragg_integration/BraggIntegrationEngineGPU.cu @@ -957,39 +957,49 @@ std::vector BraggIntegrationEngineGPU::Run(const ImagePreprocessorBu cuda_err(cudaMemsetAsync(d_rad_cnt, 0, sizeof(int) * n_rad, *stream)); } reset<<<32, 256, 0, *stream>>>(d_shell_grid, d_global_grid, d_mom, d_shell_n, d_global_n, d_invd2, GG); + cuda_err(cudaGetLastError()); dirty = true; mark_mask<<>>(d_px_x, d_px_y, d_mask, d_owner, p, n, 0); + cuda_err(cudaGetLastError()); boxsum<<>>(d_px_x, d_px_y, d_d, img, d_mask, d_owner, p, n, d_cx, d_cy, d_I, d_sigma, d_bkg, d_bkg_var, d_var_bkg, d_obs_x, d_obs_y, d_ok, d_strong, d_has_obs, d_invd2, d_isum, d_ninner, d_rbin, d_kbin, d_rad_sum, d_rad_cnt, rad_n, d_counts); + cuda_err(cudaGetLastError()); // Correct the flat annulus background for the curvature of the radial background before anything // downstream (profile fit, variance) reads it. - if (rad_n > 0) + if (rad_n > 0) { radial_correct<<<(n + threads - 1) / threads, threads, 0, *stream>>>( d_rad_sum, d_rad_cnt, rad_n, d_k_diff, k_len, k_off, d_isum, d_ninner, d_rbin, d_kbin, d_ok, d_bkg, d_I, n); + cuda_err(cudaGetLastError()); + } if (mode != IntegratorMode::BoxSum) { // Pass B: learn (shell computed inline) -> build -> fit. learn_profile<<>>(img, d_owner, d_px_x, d_px_y, d_cx, d_cy, d_d, d_invd2, d_I, d_bkg, d_ok, d_strong, d_shell_grid, d_global_grid, d_mom, d_shell_n, d_global_n, p, n); + cuda_err(cudaGetLastError()); build_profiles<<>>( d_shell_grid, d_global_grid, d_mom, d_shell_n, d_global_n, d_shell_P, d_global_P, d_sigma2_r, d_sigma2_t, p); + cuda_err(cudaGetLastError()); fit<<>>(img, d_owner, d_px_x, d_px_y, d_cx, d_cy, d_d, d_invd2, d_I, d_sigma, d_bkg, d_bkg_var, d_var_bkg, d_ok, d_shell_P, d_global_P, d_sigma2_r, d_sigma2_t, d_shell_n, d_I, d_sigma, d_var_bkg, d_ok, p, n, d_counts); + cuda_err(cudaGetLastError()); } // Nothing below reads the mask or the owner map, so put them back now, over the boxes that were // marked rather than over the frame - where that is the cheaper of the two. - if (clear_by_box) + if (clear_by_box) { mark_mask<<>>(d_px_x, d_px_y, d_mask, d_owner, p, n, 1); + cuda_err(cudaGetLastError()); + } cuda_err(cudaMemcpyAsync(h_I.get(), d_I, sizeof(float) * npredicted, cudaMemcpyDeviceToHost, *stream)); cuda_err(cudaMemcpyAsync(h_sigma.get(), d_sigma, sizeof(float) * npredicted, cudaMemcpyDeviceToHost, *stream)); diff --git a/image_analysis/bragg_prediction/BraggPredictionGPU.cu b/image_analysis/bragg_prediction/BraggPredictionGPU.cu index a42befe5f..457aa33b6 100644 --- a/image_analysis/bragg_prediction/BraggPredictionGPU.cu +++ b/image_analysis/bragg_prediction/BraggPredictionGPU.cu @@ -10,6 +10,11 @@ #include "../indexing/CUDAMemHelpers.h" #include +static inline void cuda_err(cudaError_t val) { + if (val != cudaSuccess) + throw JFJochException(JFJochExceptionCategory::GPUCUDAError, cudaGetErrorString(val)); +} + namespace { // Number of bandwidth sigmas included in the (radially thickened) Ewald-shell // acceptance window. Mirrors the CPU BraggPrediction path. @@ -245,6 +250,7 @@ int BraggPredictionGPU::Calc(const DiffractionExperiment &experiment, (2 * settings.max_l + 1 + block.z - 1) / block.z); bragg_kernel_3d<<>>(dK, settings.max_h, settings.max_k, settings.max_l, max_reflections, d_out, d_count); + cuda_err(cudaGetLastError()); // Async D2H count and synchronize cudaMemcpyAsync(h_count, d_count, sizeof(int), cudaMemcpyDeviceToHost, stream); @@ -255,6 +261,7 @@ int BraggPredictionGPU::Calc(const DiffractionExperiment &experiment, GrowCapacity(count); // see the rotation predictor cudaMemsetAsync(d_count, 0, sizeof(int), stream); bragg_kernel_3d<<>>(dK, settings.max_h, settings.max_k, settings.max_l, max_reflections, d_out, d_count); + cuda_err(cudaGetLastError()); cudaMemcpyAsync(h_count, d_count, sizeof(int), cudaMemcpyDeviceToHost, stream); cudaStreamSynchronize(stream); count = std::min(*h_count.get(), max_reflections); diff --git a/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu b/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu index d53d2c366..59855cf14 100644 --- a/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu +++ b/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu @@ -11,6 +11,11 @@ #include #include +static inline void cuda_err(cudaError_t val) { + if (val != cudaSuccess) + throw JFJochException(JFJochExceptionCategory::GPUCUDAError, cudaGetErrorString(val)); +} + namespace { __host__ __device__ inline bool is_odd(int v) { return (v & 1) != 0; } @@ -339,6 +344,7 @@ int BraggPredictionRotGPU::Calc(const DiffractionExperiment &experiment, (2 * settings.max_l + 1 + block.z - 1) / block.z); bragg_rot_kernel_3d<<>>(dK, settings.max_h, settings.max_k, settings.max_l, max_reflections, d_out, d_count); + cuda_err(cudaGetLastError()); cudaMemcpyAsync(h_count, d_count, sizeof(int), cudaMemcpyDeviceToHost, stream); cudaStreamSynchronize(stream); @@ -351,6 +357,7 @@ int BraggPredictionRotGPU::Calc(const DiffractionExperiment &experiment, GrowCapacity(count); cudaMemsetAsync(d_count, 0, sizeof(int), stream); bragg_rot_kernel_3d<<>>(dK, settings.max_h, settings.max_k, settings.max_l, max_reflections, d_out, d_count); + cuda_err(cudaGetLastError()); cudaMemcpyAsync(h_count, d_count, sizeof(int), cudaMemcpyDeviceToHost, stream); cudaStreamSynchronize(stream); count = std::min(*h_count.get(), max_reflections); diff --git a/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu b/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu index e45f97dd3..8f8fe5381 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu +++ b/image_analysis/image_preprocessing/ImagePreprocessorBufferGPU.cu @@ -3,6 +3,11 @@ #include "ImagePreprocessorBufferGPU.h" +static inline void cuda_err(cudaError_t val) { + if (val != cudaSuccess) + throw JFJochException(JFJochExceptionCategory::GPUCUDAError, cudaGetErrorString(val)); +} + __global__ void gather_kernel(const int32_t *__restrict__ image, const uint32_t *__restrict__ npixel, int32_t *__restrict__ values, @@ -39,6 +44,7 @@ void ImagePreprocessorBufferGPU::Gather(const std::vector &npixel, std cudaMemcpyHostToDevice, gather_stream); gather_kernel<<<(count + 255) / 256, 256, 0, gather_stream>>>( gpu_image.get(), gpu_gather_index.get(), gpu_gather_value.get(), count); + cuda_err(cudaGetLastError()); cudaMemcpyAsync(values.data(), gpu_gather_value.get(), count * sizeof(int32_t), cudaMemcpyDeviceToHost, gather_stream); cudaStreamSynchronize(gather_stream); diff --git a/image_analysis/roi/ROIIntegrationGPU.cu b/image_analysis/roi/ROIIntegrationGPU.cu index db215670e..d652c9297 100644 --- a/image_analysis/roi/ROIIntegrationGPU.cu +++ b/image_analysis/roi/ROIIntegrationGPU.cu @@ -140,6 +140,7 @@ void ROIIntegrationGPU::Run(const ImagePreprocessorBuffer &image, std::map>>( gpu_roi_map, image.getGPUBuffer(), npixel, width, roi_count, gpu_sum, gpu_sum2, gpu_pixels, gpu_x_weighted, gpu_y_weighted, gpu_max); + cuda_err(cudaGetLastError()); cudaMemcpyAsync(host_sum.data(), gpu_sum, sizeof(unsigned long long) * roi_count, cudaMemcpyDeviceToHost, *stream); cudaMemcpyAsync(host_sum2.data(), gpu_sum2, sizeof(unsigned long long) * roi_count, cudaMemcpyDeviceToHost, *stream); diff --git a/image_analysis/scale_merge/RotationScaleMergeGPU.cu b/image_analysis/scale_merge/RotationScaleMergeGPU.cu index 74920e3f9..40e146dc5 100644 --- a/image_analysis/scale_merge/RotationScaleMergeGPU.cu +++ b/image_analysis/scale_merge/RotationScaleMergeGPU.cu @@ -822,11 +822,14 @@ void RotationScaleMergeGPU::ScalePartials(int iters, double robust_k, double min ReduceGroupMeansKernel<<>>(d.n_groups, min_partiality, d.group_perm.get(), d.group_start.get(), d.group_count.get(), d.I.get(), d.sigma.get(), d.partiality.get(), d.corr.get(), d.group_mean.get()); + CudaCheck(cudaGetLastError(), "ReduceGroupMeansKernel launch"); PrepScaleObsKernel<<>>(d.n_obs, d.group.get(), d.partiality.get(), d.prescaling_corr.get(), d.zeta.get(), d.on_ice.get(), d.group_mean.get(), d.sigma.get(), d.inv_sigma.get(), d.sco_coeff.get(), d.sco_ok.get()); + CudaCheck(cudaGetLastError(), "PrepScaleObsKernel launch"); FitPerFrameGKernel<<>>(d.n_frames, robust_k, d.frame_start.get(), d.frame_count.get(), d.I.get(), d.inv_sigma.get(), d.sco_coeff.get(), d.sco_ok.get(), nullptr, d.g.get(), d.scaled.get()); + CudaCheck(cudaGetLastError(), "FitPerFrameGKernel launch"); UpdateCorrKernel<<>>(d.n_obs, d.frame.get(), d.prescaling_corr.get(), d.partiality.get(), d.g.get(), d.scaled.get(), d.corr.get()); } @@ -1027,6 +1030,7 @@ void RotationScaleMergeGPU::ComputePartialCC(double min_partiality, double *cc_o ReduceGroupMeansKernel<<>>(d.n_groups, min_partiality, d.group_perm.get(), d.group_start.get(), d.group_count.get(), d.I.get(), d.sigma.get(), d.partiality.get(), d.corr.get(), d.group_mean.get()); + CudaCheck(cudaGetLastError(), "ReduceGroupMeansKernel launch"); PerFrameCCKernel<<>>(d.n_frames, min_partiality, d.frame_start.get(), d.frame_count.get(), d.I.get(), d.sigma.get(), d.partiality.get(), d.corr.get(), d.on_ice.get(), d.group.get(), d.group_mean.get(), d.cc.get(), d.cc_n.get()); @@ -1180,9 +1184,13 @@ void RotationScaleMergeGPU::ScaleFulls(int iters, double robust_k, double min_pa // Unity model: partiality/prescaling_corr/zeta = 1 so coeff = mean; corr starts at 1. FillKernel<<>>(d.f_corr.get(), nf, 1.0f); + CudaCheck(cudaGetLastError(), "FillKernel launch"); FillKernel<<>>(d.f_partiality.get(), nf, 1.0f); + CudaCheck(cudaGetLastError(), "FillKernel launch"); FillKernel<<>>(d.f_rlp.get(), nf, 1.0f); + CudaCheck(cudaGetLastError(), "FillKernel launch"); FillKernel<<>>(d.f_zeta.get(), nf, 1.0f); + CudaCheck(cudaGetLastError(), "FillKernel launch"); // g/scaled reset once (a frame, once fitted, stays fitted - same as ScalePartials). CudaCheck(cudaMemset(d.scaled.get(), 0, size_t(d.n_frames) * sizeof(uint8_t)), "memset f scaled"); CudaCheck(cudaMemset(d.g.get(), 0, size_t(d.n_frames) * sizeof(double)), "memset f g"); @@ -1191,15 +1199,18 @@ void RotationScaleMergeGPU::ScaleFulls(int iters, double robust_k, double min_pa ReduceGroupMeansKernel<<>>(d.n_groups, min_partiality, d.f_gperm.get(), d.f_gstart.get(), d.f_gcount.get(), d.f_I.get(), d.f_sigma.get(), d.f_partiality.get(), d.f_corr.get(), d.group_mean.get()); + CudaCheck(cudaGetLastError(), "ReduceGroupMeansKernel launch"); // Not grid-stride, so its grid has to cover every full - unlike the grid-stride kernels // below, which the 65535 cap is there for. Capped, it would silently leave the tail of // sco_coeff/sco_ok stale above 16.8M fulls. PrepScaleObsKernel<<<(nf + BLK - 1) / BLK, BLK>>>(nf, d.f_group.get(), d.f_partiality.get(), d.f_rlp.get(), d.f_zeta.get(), d.f_on_ice.get(), d.group_mean.get(), d.f_sigma.get(), d.f_inv_sigma.get(), d.f_sco_coeff.get(), d.f_sco_ok.get()); + CudaCheck(cudaGetLastError(), "PrepScaleObsKernel launch"); FitPerFrameGKernel<<>>(d.n_frames, robust_k, d.f_frame_start.get(), d.f_frame_count.get(), d.f_I.get(), d.f_inv_sigma.get(), d.f_sco_coeff.get(), d.f_sco_ok.get(), d.f_frame_perm.get(), d.g.get(), d.scaled.get()); + CudaCheck(cudaGetLastError(), "FitPerFrameGKernel launch"); UpdateCorrKernel<<>>(nf, d.f_frame.get(), d.f_rlp.get(), d.f_partiality.get(), d.g.get(), d.scaled.get(), d.f_corr.get()); } diff --git a/image_analysis/spot_finding/AdaptiveSpotFinderGPU.cu b/image_analysis/spot_finding/AdaptiveSpotFinderGPU.cu index 2cafc5c06..e817708a3 100644 --- a/image_analysis/spot_finding/AdaptiveSpotFinderGPU.cu +++ b/image_analysis/spot_finding/AdaptiveSpotFinderGPU.cu @@ -367,11 +367,13 @@ void AdaptiveSpotFinderGPU::ReducePass(const ImagePreprocessorBuffer &image, flo gpu_pixel_to_bin->get(), gpu_corrections->get(), image.getGPUBuffer(), gpu_mean, gpu_sigma, clip_k, accumulate_corrected, gpu_sum, gpu_sum2, gpu_count, gpu_sum_corr, gpu_sum2_corr, npix, nbins); + cuda_err(cudaGetLastError()); } else { reduce_rings_global<<>>( gpu_pixel_to_bin->get(), gpu_corrections->get(), image.getGPUBuffer(), gpu_mean, gpu_sigma, clip_k, accumulate_corrected, gpu_sum, gpu_sum2, gpu_count, gpu_sum_corr, gpu_sum2_corr, npix, nbins); + cuda_err(cudaGetLastError()); } } @@ -379,6 +381,7 @@ void AdaptiveSpotFinderGPU::FinalizeStats() { const int threads = 128; const int blocks = (nbins + threads - 1) / threads; finalize_rings<<>>(gpu_sum, gpu_sum2, gpu_count, gpu_mean, gpu_sigma, nbins); + cuda_err(cudaGetLastError()); } // Host reproduction of AdaptiveSpotFinderCPU Stage B, from the clipped raw per-ring stats. @@ -478,6 +481,7 @@ void AdaptiveSpotFinderGPU::Detect(const ImagePreprocessorBuffer &image, cuda_err(cudaMemsetAsync(gpu_ring, 0, OutputByteSize(), *stream)); flag_strong<<>>( image.getGPUBuffer(), gpu_pixel_to_bin->get(), gpu_thr, gpu_ring, npix, nbins); + cuda_err(cudaGetLastError()); if (settings.signal_to_noise_threshold <= 0.0f) { // No local test asked for: the ring threshold alone decides, as the fixed photon floor @@ -492,6 +496,7 @@ void AdaptiveSpotFinderGPU::Detect(const ImagePreprocessorBuffer &image, local.photon_count_threshold = 0; ImageSpotFinderGPU::Detect(image, local); and_bits<<>>(gpu_out_1, gpu_ring, OutputSize()); + cuda_err(cudaGetLastError()); // The bit buffer stays on the device and ExtractComponents reads it there, on this same stream, // so the ordering already guarantees and_bits has finished. Waiting here only idled the host. } diff --git a/image_analysis/spot_finding/ImageSpotFinderGPU.cu b/image_analysis/spot_finding/ImageSpotFinderGPU.cu index 4270d1655..d2ec27399 100644 --- a/image_analysis/spot_finding/ImageSpotFinderGPU.cu +++ b/image_analysis/spot_finding/ImageSpotFinderGPU.cu @@ -300,8 +300,10 @@ void ImageSpotFinderGPU::Detect(const ImagePreprocessorBuffer &image, const Spot cuda_err(cudaMemsetAsync(gpu_out_1, 0, OutputSize() * sizeof(uint32_t), *stream)); analyze_pixel<<>> (image.getGPUBuffer(), gpu_out_1, gpu_out_0, spot_params); + cuda_err(cudaGetLastError()); analyze_pixel<<>> (image.getGPUBuffer(), gpu_out_0, gpu_out_1, spot_params); + cuda_err(cudaGetLastError()); // The bit buffer stays on the device - ExtractComponents reads it there. cuda_err(cudaStreamSynchronize(*stream)); } diff --git a/image_analysis/spot_finding/SpotExtractorGPU.cu b/image_analysis/spot_finding/SpotExtractorGPU.cu index d286c5586..d823ec084 100644 --- a/image_analysis/spot_finding/SpotExtractorGPU.cu +++ b/image_analysis/spot_finding/SpotExtractorGPU.cu @@ -338,20 +338,27 @@ void SpotExtractorGPU::Extract(const uint32_t *gpu_strong, const int32_t *gpu_im const int max_pix = static_cast(settings.max_pix_per_spot); count_bits<<>>(gpu_strong, gpu_res_mask, gpu_block_count, nwords); + cuda_err(cudaGetLastError()); scan_block_counts<<<1, FINISH_THREADS, 0, *stream>>>(gpu_block_count, gpu_block_offset, gpu_nstrong, compact_blocks); + cuda_err(cudaGetLastError()); scatter_bits<<>>(gpu_strong, gpu_res_mask, gpu_block_offset, gpu_image, gpu_index, gpu_value, nwords, max_strong); + cuda_err(cudaGetLastError()); // Fixed grids reading the strong-pixel count from device memory: the host never learns it, so it // never has to synchronise in the middle of the frame. init_parent<<<512, THREADS, 0, *stream>>>(gpu_parent, gpu_nstrong, max_strong); + cuda_err(cudaGetLastError()); union_neighbours<<<512, THREADS, 0, *stream>>>(gpu_index, gpu_parent, gpu_nstrong, max_strong, width); + cuda_err(cudaGetLastError()); resolve_roots<<<512, THREADS, 0, *stream>>>(gpu_parent, gpu_root, gpu_nstrong, max_strong); + cuda_err(cudaGetLastError()); finish_components<<<1, FINISH_THREADS, 0, *stream>>>(gpu_index, gpu_value, gpu_root, gpu_label, gpu_count, gpu_spot, gpu_spot_out, gpu_nspot, gpu_nstrong, max_strong, width, max_pix, static_cast(SPOT_SHAPE_FREE_PIXELS), static_cast(SPOT_MIN_FILL_PERCENT)); + cuda_err(cudaGetLastError()); // Rides along with the spot count on the frame's one synchronisation, so knowing how many strong // pixels there were costs nothing. cuda_err(cudaMemcpyAsync(host_nstrong, gpu_nstrong, sizeof(uint32_t), cudaMemcpyDeviceToHost, *stream)); diff --git a/reader/JFJochHDF5Reader.h b/reader/JFJochHDF5Reader.h index d281c0b10..175fedd89 100644 --- a/reader/JFJochHDF5Reader.h +++ b/reader/JFJochHDF5Reader.h @@ -20,6 +20,11 @@ // same images: the original _master.h5 plus optional reprocessing _process.h5 results. The active // snapshot supplies the dataset and per-image metadata; the image source always supplies pixels, // so switching snapshots never reloads image data. +// +// SWMR / growing files are NOT supported: a file opened here is expected to be final. Dataset +// extents, image counts and per-image metadata are read once at open and never re-polled, so a +// file still being written will be seen at whatever size it had at open - to pick up new frames, +// close and reopen it. class JFJochHDF5Reader : public JFJochReader { HDF5ImageSource image_source_; std::map > snapshots_; diff --git a/viewer/JFJochImageReadingWorker.cpp b/viewer/JFJochImageReadingWorker.cpp index 889098de6..3367497dc 100644 --- a/viewer/JFJochImageReadingWorker.cpp +++ b/viewer/JFJochImageReadingWorker.cpp @@ -206,6 +206,10 @@ void JFJochImageReadingWorker::LoadFile_i(const QString &filename, qint64 image_ // image. Without this, a caller that re-issues LoadFile per image instead of LoadImage (e.g. // a D-Bus client stepping through a dataset) pays a full close+reopen+reparse of the HDF5 // dataset for every single image, instead of the cheap single-image read LoadImage does. + // This also means re-opening the SAME path never re-reads it. That is deliberate: SWMR / + // growing files are not supported anywhere in the reader stack - a file we open is expected + // to be final (see JFJochHDF5Reader) - so there is no legitimate reason to re-read a path + // that is already open. A file still being written is picked up by closing and reopening. if (!http_mode && !filename.isEmpty() && filename == current_file) { LoadImage_i(image_number, summation); return; diff --git a/writer/FileWriter.cpp b/writer/FileWriter.cpp index 6a72024a9..a9fcee718 100644 --- a/writer/FileWriter.cpp +++ b/writer/FileWriter.cpp @@ -61,6 +61,14 @@ void FileWriter::WriteHDF5(const DataMessage& msg) { if (msg.number < 0) throw JFJochException(JFJochExceptionCategory::ArrayOutOfBounds, "No support for negative images"); + // The image number arrives over the network and sizes vectors downstream (file table here, + // per-image arrays in the data-file plugins as number * max_spots) - unbounded it is a huge + // allocation at best and, past the point where the product wraps, an out-of-bounds write. + if (start_message.number_of_images > 0 + && static_cast(msg.number) >= start_message.number_of_images) + throw JFJochException(JFJochExceptionCategory::ArrayOutOfBounds, + "Image number outside the collection's declared range"); + if (format == FileWriterFormat::NXmxIntegrated && master_file) { if (files.empty() ) files.resize(1);