elsa projectors_cuda

SiddonsMethodCUDA

template<typename data_t>
class elsa::SiddonsMethodCUDA

GPU-operator representing the discretized X-ray transform in 2d/3d using Siddon’s method.

The volume is traversed along the rays as specified by the

Geometry. Each ray is traversed in a contiguous fashion (i.e. along long voxel borders, not diagonally) and each traversed voxel is counted as a hit with weight according to the length of the path of the ray through the voxel.
Author

Nikola Dinev

Template Parameters

  • data_t: data type for the domain and range of the operator, defaulting to real_t

The geometry is represented as a list of projection matrices (see class Geometry), one for each acquisition pose.

Forward projection is accomplished using apply(), backward projection using applyAdjoint(). This projector is matched.

Currently only utilizes a single GPU. Volume and images should both fit in device memory at the same time.

Public Functions

SiddonsMethodCUDA(const VolumeDescriptor &domainDescriptor, const DetectorDescriptor &rangeDescriptor)

Constructor for Siddon’s method traversal.

The domain is expected to be 2 or 3 dimensional (volSizeX, volSizeY, [volSizeZ]), the range is expected to be matching the domain (detSizeX, [detSizeY], acqPoses).

Parameters

  • [in] domainDescriptor: describing the domain of the operator (the volume)

  • [in] rangeDescriptor: describing the range of the operator (the sinogram)

~SiddonsMethodCUDA() override = default

destructor

Protected Functions

SiddonsMethodCUDA(const SiddonsMethodCUDA<data_t>&) = default

default copy constructor, hidden from non-derived classes to prevent potential slicing

void forward(const BoundingBox &aabb, const DataContainer<data_t> &x, DataContainer<data_t> &Ax) const

apply Siddon’s method (i.e. forward projection)

void backward(const BoundingBox &aabb, const DataContainer<data_t> &y, DataContainer<data_t> &Aty) const

apply the adjoint of Siddon’s method (i.e. backward projection)

SiddonsMethodCUDA<data_t> *_cloneImpl() const

implement the polymorphic clone operation

bool _isEqual(const LinearOperator<data_t> &other) const

implement the polymorphic comparison operation

Private Members

thrust::device_vector<real_t> _invProjMatrices

inverse of of projection matrices; stored column-wise on GPU

thrust::device_vector<real_t> _rayOrigins

ray origins for each acquisition angle, TODO: enable kernel to have this as data_t

Private Static Attributes

const unsigned int THREADS_PER_BLOCK = TraverseSiddonsCUDA<data_t>::MAX_THREADS_PER_BLOCK

threads per block used in the kernel execution configuration

JosephsMethodCUDA

template<typename data_t>
class elsa::JosephsMethodCUDA

GPU-operator representing the discretized X-ray transform in 2d/3d using Joseph’s method.

The volume is traversed along the rays as specified by the

Geometry. For interior voxels the sampling point is located in the middle of the two planes orthogonal to the main direction of the ray. For boundary voxels the sampling point is located at the center of the ray intersection with the voxel.
Author

Nikola Dinev

Template Parameters

  • data_t: data type for the domain and range of the operator, defaulting to real_t

The geometry is represented as a list of projection matrices (see class Geometry), one for each acquisition pose.

Forward projection is accomplished using apply(), backward projection using applyAdjoint(). The projector provides two implementations for the backward projection. The slow version is matched, while the fast one is not.

Currently only utilizes a single GPU. Volume and images should both fit in device memory at the same time.

Warning

Hardware interpolation is only supported for JosephsMethodCUDA<float>

Warning

Hardware interpolation is significantly less accurate than the software interpolation

Public Functions

JosephsMethodCUDA(const VolumeDescriptor &domainDescriptor, const DetectorDescriptor &rangeDescriptor, bool fast = true)

Constructor for Joseph’s traversal method.

The domain is expected to be 2 or 3 dimensional (volSizeX, volSizeY, [volSizeZ]), the range is expected to be matching the domain (detSizeX, [detSizeY], acqPoses).

Parameters

  • [in] domainDescriptor: describing the domain of the operator (the volume)

  • [in] rangeDescriptor: describing the range of the operator (the sinogram)

  • [in] fast: performs fast backward projection if set, otherwise matched; forward projection is unaffected

~JosephsMethodCUDA() override

destructor

Protected Functions

JosephsMethodCUDA(const JosephsMethodCUDA<data_t> &other)

copy constructor, used for cloning

void forward(const BoundingBox &aabb, const DataContainer<data_t> &x, DataContainer<data_t> &Ax) const

apply Siddon’s method (i.e. forward projection)

void backward(const BoundingBox &aabb, const DataContainer<data_t> &y, DataContainer<data_t> &Aty) const

apply the adjoint of Siddon’s method (i.e. backward projection)

JosephsMethodCUDA<data_t> *_cloneImpl() const

implement the polymorphic clone operation

bool _isEqual(const LinearOperator<data_t> &other) const

implement the polymorphic comparison operation

Private Types

using cudaArrayFlags = unsigned int

convenience typedef for cuda array flags

Private Functions

template<ContainerCpyKind direction, bool async = true>
void copy3DDataContainer(void *hostData, const cudaPitchedPtr &gpuData, const cudaExtent &extent) const

Copies contents of a 3D data container between GPU and host memory.

Note that hostData is expected to be a pointer to a linear memory region with no padding between dimensions - e.g. the data in

DataContainer is stored as a vector with no extra padding, and the pointer to the start of the memory region can be retrieved as follows:
Template Parameters

  • direction: specifies the direction of the copy operation

  • async: whether the copy should be performed asynchronously wrt. the host

Parameters

  • hostData: pointer to host data

  • gpuData: pointer to gpu data

  • [in] extent: specifies the amount of data to be copied

DataContainer x; void* hostData = (void*)&x[0];

template<cudaArrayFlags flags = 0U>
std::pair<cudaTextureObject_t, cudaArray*> copyTextureToGPU(const DataContainer<data_t> &hostData) const

Copies the entire contents of DataContainer to the GPU texture memory.

Return

a pair of the created texture object and its associated cudaArray

Template Parameters

  • cudaArrayFlags: flags used for the creation of the cudaArray which will contain the data

Parameters

  • [in] hostData: the host data container

Private Members

DetectorDescriptor &_detectorDescriptor

Reference to DetectorDescriptor stored in LinearOperator.

VolumeDescriptor &_volumeDescriptor

Reference to VolumeDescriptor stored in LinearOperator.

const bool _fast

flag specifying which version of the backward projection should be used

cudaPitchedPtr _projInvMatrices

inverse of of projection matrices; stored column-wise on GPU

cudaPitchedPtr _projMatrices

projection matrices; stored column-wise on GPU

cudaPitchedPtr _rayOrigins

ray origins for each acquisition angle

Private Static Attributes

const unsigned int THREADS_PER_BLOCK = TraverseJosephsCUDA<data_t>::MAX_THREADS_PER_BLOCK

threads per block used in the kernel execution configuration