mx/Jungfraujoch
2
1
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases 68 Wiki Activity

Build the GPU engines a worker never uses on first use, not always

Browse Source 
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>
This commit is contained in:
jungfrau
2026-08-15 21:42:41 -04:00
co-authored by Claude Opus 5
parent eb634400ea
commit 2cbb3fc8b4
6 changed files with 84 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files
image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu
+13 -2
View File
@@ -231,6 +231,18 @@ BSLZ4DecoderGPU::BSLZ4DecoderGPU(size_t in_max_uncompressed_bytes, std::shared_p
// tens; sizing this from the UNCOMPRESSED size cost ~73 MB per worker to hold ~4 MB.
}
// The stored pixel depth is fixed within a dataset, so in practice this runs once - but it is not
// promised anywhere, and sizing for the widest type instead would hold twice the memory a 16-bit
// detector needs. Grown with slack because cudaMalloc and cudaFree synchronise the whole device.
void BSLZ4DecoderGPU::EnsureUncompressedCapacity(size_t bytes) {
if (bytes <= max_uncompressed_bytes)
return;
const size_t want = std::max(bytes, max_uncompressed_bytes + max_uncompressed_bytes / 2);
cuda_err(cudaStreamSynchronize(*stream));
gpu_shuffled = CudaDevicePtr<uint8_t>(want);
max_uncompressed_bytes = want;
}
void BSLZ4DecoderGPU::EnsureCompressedCapacity(size_t bytes) {
if (bytes <= compressed_capacity)
return;
@@ -259,8 +271,7 @@ BSLZ4ShuffledImage BSLZ4DecoderGPU::DecodeShuffled(const CompressedImage &image)
if (clen < 12)
throw JFJochException(JFJochExceptionCategory::Compression, "bslz4 chunk shorter than its header");
if (total_bytes > max_uncompressed_bytes)
throw JFJochException(JFJochExceptionCategory::Compression, "bslz4 image larger than the decoder was sized for");
EnsureUncompressedCapacity(total_bytes);
if (be64(src) != total_bytes)
throw JFJochException(JFJochExceptionCategory::Compression, "bslz4 header size does not match the image");