Upgrade to parallelproj 2.0 and use cuvec

.appveyor.yml

@@ -51,7 +51,7 @@ build_script:
   - "set PATH=%MINICONDA%;%MINICONDA%\\Scripts;%MINICONDA%\\Library\\bin;%PATH%"
   # install parallelproj and Python stuff
   # don't do numpy here due to https://github.com/conda-forge/numpy-feedstock/issues/350
-  - conda create --name stirbuild -c conda-forge  -yq libparallelproj=1 swig pytest ccache ninja cmake
+  - conda create --name stirbuild -c conda-forge  -yq libparallelproj=2 swig pytest ccache ninja cmake
   - CALL conda.bat activate stirbuild
   - python --version
   - pip install numpy matplotlib

.github/workflows/build-test.yml

@@ -140,7 +140,9 @@ jobs:
           compiler: clang
           compiler_version: 21
           cuda_version: "0"
+          # openMP off as problems again with brew
           BUILD_FLAGS: "-DSTIR_OPENMP=OFF"
+          # parallelproj off as it needs openMP
           parallelproj: "OFF"
           BUILD_TYPE: "Release"
           ROOT: "OFF"
@@ -303,16 +305,20 @@ jobs:
           esac
 
           if test "${{matrix.parallelproj}}XX" == "ONXX"; then
-             git clone --depth 1 --branch v1.7.3 https://github.com/gschramm/parallelproj
-             mkdir parallelproj/build
-             cd parallelproj/build
+             git clone --depth 20 https://github.com/KUL-recon-lab/libparallelproj/
+             cd libparallelproj
+             git checkout v2.0.3
+             mkdir build
+             cd build
              if test "${{matrix.cuda_version}}" == "0"; then
-               extra_args="-DSKIP_CUDA_LIB:BOOL=ON"
+                 extra_args="-DUSE_CUDA:BOOL=OFF"
+             else
+                 extra_args="-DUSE_CUDA:BOOL=ON"
              fi
              cmake .. -DCMAKE_INSTALL_PREFIX=${CMAKE_INSTALL_PREFIX} ${extra_args}
              cmake --build . --target install --config Release
              cd ../..
-             rm -rf parallelproj
+             rm -rf libparallelproj
           fi
 
           # Install ROOT (warning: brittle due to OS versions etc)

.gitmodules

@@ -1,3 +1,6 @@
 [submodule "external_helpers/fmt"]
 	path = external_helpers/fmt
 	url = https://github.com/fmtlib/fmt.git
+[submodule "cuvec"]
+	path = external_helpers/CuVec
+	url = https://github.com/AMYPAD/CuVec

CMakeLists.txt

@@ -261,17 +261,18 @@ endif()
 
 # Parallelproj
 if(NOT DISABLE_Parallelproj_PROJECTOR)
-    find_package(parallelproj 1.3.4 CONFIG)
+    find_package(parallelproj CONFIG)
     if (parallelproj_FOUND)
       set(STIR_WITH_Parallelproj_PROJECTOR ON)
-      if (parallelproj_built_with_CUDA)
+      if (PARALLELPROJ_VERSION VERSION_LESS 2.0)
+        message(WARNING "Found parallelproj ${parallelproj_VERSION}. This will not be supported anymore in the near future.")        
+      endif()
+      if (PARALLELPROJ_CUDA)
         message(STATUS "Found parallelproj ${parallelproj_VERSION} (will use its CUDA support)")
+        set(parallelproj_built_with_CUDA ON)
       else()
         message(STATUS "Found parallelproj ${parallelproj_VERSION} (but using its OpenMP version as it wasn't built with CUDA)")
       endif()
-      if (parallelproj_VERSION VERSION_LESS 1.0.1)
-        message(STATUS "If the above parallelproj info looks incorrect, upgrade it to at least 1.0.1 (but 1.2.13 or later is recommended)")
-      endif()
     endif()
 endif()
 if (STIR_WITH_Parallelproj_PROJECTOR)

documentation/release_6.4.htm

@@ -12,15 +12,14 @@ <h1>Summary of changes in STIR release 6.4</h1>
   <h2>Overall summary</h2>
 
     <p>
-      This version is 100% backwards compatible with STIR 6.3, aside from three important bug fixes,
-      which will change results (see below).
+      This version is 100% backwards compatible with STIR 6.3.
     </p>
     <p>
       This is a release with many new features, as listed below. Highlights are
       <ul>
 
       </ul>
-      But there are many other changes from other contibutors as well, see below.
+      But there are many other changes from other contributors as well, see below.
       Of course, there is also the usual code-cleanup and some improvements to the documentation.
       Overall overview and release management was by Kris Thielemans (UCL) with help
       from Daniel Deidda (NPL).
@@ -29,8 +28,8 @@ <h2>Overall summary</h2>
     <h2>Patch release info</h2>
     <ul>
       <li>
-        6.4.0 released x/2025<br>
-        <a href="https://github.com/UCL/STIR/milestone/12">GitHub Milestone 6.3</a>
+        6.4.0 released XXX/2026<br>
+        <a href="https://github.com/UCL/STIR/milestone/13">GitHub Milestone 6.4</a>
       </li>
     </ul>
 
@@ -61,7 +60,25 @@ <h3>Changed functionality</h3>
       <br>
       <a href=https://github.com/UCL/STIR/pull/1667>PR #1667</a>.
     </li>
+    <li>
+      Several optimisations were contributed as part of the <a href="https://www.ccpsynerbi.ac.uk/events/airbi-hackathon/">SyneRBI AI-RBI hackathon</a>:
+      <ul>
+        <li>
+          Compatibility with <a href="https://github.com/KUL-recon-lab/libparallelproj/">libparallelproj 2.0</a> and usage of CUDA managed pointers via
+          <a href="https://github.com/AMYPAD/cuvec">CuVec</a> for internal variables in our parallelproj interface and <code>CudaGibbsPenalty</code>.
+          This results in a ~20% speed-up, but also code simplification.
+          <br>
+          <a href=https://github.com/UCL/STIR/pull/1689>PR #1689</a>.
+        </li>
+        <li>
+          Extra constructors for array, image and projdata classes that allow <code>std::move</code> for input arrays.
+          <br>
+          <a href=https://github.com/UCL/STIR/pull/1693>PR #1693</a> and <a href=https://github.com/UCL/STIR/pull/1694>PR #1694</a>.
+        </li>
+      </ul>
+    </li>
   </ul>
+
   <h4>Changes to examples</h4>
   <ul>
   </ul>

external_helpers/CMakeLists.txt

@@ -39,3 +39,11 @@ else()
   set(FMT_INCLUDE_DIRS ${FMT_INCLUDE_DIRS} PARENT_SCOPE)
 endif()
 
+################ CuVec
+if(NOT EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/CuVec/Makefile")
+  message(FATAL_ERROR "The CuVec submodule was not downloaded! GIT_SUBMODULE was turned off or failed. Please update submodules and try again.")
+endif()
+set(CUVEC_INCLUDE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/CuVec/cuvec/include")
+set(CUVEC_INCLUDE_DIR ${CUVEC_INCLUDE_DIR} PARENT_SCOPE)
+
+install(FILES "${CUVEC_INCLUDE_DIR}/cuvec.cuh" DESTINATION "${STIR_INCLUDE_INSTALL_DIR}/")

src/CMakeLists.txt

@@ -269,6 +269,14 @@ target_link_libraries(stir_registries)
 # TODO, really should use stir_libs.cmake
 target_include_directories(stir_registries PRIVATE ${STIR_INCLUDE_DIR})
 target_include_directories(stir_registries PRIVATE ${Boost_INCLUDE_DIR})
+target_include_directories(stir_registries PUBLIC
+  $<BUILD_INTERFACE:${CUVEC_INCLUDE_DIR}>
+  $<INSTALL_INTERFACE:include>)
+if (NOT STIR_WITH_CUDA)
+  target_compile_definitions(stir_registries PUBLIC CUVEC_DISABLE_CUDA)
+else()
+  target_link_libraries(stir_registries PUBLIC CUDA::cudart)
+endif()
 
 # go and look for CMakeLists.txt files in all those directories
 foreach(STIR_DIR ${STIR_DIRS} ${STIR_TEST_DIRS})

src/include/stir/cuda_utilities.h

@@ -1,5 +1,5 @@
 /*
-    Copyright (C) 2024, University College London
+    Copyright (C) 2024, 2026, University College London
     Copyright (C) 2025, University of Milano-Bicocca
     This file is part of STIR.
 
@@ -22,6 +22,10 @@
 #include "stir/Array.h"
 #include "stir/info.h"
 #include "stir/error.h"
+#ifdef __CUDACC__
+#  include <cuda_runtime.h>
+#  include "cuvec.cuh"
+#endif
 #include <vector>
 
 START_NAMESPACE_STIR
@@ -45,12 +49,16 @@ struct cuda_int3
   int x = 0, y = 0, z = 0;
 };
 #else
-#  include <cuda_runtime.h>
 typedef dim3 cuda_dim3;
 typedef int3 cuda_int3;
 #endif
 
 #ifdef __CUDACC__
+
+//! copy an `Array` to pre-allocated device memory
+/*!
+  \ingroup CUDA
+*/
 template <int num_dimensions, typename elemT>
 inline void
 array_to_device(elemT* dev_data, const Array<num_dimensions, elemT>& stir_array)
@@ -71,9 +79,25 @@ array_to_device(elemT* dev_data, const Array<num_dimensions, elemT>& stir_array)
     }
 }
 
+//! copy an `Array` to pre-allocated CuVec
+/*!
+  \ingroup CUDA
+*/
 template <int num_dimensions, typename elemT>
 inline void
-array_to_host(Array<num_dimensions, elemT>& stir_array, const elemT* dev_data)
+array_to_device(CuVec<elemT>& dev_data, const Array<num_dimensions, elemT>& stir_array)
+{
+  std::copy(stir_array.begin_all(), stir_array.end_all(), dev_data.begin());
+}
+
+//! copy CUDA pointer to `Array`
+/*!
+  \ingroup CUDA
+  The third argument is ignored, as `cudaMemcpy` always syncs device and host.
+*/
+template <int num_dimensions, typename elemT>
+inline void
+array_to_host(Array<num_dimensions, elemT>& stir_array, const elemT* dev_data, bool /* sync */ = true)
 {
   if (stir_array.is_contiguous())
     {
@@ -92,6 +116,20 @@ array_to_host(Array<num_dimensions, elemT>& stir_array, const elemT* dev_data)
     }
 }
 
+//! copy CuVec to `Array`
+/*!
+  \ingroup CUDA
+  If \a sync = \c true, the function will call `cudaDeviceSynchronize()` before copying.
+*/
+template <int num_dimensions, typename elemT>
+inline void
+array_to_host(Array<num_dimensions, elemT>& stir_array, const CuVec<elemT>& dev_data, bool sync = true)
+{
+  if (sync)
+    cudaDeviceSynchronize();
+  std::copy(dev_data.begin(), dev_data.end(), stir_array.begin_all());
+}
+
 //! \brief Performs a parallel reduction sum on shared memory within a CUDA thread block, final value stored in shared_mem[0].
 template <typename elemT>
 __device__ inline void