From 196c72a7fedc3aa78a8b28a71272a6c09d0b779d Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Fri, 31 Jul 2026 11:50:10 +0200 Subject: [PATCH] Bragg prediction: the rotation GPU launch was one plane short in each direction The kernel guards against 2*max_hkl+1 and maps thread i to h = i - max_hkl, but the host launched a grid sized 2*max_hkl. The h = k = l = +max_hkl planes were therefore never launched while -max_hkl was, so the GPU predicted an asymmetric subset of what the CPU loop (inclusive on both ends) does. The same bug was fixed on the stills twin when the whole hkl range moved to the GPU; the rotation predictor kept the old expression. It only bites where the cell actually reaches |h| = 100 inside d_min - a ~150 A axis at 1.5 A - so most data never noticed. Over the 33-crystal rotation battery 29 crystals are bit-identical and 4 gain observations, all of them large-cell or high-resolution: +8519, +4693, +901 and +758 observations, with the high-shell CC1/2 up 15.0->15.1%, 52.0->52.2%, 76.3->76.6% and 51.6->52.1%. Nothing is lost anywhere, and R-meas and ISa move by at most 0.01. Co-Authored-By: Claude Opus 5 (1M context) --- image_analysis/bragg_prediction/BraggPredictionRotGPU.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu b/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu index 46968ddf..5a0ed52e 100644 --- a/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu +++ b/image_analysis/bragg_prediction/BraggPredictionRotGPU.cu @@ -286,7 +286,8 @@ int BraggPredictionRotGPU::Calc(const DiffractionExperiment &experiment, cudaMemcpyAsync(dK, &hK, sizeof(KernelConstsRot), cudaMemcpyHostToDevice, stream); cudaMemsetAsync(d_count, 0, sizeof(int), stream); - const int range = 2 * settings.max_hkl; + // Inclusive on both ends, matching the kernel's own range and the CPU loop (-max_hkl .. +max_hkl). + const int range = 2 * settings.max_hkl + 1; dim3 block(8, 8, 8); dim3 grid((range + block.x - 1) / block.x, (range + block.y - 1) / block.y,