kernelkit.kernels.VoxelDrivenConeBP

class kernelkit.kernels.VoxelDrivenConeBP(voxels_per_block: tuple = None, projs_per_block: int = None, volume_axes: tuple = (0, 1, 2), projection_axes: tuple = (0, 1, 2))

Voxel-driven conebeam backprojection kernel.

Attributes:

cuda_source

Returns the CUDA source code of the kernel

is_compiled

projection_axes

volume_axes

Methods

FLOAT_DTYPE	alias of float32
InterpolationMethod(value[, names, module, ...])	Texture fetching approach.
__call__(textures, volume, volume_geometry)	XrayBackprojection with conebeam geometry.
compile([max_projs, texture])	Compile the kernel.
geoms2params(projection_geometry, ...[, ...])	Compute kernel parameters
set_params(params)	Copy parameters to constant memory of the kernel.

__init__(voxels_per_block: tuple = None, projs_per_block: int = None, volume_axes: tuple = (0, 1, 2), projection_axes: tuple = (0, 1, 2))

Conebeam backprojection kernel.

Parameters:

voxels_per_blocktuple, optional

Number of voxels computed in one thread block, corresponding to the reversed volume.shape, not taking into account the mapping to world coordinates via volume_axes. I.e.: voxels_per_block[2] describes CUDA blockDim.x and voxels_per_block[0] describes how many slices i of volume[i, …] are processed in one thread. If your input volume is (8, 16, 32) voxels, maybe a good setting would be (32, 16, 8), regardless to which world axes (x, y, z) the 8, 16, and 32 voxels belong. Note voxels_per_block[0] * voxels_per_block[1] is constrained to the maximum number of CUDA threads, typically 1024.

projs_per_blockint, optional

Maximum number of projections processed in one thread block. If None defaults to LIMIT_PROJS_PER_BLOCK.

volume_axestuple, optional

Axes of the backprojection volume. Defaults to (0, 1, 2). The axes are in the order x, y, z. The first axis is the source- detector axis, the second is the horizontal orthogonal axis and the third is the vertical axis. ASTRA Toolbox uses (2, 1, 0).

projection_axestuple, optional

Axes of the projections to be backprojected. Defaults to (0, 1, 2). The axes are in the order angle, row, column.

See also

kernelkit.kernel.VoxelDrivenConeFP

Conebeam forward projection kernel

kernelkit.kernel.BaseKernel

Base class for CUDA kernels

Methods

__init__([voxels_per_block, ...])	Conebeam backprojection kernel.
compile([max_projs, texture])	Compile the kernel.
geoms2params(projection_geometry, ...[, ...])	Compute kernel parameters
set_params(params)	Copy parameters to constant memory of the kernel.

Attributes

SUPPORTED_DTYPES
cuda_source	Returns the CUDA source code of the kernel
is_compiled
projection_axes
projs_per_block
volume_axes
voxels_per_block