Repository navigation
v1.2.11
Summary
Maintenance release on top of 1.2.10. No public-API changes; pure-adjoint
kernels and autograd Function classes are untouched, so autograd and
iterative-reconstruction paths are byte-for-byte equivalent to 1.2.10.
This release adds a new performance benchmark suite, coalesces the cone
FDK gather kernel's write path, documents several analytical-API
enhancements that were merged since 1.2.10, and ships a CHANGELOG.md
at the repo root.
Added
Parker short-scan demos in examples/fdk_cone.py and examples/fbp_fan.py
Both examples now expose an apply_parker switch that not only multiplies
by parker_weights but also samples angles on the minimal short-scan range
pi + 2 * gamma_max, where gamma_max is computed from the detector
half-width and sdd. Before this change the apply_parker flag in
fbp_fan.py was silently a no-op because the angular range stayed at
2*pi regardless of the flag.
tests/test_gradcheck.py
torch.autograd.gradcheck coverage for ParallelProjectorFunction,
FanProjectorFunction, and ConeProjectorFunction. Compares the
analytical backward Jacobian to a finite-difference numerical Jacobian
on small inputs with float32-calibrated tolerances. Strongest possible
autograd-correctness guard beyond the inner-product adjoint tests
shipped in 1.2.10.
tests/test_ramp_filter_windows.py
27 parametrised tests covering every ramp_filter_1d window option
(None/"ram-lak", "hann", "hanning", "hamming", "cosine",
"shepp-logan") at three layers: direct _ramp_window helper
(DC gain, Nyquist value, monotonicity), full ramp_filter_1d
end-to-end (shape, DC annihilation, padding correctness, rfft vs
complex-fft parity, sample_spacing scaling), and a high-frequency
attenuation sanity check on a step input.
tests/benchmarks/ — pytest-benchmark suite
New opt-in performance suite covering every CUDA kernel in the library:
forward projector, pure-adjoint backprojector, and the full analytical
FBP / FDK pipeline (cosine pre-weight + ramp + angular weights +
voxel-gather + analytical scale). Three sizes for each of the three
geometries — 27 benchmarks total.
Excluded from the default pytest tests/ run via
--ignore=tests/benchmarks in pytest.ini so the regular 57-test
suite stays fast (~15 s). Run explicitly with:
pytest tests/benchmarks/ --benchmark-only \
--benchmark-columns=min,median,mean,stddev,roundsSupports before/after baselines via --benchmark-save=<name> and
--benchmark-compare=<name>, which makes quantitative perf work
trivial going forward.
CHANGELOG.md
Keep-a-Changelog file at the repo root, detailed 1.2.11 and 1.2.10
entries, pointers to earlier versions on GitHub releases.
Changed
_cone_3d_fdk_backproject_kernel memory-access layout
The kernel now unpacks cuda.grid(3) as (iz, iy, ix) instead of
(ix, iy, iz), and cone_weighted_backproject launches the grid as
_grid_3d(D, H, W) instead of _grid_3d(W, H, D). This makes iz
warp-adjacent, which matches the innermost stride-1 axis of the WHD
output buffer, so the final d_vol[ix, iy, iz] = accum write becomes
fully coalesced.
Measured speedup on a 64^3 – 160^3 volume sweep at 360 views:
1.0x – 1.14x, within run-to-run noise. The kernel is read-bound
on the four scattered sinogram bilinear reads per view, not
write-bound, so the L2 cache was already absorbing the old
uncoalesced writes. The fix is kept anyway because it is the right
default and will benefit immediately if a future kernel edit shifts
work toward the write path.
Coalescing-behavior comments were added to _parallel_2d_fbp_backproject_kernel
and _fan_2d_fbp_backproject_kernel as well. Those 2D kernels were
already coalesced (their first grid index is warp-adjacent and matches
the innermost Nx axis of the (Ny, Nx) output buffer), but the
comments make the invariant explicit and guard against regressions.
Notes — performance research (no code change)
A full audit of all seven numba-cuda memory types (device, pinned,
mapped, managed, shared, local, constant) was performed against the
gather kernels. Conclusion: the kernels are already close to optimal
within numba-cuda's available primitives.
Key findings (documented in full in the project's memory notes):
cuda.const.array_likeis compile-time only in numba — it bakes
values at JIT time, so it cannot be used for the runtime-varying
cos/sin tables. A direct experiment confirmed this (mutating the
Python source array between kernel calls had no effect on the output).- A shared-memory cos/sin cache prototype was implemented and
benchmarked against the current global-memory version. Speedup was
0.98x–1.01x — within noise. The hardware L1 cache already handles
uniform-address reads (all threads in a warp read the sameiview
index) via broadcast coalescing, so the theoretical win from shared
memory does not materialise. - A TPB (thread-per-block) sweep across 12 configurations found
candidates that looked ~5% faster than the current(8, 8, 8)
default, but the 5% delta was smaller than the 10–30% measurement
noise on a thermally-variable consumer GPU, so no TPB change was
committed.
The cone FDK kernel runs at approximately 3.75 G voxel-view
iterations / second on a 128^3 / 360-view reconstruction, which is
only ~1% of peak compute and ~90 GB/s effective memory traffic — the
kernel is L1/L2 latency bound on the scattered sinogram bilinear
reads, not bandwidth or compute bound. Getting significantly more
speed from this access pattern inside numba-cuda is not currently
feasible with the memory-hierarchy primitives available to us.
Compatibility
Source-compatible with 1.2.10 for every public entry point. No
argument changes, no removals, no renames.
PyPI
Install with pip install diffct==1.2.11
(PyPI page).