v1.0.0-rc.178 (#88)

* Rugnux: Processes serial crystallography data from XFELs - SACLA (MPCCD), PAL-XFEL and European XFEL (CXI), LCLS (XTC) and SwissFEL Alvra (JUNGFRAU with a CrystFEL module geometry) - with several files or Jungfraujoch master files merged in one run (`rugnux *_master.h5`).
* Rugnux: Stills processing takes a known unit cell and space group as a starting point, finds the beam centre (and, with a cell, the detector distance) from the images, detects the saturation level in other facilities' files and refines the placement of individual detector panels.
* Rugnux: The stills report gives R_split and CC1/2 from frame half-sets, comparable with CrystFEL.
* jfjoch_viewer: Opens the same XFEL file formats.
* Validation: A stills battery checks Rugnux against deposited structures from SACLA, PAL-XFEL, LCLS, European XFEL and SwissFEL.
* Rugnux: More reliable convergence of rotation scaling, and live molecular replacement starts as soon as the model is uploaded.
* jfjoch_broker: Initialize clears the previous dataset's bearer tokens, image buffer and live Rugnux status, and a slow CUDA JIT cache is reported at start.

Reviewed-on: #88
Co-authored-by: Filip Leonarski <filip.leonarski@psi.ch>

tests/CUDAMemHelpersTest.cpp

@@ -138,11 +138,11 @@ TEST_CASE("CudaPinnedPool_ReusedByTheNextEngine", "[CUDAMemHelpers]") {
     CHECK(cuda_release_pinned_pool() == 0);
 }
 
-// Within a scope a cuFFT plan is kept once made and lent to one user at a time: two users at once get
-// two plans, the next one after them gets one of those back, and a transform on a lent plan - another
-// stream, a work area of the caller's - gives the same result bit for bit. When the scope ends, every
-// kept plan is destroyed; outside a scope nothing is kept.
-TEST_CASE("CudaFFTPlan_KeptAndLentToOneUserAtATime", "[CUDAMemHelpers]") {
+// Within a scope a cuFFT plan is made once per shape and shared: two users at once get the same plan,
+// each transforming on its own stream with its own work area, bit for bit as on a plan of its own; the
+// plan is kept when they are done, and the next user gets it. When the scope ends, every kept plan is
+// destroyed; outside a scope a plan goes with its last user.
+TEST_CASE("CudaFFTPlan_SharedAndKept", "[CUDAMemHelpers]") {
     if (get_gpu_count() == 0)
         SKIP("No CUDA GPU present");
     REQUIRE(jfjoch_cufft_plans::kept() == 0);
@@ -154,11 +154,11 @@ TEST_CASE("CudaFFTPlan_KeptAndLentToOneUserAtATime", "[CUDAMemHelpers]") {
     for (size_t i = 0; i < input.size(); i++)
         input[i] = static_cast<float>((i * 7919) % 1013) - 500.0f;
     CudaDevicePtr<float> d_in(input.size());
-    CudaDevicePtr<cufftComplex> d_out(nc * batch);
     REQUIRE(cudaMemcpy(d_in, input.data(), input.size() * sizeof(float), cudaMemcpyHostToDevice) == cudaSuccess);
 
     auto transform = [&](const CudaFFTPlan &plan, cudaStream_t stream) {
-        REQUIRE(cufftExecR2C(plan, d_in, d_out) == CUFFT_SUCCESS);
+        CudaDevicePtr<cufftComplex> d_out(nc * batch);
+        REQUIRE(plan.R2C(d_in, d_out) == CUFFT_SUCCESS);
         std::vector<cufftComplex> out(nc * batch);
         REQUIRE(cudaMemcpyAsync(out.data(), d_out, out.size() * sizeof(cufftComplex), cudaMemcpyDeviceToHost,
                                 stream) == cudaSuccess);
@@ -166,30 +166,34 @@ TEST_CASE("CudaFFTPlan_KeptAndLentToOneUserAtATime", "[CUDAMemHelpers]") {
         return out;
     };
 
-    CudaStream stream_a, stream_c;
+    CudaStream stream_a, stream_b, stream_c;
     std::vector<cufftComplex> first;
-    cufftHandle handle_a = 0;
+    cufftHandle handle = 0;
     {
         CudaFFTPlan a(shape, stream_a);
-        CudaFFTPlan b(shape);
-        handle_a = a.get();
-        CHECK(handle_a != b.get());
+        CudaFFTPlan b(shape, stream_b);
+        handle = a.Handle();
+        CHECK(b.Handle() == handle);
+        CHECK(jfjoch_cufft_plans::kept() == 1);
         first = transform(a, stream_a);
+        const std::vector<cufftComplex> second = transform(b, stream_b);
+        CHECK(memcmp(second.data(), first.data(), first.size() * sizeof(cufftComplex)) == 0);
     }
+    CHECK(jfjoch_cufft_plans::kept() == 1);
 
     const size_t work = CudaFFTPlan::WorkBytes(shape);
     CudaDevicePtr<char> work_area(std::max<size_t>(work, 1));
     CudaFFTPlan c(shape, stream_c, work_area);
-    CHECK(c.get() == handle_a);
+    CHECK(c.Handle() == handle);
     CHECK(c.WorkBytes() == work);
     const std::vector<cufftComplex> again = transform(c, stream_c);
     REQUIRE(again.size() == first.size());
     CHECK(memcmp(again.data(), first.data(), first.size() * sizeof(cufftComplex)) == 0);
-    CHECK(jfjoch_cufft_plans::kept() == 1);   // b; a is lent to c
 
-    c = CudaFFTPlan();
-    CHECK(jfjoch_cufft_plans::kept() == 2);
+    // The scope ends while c still uses the plan: it goes with c.
     scope.reset();
+    CHECK(jfjoch_cufft_plans::kept() == 1);
+    c = CudaFFTPlan();
     CHECK(jfjoch_cufft_plans::kept() == 0);
 
     CHECK(CudaFFTPlan::WorkBytes(shape) == work);