Data on the GPU: `PreferredDevice`, and what honours it¶

Every object in the library that holds bulk numbers now says where those numbers should live. One property, two words, the same on all of them:

Object	Property	Where it comes from
`Abstract::VoxImage` (so `VoxImage<T>`, `MultiVoxImage`)	`PreferredDevice`	this page
any `Collection<T>`, `MuonScatterCollection`, `EventTable`, `GeantEventCollection`	`PreferredDevice`	`CollectionControl`
any `Algorithm`	`PreferredDevice`	`uprops::Algorithm` (`IsAccelerated()`)

Set the data and the algorithm that reads it both to Device and nothing crosses the bus between them. That is the whole point; the rest of this page is what it took to be able to say it truthfully.

1. Device memory without a CUDA toolkit¶

DataAllocator<T> has had a VRAM half since the beginning, and every line of it sat inside #ifdef USE_CUDA — which is off in nearly every build, because switching it on needs nvcc and a CUDA toolkit. The effect was worse than the feature being absent:

void MoveToVRAM() {          // the old one, with CUDA off
  if (m_Device == MemoryDevice::VRAM) return;
  m_Device = MemoryDevice::VRAM;   // ... and nothing else
}
T *data() { return m_Device == MemoryDevice::RAM ? m_RamData : m_VramData; }

The move reported success, allocated nothing, and data() then handed back a null pointer. A promise that compiles and lies.

src/Core/DeviceMemory.h is the other way round. Allocating device memory and copying to and from it is the driver's job, and the driver is present wherever an NVIDIA GPU is usable — libcuda.so, the thing nvidia-smi talks to. Compiling kernels needs the toolkit, and nothing in DeviceMemory compiles a kernel. So the driver API is loaded with dlopen at run time and device memory works in an ordinary build:

DeviceMemory &device = DeviceMemory::Instance();
if (device.IsAvailable()) {
  void *block = device.Allocate(bytes);     // cuMemAlloc, through the driver
  device.CopyToDevice(block, host, bytes);  // cuMemcpyHtoD
  ...
  device.Free(block);
}

Three details that matter:

It retains the device's primary context (cuDevicePrimaryCtxRetain), not a context of its own — so a pointer from here is a pointer numba, CuPy and PyTorch can use inside the same process. A private context would hand Python addresses from a context it is not in.
The current context is per thread, so every operation sets it first. This library fills collections from worker threads.
GetHeldBytes() / GetAllocationCount() are the accounting a long session's leak shows up in. The driver will not say how big a pointer was, so the sizes are remembered here and given back on Free.

Everything answers honestly when there is no GPU. IsAvailable() is false, GetUnavailableReason() says why in words, Allocate returns null, and no caller is handed a pointer that is not one.

A build with -DUSE_CUDA=ON is still worth having — it gets kernels, and they reach the same memory. A device pointer is a device pointer, whoever allocated it.

2. What may go to the device: `DeviceStorable`¶

The first version of this gate was std::is_trivially_copyable<T>, which is wrong here, and the test suite said so: nothing in the library could go to the device. std::is_trivially_copyable also demands a trivial assignment, and an Eigen::Matrix declares its own operator= — so every vector, point, line and MuonScatter in this library failed it, while a memcpy of any of them is exactly right.

What actually has to hold is that a copy of the bytes is a copy of the object, and that letting it go needs no work:

template <typename T>
struct DeviceStorable
    : std::integral_constant<bool,
          std::is_trivially_copy_constructible<T>::value &&
          std::is_trivially_destructible<T>::value> {};

Specialise it for a type that is byte-copyable for a reason the compiler cannot see; never for one that owns memory. A Geant::GeantEvent holds a path and a name, fails it, and stays on the host however its collection's PreferredDevice is set — which is what GetCollectionDevice() reports.

The trait is worth keeping in mind when writing a type that is meant to go to the device. VoxRaytracer::RayData::Element is two numbers and carried an empty ~Element() {}, which is enough to make it non-trivially-destructible; the CUDA raytracer then asked for its output buffer in device memory, was honestly refused, and launched a kernel on the null pointer it got back. The destructor is gone (VoxRaytracer.h says why), and the launch paths in VoxRaytracerCUDA.hpp now check the pointer instead of trusting it.

3. Placement: when the data actually moves¶

CollectionControl::PlaceCollection() and Abstract::VoxImage::PlaceImage() are the same shape:

const bool onDevice = GetPreferredDevice() == MemoryDevice::VRAM &&
                      CanPlaceOnDevice() && IsDeviceMemoryAvailable();
const MemoryDevice target = onDevice ? MemoryDevice::VRAM : MemoryDevice::RAM;
if (ResidentDevice() != target) MoveImageTo(target);

and both run on every Object::Updated() — after a run filled the data, a document loaded it, the dimensions changed, or the preference itself moved. That is also what undoes a host read: At(), operator[], begin(), an export and a serialize() all bring the block home first (so host code is always correct), and the next announcement puts it back.

So Device pays off when what consumes the data works on the device too, and costs two copies a round when it does not. GetDataLocation() is the read-only property that says which of those is happening:

Device: Tesla P40
Host (no NVIDIA driver library: libcuda.so.1)
Host (these voxels cannot be copied to a device)

It is deliberately not serialized: it is a fact about the machine the process is running on, and a document is meant to open on any machine.

PreferredDevice is saved — but not through the property walk, and that distinction is the whole of the compatibility story. The inherited ContainerBox::serialize ends in SerializeValues(), which writes every serialized property in declaration order; the binary archive is positional, so a property added there would have shifted every field after it in every image ever saved. So the property is marked SetSerialized(false) and the field is written by the image's own serialize() behind the class version:

if (version >= 1) this->SerializeDeviceSettings(ar);

boost::serialization::version<uLib::VoxImage<T>> is 1 (and BOOST_CLASS_VERSION(uLib::MultiVoxImage, 1)), so a version 0 document has no field here to take and still loads — which VoxImageDeviceTest::TestAnOlderDocumentStillLoads proves by reading a shipped .gcx from before the change. (EventTable does the same with if (version >= 1).)

On loading, SerializeDeviceSettings only records the preference: the voxels are read into host memory right after it, and moving them to the device first would leave the block device-resident while the payload landed in the host buffer, so the next host read would fetch the stale device copy over it. The move happens once the voxels are in.

4. Reading it where it is: `IDeviceArray`¶

A device pointer is not useful on its own; a consumer needs to know how many scalars, of what type, how far apart. DeviceArrayView is that, and IDeviceArray is the one interface everything holding bulk data implements:

struct DeviceArrayView {
  void *data;              // the device address of the first scalar
  std::size_t count;       // how many
  std::size_t itemSize;    // bytes of one (4 for a float)
  std::size_t stride;      // bytes to the next -- itemSize when packed
  std::string typeString;  // "<f4", "<f8", "|u1"
  bool writable;
};

class IDeviceArray {
  virtual bool GetDeviceArray(DeviceArrayView &view) = 0;
  virtual bool IsDeviceResident() const = 0;
};

stride is what makes it work for this library's own types rather than only for arrays of floats:

Object	What is exposed	`typestr`	`stride`
`VoxImage<Voxel>`	each voxel's `Value`	`<f4`	`sizeof(Voxel)` = 8
`MultiVoxImage`	the active channel and component	`<f4`	`components * 4`
`EventTable`	the values, row-major	`<f8`	8 (packed)
`Collection<float>` and friends	the elements	`<f4`	packed
`MuonScatterCollection`	the events' bytes	`\|u1`	1

A MuonScatter has no single numpy spelling, so what is offered is the right pointer and the right length with the layout left to a caller that knows the type — honest, and enough for a kernel written against MuonScatter. MultiVoxImage refuses Magnitude: it is computed from three components, not stored, so there is no run of bytes to point at.

GetDeviceArray() moves the data to the device to answer, so a GPU consumer always gets something it can use; the point of PreferredDevice is that by then it is usually already there and nothing is copied.

5. From Python¶

The same interface, under the name the ecosystem already uses: GPU memory from Python. Short version:

image.SetPreferredDevice("Device")
array = numba.cuda.as_cuda_array(image)   # a view of VRAM, not a copy

6. The CUDA build¶

Everything above works with the driver alone. Compiling kernels needs the toolkit, which the environment carries (condaenv.yml):

micromamba install -n uLib -c conda-forge \
    cuda-version=12.9 cuda-nvcc=12.9 cuda-cudart-dev cuda-cccl \
    cuda-driver-dev cuda-nvrtc-dev

12.9 is what the other environments on this machine use, and CUDA 12.x covers the Tesla P40 (sm_61, which is CMAKE_CUDA_ARCHITECTURES in the top CMakeLists.txt). cuda-cccl is not optional: Core/Vector.h includes thrust/device_vector.h under USE_CUDA.

Then the cuda preset, which is the clang/ccache preset plus USE_CUDA=ON:

conan install . --output-folder=build/cuda --build=missing --profile=fast \
    -c tools.cmake.cmaketoolchain:user_presets=False \
    -c tools.cmake.cmaketoolchain:generator="Unix Makefiles"
cmake --preset cuda
cmake --build build/cuda -j"$(nproc)"

If the toolkit lives somewhere other than the environment, add -DCUDAToolkit_ROOT=<prefix> -DCMAKE_CUDA_COMPILER=<prefix>/bin/nvcc.

What changes in that build:

Math/VoxImage.cpp and Math/VoxRaytracer.cpp are compiled by nvcc (set_source_files_properties(... LANGUAGE CUDA)), as are the VoxImage and raytracer tests and Core/testing/VectorMetaAllocatorTest.cpp.
Core/Vector.h grows DeviceData(), DeviceBegin(), DeviceEnd() — thrust::device_ptr, so a collection's elements go straight into a thrust algorithm.
VoxRaytracerCUDA.hpp is compiled in.
The two targets below exist.
@cuda.jit compiles in a Python recipe. numba needs two pieces of the toolkit for that: libNVVM, and libcudart — it asks the runtime for its version to work out which architectures NVVM accepts, and reports none if it cannot find it, which looks exactly like "your GPU is too old". Both are on the loader path when the toolkit is installed in the environment; a toolkit in another prefix needs CUDA_HOME=<prefix> and that prefix's targets/x86_64-linux/lib on LD_LIBRARY_PATH. uLib.device.kernels_available() checks the whole chain, including that NVVM accepts this device's compute capability.

-allow-unsupported-compiler is in CMAKE_CUDA_FLAGS because the preset's host compiler is clang 21, which 12.9 does not list; nvcc then accepts it. CMake's own default host compiler is the environment's g++ 14.3, which 12.9 does support — that is what it picks, and what these builds used.

7. The example, and the tests¶

src/Math/examples/02_images_on_the_gpu.cpp (run by ctest, and shipped as assets/examples/math/02_images_on_the_gpu.gcx) draws a sphere into a 32³ image, moves it to the GPU, writes a value into device memory through the driver, reads it back through the host, and gives the memory back — printing what it holds at each step. With no GPU it runs the other branch and checks the honest answer. Its recipe does the slice profile on the device.

src/Math/examples/03_a_kernel_on_the_voxels.cpp is the CUDA build's example (and test): a 96³ image, a kernel that sums it per slice straight out of VRAM, a second that masks it in place, both checked against the host's own answer, with the two timings printed beside each other.

src/Math/testing/VoxImageKernelTest.cpp is the test that needs a kernel to ask its questions: that a kernel sees exactly what the host wrote; that the stride is right (the Count beside every Value comes back untouched); that one component of a vector channel can be written without touching the other two; that an EventTable reduces on the device; and — the one that matters most — that memory from DeviceMemory is usable by cudaMemcpy and by a kernel, which is the proof that retaining the primary context was right. Had it taken a context of its own, every pointer this library hands out would belong to a context nothing else is in.

src/Math/testing/VoxImageDeviceTest.cpp covers the property shape, the placement rules both ways, the strided view, the save/load round trip and DeviceMemory's own accounting; src/HEP/Geant/testing/CollectionControlTest.cpp covers the collections; src/Python/testing/device_memory_test.py covers the Python side. All three pass with a GPU and without one.

Data on the GPU: PreferredDevice, and what honours it¶