7 Commits

Include renames

Author	SHA1	Message	Date
jungfrauandClaude Opus 5	f36cd88795	Give a 16-bit image a saturation code when its limit is 65534 ... The narrow encoding picked NARROW_BAD as the saturation code whenever the saturation limit was above UINT16_MAX - 2. That is one too strict. A real value is STRICTLY below the limit, so a limit of UINT16_MAX - 1 still leaves UINT16_MAX - 1 free to be the code; only a limit of the whole range has nothing to spare, and that is exactly the case where no pixel can be saturated, because the "is error" test claims UINT16_MAX first. At a limit of exactly 65534 the old condition therefore stored a saturated pixel as the masked code, and it widened back to INT32_MIN instead of INT32_MAX. Masked and saturated are not interchangeable: the strong-pixel search flags a saturated pixel unconditionally and a masked one never, so the overloaded core of the strongest spots dropped out of the strong-pixel mask. Bragg integration treats the two alike and the image statistics are taken from the raw value, so nothing downstream of those moved - which is why a battery over 24 crystals showed nothing. That value is not a corner case. GetByteDepthImage()-driven writing stores saturation_value = GetSaturationLimit() - 1 and the readers take it back as-is, so a 16-bit acquisition whose detector cutoff is at or above the full range comes back with a limit of exactly 65534. The one 16-bit dataset in the rotation test set declares 11963, which is safe, so it could not have caught this. Found by review, not by testing, because nothing tested the narrow path at all: every GPU test writes into the wide buffer directly and never asks for the narrow one, and every preprocessor test copies the image back to the host, which forces the wide path. So the test comes with the fix. It runs the wide and the narrow preprocessor over the same synthetic frame - values around each boundary, masked on a stride coprime with the value cycle so every value appears both masked and unmasked - across saturation limits of none, 5000, 0xFFFD, 0xFFFE and 0xFFFF, on both the host-upload and the device-decode entry point, and compares the statistics and every pixel. Against the old condition it fails with 625860 differing pixels; against this one it passes. Also makes the header self-contained: it uses __host__/__device__ and the CUDA vector types and only compiled because every includer happened to pull in cuda_runtime.h first. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>	2026-08-18 18:41:53 -04:00
jungfrauandClaude Opus 5	9f49e5abb7	Keep 16-bit images 16-bit through the GPU pipeline ... A detector reading out 16 bits had its frame widened to int32 the moment it was decoded, and every per-pixel pass over that frame then moved four bytes a pixel to carry two. Those passes - the ring statistics three times over, the strong-pixel search, spot extraction, the azimuthal and ROI integrators, Bragg integration - are the bulk of the image loop's device traffic, and 16 bits is the mode a fast acquisition runs in, which is exactly where throughput matters. The preprocessed image now keeps the width of its source. Two codes at the top of the 16-bit range carry the two special states, and they cannot collide with a real value: 0xFFFF masked, or the source's own bad-pixel marker. saturation a pixel at or above the saturation limit. 0xFFFE where the limit code leaves room - a 16-bit EIGER declares a count-rate limit of a few thousand, so there is room to spare - and 0xFFFF where the limit is the whole range, in which case the "is error" test has already claimed 0xFFFF, nothing can be saturated, and 0xFFFE stays a real value. Either way a real value is strictly below the saturation limit and so below both codes. Nothing is clipped and nothing is lost, and which code is in force is carried with the image rather than assumed. No engine learns a second convention. PixelView widens on load, so a masked pixel still reads as INT32_MIN and a saturated one as INT32_MAX, and every existing `v != INT32_MIN && v != INT32_MAX` test keeps its meaning. One code path, not two instantiations that can drift apart; the branch is on a pointer that is the same for every thread of every block, on kernels whose time is the loads it selects between. The vector loads are kept - four pixels still arrive in one transaction, 16 bytes wide or 8, whichever the image is. The wide path is unchanged, and is still taken for anything that is not a 16-bit source, and for any caller that wants the preprocessed image copied back to the host - that mirror is int32 and the CPU engines know only that convention. Measured on the one 16-bit dataset in the rotation test set, which is also the smallest detector in it (2.5M pixels, where per-pixel work is a small part of the loop): image loop 1.025 s -> 1.005 s at one GPU, whole run 5.64 s -> 5.52 s. The gain scales with the frame, so a 16M-pixel detector - where six full-frame passes are 86 % of the loop's GPU time - has much more to gain, and nothing here can measure that: every other dataset in the test set is stored 32-bit. Correctness on that dataset is exact where it can be: indexing rate, first-pass validation score and the integrated partial count are identical to the wide path, and its whole battery row - reflections, observations, space group, R_meas, CC1/2, ISa, mosaicity - is unchanged. Battery 6m17s -> 6m16s, 21/24 space groups, no failures. Also logs, once per run, the width the images are stored in, since it decides how much of the frame moves through every pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>	2026-08-18 01:52:27 -04:00
jungfrauandClaude Opus 5	2cbb3fc8b4	Build the GPU engines a worker never uses on first use, not always ... Every worker thread built a full set of analysis engines. Two of them are never asked for on the offline path: the fixed-threshold spot finder, because detection is adaptive by default, and the azimuthal integrator, because the fused adaptive finder produces the profile as a by-product. They are still needed elsewhere - the broker defaults to non-adaptive detection, and --no-adaptive-spots asks for the finder - so they are built on first use rather than removed. A lazily built finder takes the current resolution mask on construction; without that it would find spots outside the limits it was never told about. The bitshuffle decoder sized its output buffer for the widest pixel type there is rather than the one the images actually have, holding a second full frame per worker on 16-bit data. It is sized from the image now and grows if a later frame needs more. The shared-table checksum runs over eight interleaved lanes. FNV's multiply is a loop-carried dependency, so one chain retires a byte every few cycles whatever memory bandwidth is spare, and every worker hashes tens of megabytes of geometry tables as it builds its engines - about 5% of all CPU samples on a 16M-pixel detector. Measured on a 16M-pixel rotation dataset: cudaMalloc 11314 -> 9474 calls and, with cudaFree, 117 s -> 78 s of aggregate thread time; both synchronise the whole device, so that time is spent blocking every other worker. Whole battery 15m32s -> 12m30s. Data quality against main, over 24 crystals and eight statistics each: the same space group on all 24, and every difference smaller than what two runs of an IDENTICAL binary produce (measured: 13 of 24 crystals reproduce exactly run to run, worst R_meas swing 5.5 points, against 4.6 points for main vs this branch). The float atomics in the reductions have always made this so. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-15 21:42:41 -04:00
jungfrauandClaude Opus 5	2f586d2267	Stop allocating GPU and pinned memory nothing reads ... Three resource fixes and two latent bugs, none of which changes a computed number. The preprocessed image has a host copy that only a CPU engine ever reads. On the GPU path every engine reads the device buffer instead, and rugnux always runs the fused adaptive finder, so that host copy is allocated, zeroed and PAGE-LOCKED for nothing - 72 MB per worker, 3.5 GB over 48 of them, and a cudaHostRegister each, which the driver serializes. It is now skipped by the same condition that already decides whether the device copies the image back. ImagePreprocessorBuffer keeps the pixel count separately so size() still answers when the mirror was not allocated. ROIIntegrationGPU asked device 0 for the SM count it sizes its grid from, while workers are pinned round-robin across the GPUs - so on a multi-GPU node it could size a grid from a card it never launches on. It asks the current device now, like every other engine. ~CudaRegisteredVector called a function that throws out of a destructor, and the move-assignment did the same from a noexcept function. Either would abort the process rather than report the failure, and teardown - after a device reset, or while another exception unwinds - is exactly where cudaHostUnregister fails. Both now use an unchecked unregister, as every other destructor in that header already does for its own teardown call. The throwing form stays for rebind()/unregister(), which are called from live code. Measured on a 16M-pixel rotation dataset: unchanged space group, merged reflection count and merging statistics. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>	2026-08-15 17:50:19 -04:00
leonarski_f	538f3504d3	v1.0.0.rc-161 (#71) ... This is an UNSTABLE release. It includes many experimental features, as well as many AI generated fixes. We recommend using rc.152 for production use. * **rugnux: significantly better quality of results, and faster.** A large rework of integration, scaling, merging, geometry refinement and space-group determination, together with measurements the program previously made no attempt at - the direct beam before indexing, the beam stop, the goniometer rotation scale, and the stretches of a sweep the crystal did not deliver. A rotation dataset typically gains observations at better <I/sigma> and R_meas, and every `mx` and `scale` run writes a `<prefix>_report.txt` results report modelled on XDS's `CORRECT.LP`. Many defaults moved with it: spot detection is self-calibrating, beam-stop detection and rotation geometry post-refinement are on, resolution limits default to as far as the detector reaches, and ice-ring handling engages only where the crystal is measured to have ice. * **jfjoch_viewer:** the beam-stop shadow, the detector calibration and the beam-centre measurement are reachable from "Analyze dataset"; the settings panel reports how the sample moved and how polarized the beam was; image rendering and interaction are faster. * **Performance:** bitshuffle+LZ4 images are decoded on the GPU rather than on the host, with the bitshuffle inverse fused into preprocessing so the decompressed frame is never held in device memory. * **Broker, writer, packaging and build:** image-slot lifetime and locking fixes, per-image datasets sized by the images actually written, the Debian/Ubuntu broker package renamed to `jfjoch`, and `image_analysis` compiling under MSVC again. **Breaking change to the rugnux command line:** * `--azint-only` and `--scale` are **removed**, replaced by `--mode azint` and `--mode scale`; the full pipeline is `--mode mx` and remains the default. A script passing the old flags now fails with the list of valid modes rather than silently running the wrong one. * `-t`/`--stride` is **refused on rotation data**: skipping frames cuts every reflection's rocking curve, so the combined fulls and their partiality would be measured over frames the sweep never recorded. Select a contiguous range with `-s`/`-e` instead. `--mode azint` and `--force-still` still take a stride. **Breaking changes to OpenAPI** - regenerate the client (`jfjoch-client` 1.0.0-rc.161, `frontend/src/client`) or read the affected fields as optional: * `image_scale_b` is removed from the `plot_type` enum, so a client requesting that plot now gets an error rather than a curve. * `azim_int_settings.high_q_recipA`, `spot_finding_settings.high_resolution_limit` and `spot_finding_settings.low_resolution_limit` are no longer `required`. All three mean "no limit at that end" when unset and are omitted from the response instead of carrying a placeholder value, which raises in a client generated from an rc.160-or-earlier spec. A value of 0 is still accepted and means the same thing. **Breaking changes to the stored formats** - a consumer reading these fields must treat them as optional: * The per-image image-scale B factor is no longer computed, so `/entry/MX/imageScaleBFactor` is absent from newly written HDF5 files and the corresponding key is absent from the CBOR DataMessage and END blocks. Files written by rc.160 and earlier still contain it and still open; nothing in the pipeline reads it any more. * `_reflns.jfjoch_diffrn_ISa` now carries the whole-range `1/sqrt(a*b)` that XDS's ISa denotes, and the error-model `a` and `b` are reported in XDS's convention; the strong-reflection asymptote moves to `_reflns.jfjoch_diffrn_ISa_asymptotic`. **A file written by an earlier version carries the asymptote under the plain `ISa` name.** Reviewed-on: #71 Co-authored-by: Filip Leonarski <filip.leonarski@psi.ch>	2026-08-13 17:03:10 +02:00
leonarski_f	54c0100e8e	v1.0.0-rc.157 (#67) ... This is an UNSTABLE release. It includes many experimental features, as well as many AI generated fixes. We recommend using rc.152 for production use. * rugnux: Rebrand the offline data-processing subsystem as `rugnux` and consolidate all offline analysis into the single `rugnux` binary - `jfjoch_process` is now `rugnux`, the former `jfjoch_azint` is now `rugnux --azint-only`, and `jfjoch_scale` is now `rugnux --scale` (see the new docs/NAMING.md and docs/RUGNUX.md). Scaling and merging are on by default for rotation and stills (`--no-merge` disables them), replacing the previous opt-in `-M, --scale-merge`. * rugnux: CLI fixes - default `-N` to all hardware threads, parse numeric option arguments strictly (reject non-numeric or trailing input instead of silently yielding 0), require `--wavelength > 0`, and correct the reproduced command line and `--scale` reference-cell handling. * rugnux: De-novo space-group improvements - recover genuine high symmetry and centred Bravais lattices from intensities, add an automatic CC1/2 high-resolution cutoff, and report L-test twinning statistics. * rugnux: Index weakly-diffracting low-resolution rotation data that previously failed (e.g. F-cubic crystals that diffract only to ~4 A on a detector reaching ~1.5 A). The per-frame indexing gate now measures the indexed fraction only within the resolution range the lattice actually diffracts to, so the many sub-diffraction ice/noise spots no longer make the fraction floor unreachable; the two-pass first pass tries several image-sampling schemes (spread across the whole rotation vs a consecutive wedge whose native stride keeps a reflection's rocking curve continuous, letting the FFT resolve a long axis) and keeps the one that indexes the most frames; and the de-novo space-group search no longer discards all reflections (and crashes) when every resolution shell falls below <I/sigma> = 1. * rugnux: Lower the low-resolution R-meas for strongly-diffracting rotation data - drop edge-of-sweep truncated fulls whose rocking curve was captured below `--min-captured-fraction` (default 0.7 for rotation), and report R-meas only over the observations kept by outlier rejection (matching XDS). The 0.7 default also strips the partiality-extrapolated fulls that dominate the intensity second moment on weakly-diffracting crystals, so the de-novo space-group search is no longer starved by the error-model I/sigma floor and recovers the correct symmetry (e.g. the F-cubic Benas crystals: Benas_3 -> F432, Benas_7 -> P6122, instead of P4/P1); on the reference battery every other crystal keeps its space group. * rugnux: Write the refined geometry (beam, tilt, axis) to _process.h5 and place non-standard mmCIF items under a reserved `jfjoch` prefix. * jfjoch_broker: Ordinary acquisition failures (receiver/writer/analysis problems, missed packets, writer disconnect) now return to the Idle state with an Error-severity message, so a run can be retried without an expensive re-initialisation; only failures that leave the detector in an undefined state (new JFJochCriticalException, e.g. PCIe/FPGA faults) go to the Error state and force re-initialisation. * jfjoch_broker: A synchronous /start now reports its failure to the HTTP caller instead of returning HTTP 200, and an incomplete or truncated dataset (missing packets, writer disconnect) is reported as an error rather than a "reduce frame rate" warning. * jfjoch_broker: Drop uncollected placeholder rows (number = -1) from the scan_result REST endpoint. * jfjoch_broker: Fix the inverted per-image compression ratio reported by the Lite receiver (was compressed/uncompressed instead of uncompressed/compressed). * jfjoch_broker: Bragg integration adds a quantization-noise variance floor with a box-sum fallback, and treats the type-maximum marker as an invalid pixel for unsigned image types. * jfjoch_writer: Detect file-overwrite conflicts at start for back-channel transports, and reset the writer when end-of-collection finalisation fails. * jfjoch_viewer: Preview overlays follow the geometry (resolution/ROI arcs, true beam centre, predictions, coral secondary-lattice spots, legend), add save-as-JPEG, and fix an HTTP live-follow memory leak. * Frontend: Improved aesthetics and usability, and added in-browser pixel-mask and JUNGFRAU-pedestal visualisation. * CI: Name the Windows installer jfjoch-viewer-* instead of jfjoch-*.Reviewed-on: #67 Co-authored-by: Filip Leonarski <filip.leonarski@psi.ch>	2026-07-11 07:19:11 +02:00
leonarski_f	c981e1b91c	v1.0.0-rc.137 (#46) ... This is an UNSTABLE release. The release has significant modifications and bug fixes, if things go wrong, it is better to revert to 1.0.0-rc.132. * jfjoch_broker: Better track time for each operation in the processing stack * jfjoch_broker: Rewrite preprocessing of diffraction images in the non-FPGA workflow to better use GPUs (work in progress) * jfjoch_broker: Remove ROI calculation in the non-FPGA workflow (work in progress) * jfjoch_viewer: Toolbar displays image number starting from 1 (instead of 0) Reviewed-on: #46	2026-04-25 19:59:21 +02:00