SPECTGPU by danieldeidda · Pull Request #1687 · UCL/STIR

danieldeidda · 2026-03-04T17:02:27Z

…eton classe, atm a lot of copy from niftyGPU

Changes in this pull request

Testing performed

Related issues

Checklist before requesting a review

[] I have performed a self-review of my code
[] I have added docstrings/doxygen in line with the guidance in the developer guide
[] I have implemented unit tests that cover any new or modified functionality (if applicable)
[] The code builds and runs on my machine
[] documentation/release_XXX.md has been updated with any functionality change (if applicable)

Contribution Notes

Please tick the following:

The content of this Pull Request (the Contribution) is intentionally submitted for inclusion in STIR (the Work) under the terms and conditions of the Apache-2.0 License.
I (or my institution) have signed the STIR Contribution License Agreement (not required for small changes).

…eton classe, atm a lot of copy from niftyGPU

…calls the kernels; we need to double check that the kernel can write into the sino but probably we need to dodevice_to_host(stir_sino,cuda_array_created_by_forward)

… that calls the kernels; we need to double check that the kernel can write into the sino but probably we need to dodevice_to_host(stir_sino,cuda_array_created_by_forward)

addition of wrappers for actual_FP/BP removal of most of helper currently FP and BP work pretty well and are very close to spectUB rough speed estimate on debug mode were: SPECTUB 3min36s SPECTGPU 20s 3x3x3 gaussian kernel was fine for FP but introduced holes artifact in BP it was necessary to change the kernel to a 5x5x7 with sigma 2 to make a BPimage very close to SPECTUB. Speed above is related to 5x5x7.

codacy-production · 2026-09-09T17:26:28Z

Up to standards ✅

🟢 Issues 0 issues

Results:
0 new issues

View in Codacy

🟢 Metrics 65 complexity · 15 duplication

Metric Results

Complexity 65

Duplication 15

View in Codacy

_{NEW Get contextual insights on your PRs based on Codacy's metrics, along with PR and Jira context, without leaving GitHub. Enable AI reviewer}
_{TIP This summary will be updated as you push new changes.}

danieldeidda · 2026-09-10T10:18:22Z

the forward projction of a 150x150x200 image into a sino of 150, 80, 200 (tg,view,ax) values are pretty much identical.

the back projection on the other hand shows some scaling factor which I suppose we expect?

adapted PairProjectorByBinUsingSPECTGPU to be able to run reconstruction

danieldeidda · 2026-09-14T09:16:14Z

Very shallow tests on Attenuation and PSF but is seems that it is doing what I would expect it to be doing. But of course more tests are needed. Actual unit test are also needed like adjointness, etc.

reconstruction also seem to work fairly well:
the following is a reconstruction of simulated data (no noise) 2 iteration with

PSF+AC (PSF is just random to see the effect on recon and speed)
AC
no corrections

currently running with OMP off. when turning OMP on I get an assert aerror on ptr_access

KrisThielemans · 2026-09-14T09:31:28Z

Hi. looks like great progress! Obviously, I'm interested to know some timings :-)

the back projection on the other hand shows some scaling factor which I suppose we expect?

no, we don't. The SPECTUB projector backprojector is adjoint (as it uses the matrix), so has the correct scale.

danieldeidda · 2026-09-14T10:53:55Z

I thought there was some scaling in SPECTUB as I always need to multiply by the number of views to get images in the right scale. Anyway in debug mode with no OMP around 300s 2 iterations fully 3D PSF and AC I am reconstructing with SPECTUB now but it's been running for a while and not finished the first iter yet. GPU recon with no correction is 90 s and 160s only AC

danieldeidda · 2026-09-14T14:00:45Z

comparison with SPECTUB

I though SPECTUB would have some sort of scaling factor as we always need to multiply by the number of views to get "quantitative" values interestingly enough the reconstruction has the same scale.

Rough times in debug mode with OMP=off 2 iterations are:
GPU:
PSFAC 300s (CYAN)
AC 150s
no corr 90s
SPECTUB:
PSFAC 3400s (purple)

Added Hei Yin Jowett Chan as an author to the file.

pdate authorship

Jowett: Revert authorship added to the wrong file

danieldeidda · 2026-09-14T15:22:49Z

when OMP is on: the following debugging track is followed:

`array_to_device(elemT* dev_data, const Array<num_dimensions, elemT>& stir_array)
{
if (stir_array.is_contiguous())

{
  info("array_to_device contiguous", 100);

  cudaMemcpy(dev_data, stir_array.get_const_full_data_ptr(), stir_array.size_all() * sizeof(elemT), cudaMemcpyHostToDevice);

  **stir_array.release_const_full_data_ptr();**

}`

the assert happens when calling realease_const_full_data_ptr() which is called by SPECTGPUBackwardProjectorCUDA.cu: array_to_device(dev_image, stir_image);

KrisThielemans · 2026-09-14T16:11:41Z

interesting. only with OMP?

I think it's indeed the case that the get/release pair currently is not thread-safe, i.e. would have to be in a critical section to make it safe. I don't know about cudaMemCpy and thread-safety.

There's various issues here:

do we want to make it thread-safe?
does it matter for const access?
why are we copying the same image multiple times to the GPU? (That seems wrong/inefficient).

I guess the last one is the main one.

danieldeidda · 2026-09-15T11:20:47Z

Yes that's it! I think that is because the image needs to be updated as the function is called for every view. But then I should copy outside of this call maybe in start_accumulating and free in get_output? so that the BP just updates the output image

KrisThielemans · 2026-10-07T08:05:58Z

Not up to standards ⛔

🔴 Issues 2 high · 1 medium · 3 minor

Alerts:
⚠ 6 issues (≤ 0 issues of at least minor severity)

Results:
6 new issues

Category Results

BestPractice 1 medium

ErrorProne 2 high

CodeStyle 3 minor

View in Codacy

🟢 Metrics 98 complexity · 25 duplication

Metric Results

Complexity 98

Duplication 25

View in Codacy

_{NEW Get contextual insights on your PRs based on Codacy's metrics, along with PR and Jira context, without leaving GitHub. Enable AI reviewer}
_{TIP This summary will be updated as you push new changes.}

danieldeidda · 2026-10-08T11:22:25Z

I guess I will put a comment because from the commit it's a bit weird to read:
To make things faster and given that all arrays are of the same size a AllocatedStack structure was created:

all the allocations happen at setup so that in the loop we only initialise and assigne values. this has significantly speeded up performonce.
all arrays are now CuVec but at the end of the reconstruction we still have a free problem. The structure has a sestructor that calls resize(0) for all the cuvecs but the error is still there
the new stack structure create some race problem when using OMP because many threads where accessing the same object, so a thread_local AllocatedStack was used for everythread
in the BP part given that each thread contribute to the the same image it was also necessary to add an accumulation step as critical section, so a kernel for arrays accumulation was created;
updated timing given that spectUB becomes really fast once the matrix is created:
SPECTUB 2 iter 231 s (223 matrix + 4.2 s per full iter)
SPECTGPU 2 iter: 8.7s (6.4 matrix + 2.3 per full iter)

@casperdcl maybe have a look at my resize(0) destructor to see if I misinterpreted your comment.

creating OSMAPOSL and PpoissonLogLikelihood tests (these are running with the caveat that i had to make epsilon for the hessian product with input 0.001)

KrisThielemans

some changes in cuda_utilites for cuvec, which you're no longer using of course, but best to add them in anyway

KrisThielemans

some comments. mostly house-keeping.

I'm still a bit puzzled by the threading and GPU interaction, but need to do other things...

general clean up; moving cu_allocate and cu_free into cuda utilities; adressing all comments

danieldeidda · 2026-10-16T16:13:45Z

@KrisThielemans I think I am done

KrisThielemans · 2026-10-16T18:47:54Z

I think I am done

all tests are failing, so I guess not yet :-/

did you add the STIR_WITH_CUDA to the test?

KrisThielemans

we have src/buildblock/cuda_utilities.cu and src/buildblock/CUDA/cuda_utilities.cu

danieldeidda · 2026-10-18T06:18:09Z

I think I am done

all tests are failing, so I guess not yet :-/

did you add the STIR_WITH_CUDA to the test?

I thought i did bu maybe not everywher. will need to check

…g recontest data due to previous bag in setting number of sigma incorrectly; fixing test error for SPECTUB

KrisThielemans · 2026-10-19T13:01:13Z

Compiling fine, but SPECTUB test segfaulting. Do you have the same problem when running it locally?

danieldeidda · 2026-10-19T14:44:31Z

Compiling fine, but SPECTUB test segfaulting. Do you have the same problem when running it locally?

I actually never run it till the end, because is quite slow but running now and getting some difference in the gradient test

danieldeidda · 2026-10-19T15:15:05Z

Compiling fine, but SPECTUB test segfaulting. Do you have the same problem when running it locally?

it is giving seg fault locally too

danieldeidda · 2026-10-19T16:08:26Z

i found why the segfault happens but the issue will still be that the gradient test does not pass

danieldeidda and others added 9 commits March 4, 2026 17:01

SPECTGPU:starting point; created SPECTGPU_projector folder and schel…

b06adca

…eton classe, atm a lot of copy from niftyGPU

forward line integrals

c948b16

$JCHAN\jchan$

implemented a backward projection kernel in cuda

5d300cf

GPUSPECT: forward rotation and interpolation in cuda

d4cd638

Adding kernel push

9ca0e11

fixup! Adding kernel push

d731b5c

SPECTGPU: run pre-commit hooks on SPECTGPURotateAndGaussianInterpolate

8dfbdd1

joining all contributions; created an actual_forward_project() that …

aec9adf

…calls the kernels; we need to double check that the kernel can write into the sino but probably we need to dodevice_to_host(stir_sino,cuda_array_created_by_forward)

fixup! joining all contributions; created an actual_forward_project()…

edc60a8

… that calls the kernels; we need to double check that the kernel can write into the sino but probably we need to dodevice_to_host(stir_sino,cuda_array_created_by_forward)

casperdcl moved this to In Progress in 2026-03 AIRBI Hackathon Mar 16, 2026

casperdcl added this to 2026-03 AIRBI Hackathon Mar 16, 2026

KrisThielemans reviewed Mar 23, 2026

View reviewed changes

Comment thread src/recon_buildblock/SPECTGPU_projector/ForwardProjectorByBinSPECTGPU.cxx Outdated

KrisThielemans marked this pull request as draft March 23, 2026 15:46

KrisThielemans self-assigned this Mar 23, 2026

danieldeidda added 2 commits September 9, 2026 18:21

danieldeidda marked this pull request as ready for review September 9, 2026 17:28

SPECTGPU added Attenuation and depth dependent PSF;

74a840c

adapted PairProjectorByBinUsingSPECTGPU to be able to run reconstruction

jchan192 added 3 commits September 14, 2026 10:48

Update authorship in SPECTGPURotateAndGaussianInterpolate.cu

e7cbf28

Added Hei Yin Jowett Chan as an author to the file.

Update SPECTGPUProjection.cu

5d162c3

pdate authorship

Update SPECTGPURotateAndGaussianInterpolate.cu

864c871

Jowett: Revert authorship added to the wrong file

danieldeidda added 2 commits October 8, 2026 12:00

danieldeidda commented Oct 8, 2026

View reviewed changes

Comment thread src/include/stir/recon_buildblock/SPECTGPU_projector/ForwardProjectorByBinSPECTGPU.h Outdated

danieldeidda added 2 commits October 12, 2026 23:14

reducing double exp() calculation in pull kernel for a little speed up;

708ddfe

creating OSMAPOSL and PpoissonLogLikelihood tests (these are running with the caveat that i had to make epsilon for the hessian product with input 0.001)

reverting to cuda malloc(free) fixes problem with free()

dd22df7

danieldeidda commented Oct 15, 2026

View reviewed changes

Comment thread src/recon_buildblock/SPECTGPU_projector/SPECTGPUForwardProjectorCUDA.cu Outdated

KrisThielemans requested changes Oct 15, 2026

View reviewed changes

Comment thread src/buildblock/cuda_utilities.cu

Comment thread src/buildblock/cuda_utilities.cu

Comment thread src/buildblock/cuda_utilities.cu

Comment thread src/buildblock/cuda_utilities.cu

KrisThielemans reviewed Oct 15, 2026

View reviewed changes

Comment thread src/recon_test/test_PoissonLogLikelihoodWithLinearModelForMeanAndProjData.cxx Outdated

KrisThielemans requested changes Oct 15, 2026

View reviewed changes

samdporter reviewed Oct 15, 2026

View reviewed changes

Comment thread src/recon_buildblock/SPECTGPU_projector/ProjectorByBinPairUsingSPECTGPU.cxx Outdated

danieldeidda added 2 commits October 16, 2026 16:00

adding working demo using SIMIND;

eeabb81

general clean up; moving cu_allocate and cu_free into cuda utilities; adressing all comments

exposing SPECTGPU to SWIGPy

ecff145

KrisThielemans requested changes Oct 16, 2026

View reviewed changes

danieldeidda added 5 commits October 19, 2026 09:38

making allocatedSatck it's own file; running FBP for SPECTGPU updatin…

b2276d2

…g recontest data due to previous bag in setting number of sigma incorrectly; fixing test error for SPECTUB

moving SPECTGPU test within if stir_with_cuda

e487850

moving SPECTGPU test within if stir_with_cuda

cf5d756

some clean up

a0d8edb

hopefully fixing test error related to no GPU

208f897

Conversation

danieldeidda commented Mar 4, 2026

Changes in this pull request

Testing performed

Related issues

Checklist before requesting a review

Contribution Notes

Uh oh!

Uh oh!

codacy-production Bot commented Sep 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Up to standards ✅

Uh oh!

danieldeidda commented Sep 10, 2026

Uh oh!

danieldeidda commented Sep 14, 2026

Uh oh!

KrisThielemans commented Sep 14, 2026

Uh oh!

danieldeidda commented Sep 14, 2026

Uh oh!

danieldeidda commented Sep 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

danieldeidda commented Sep 14, 2026

Uh oh!

KrisThielemans commented Sep 14, 2026

Uh oh!

danieldeidda commented Sep 15, 2026

Uh oh!

KrisThielemans commented Oct 7, 2026

Uh oh!

casperdcl commented Oct 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

KrisThielemans commented Oct 7, 2026

Uh oh!

codacy-production Bot commented Oct 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Not up to standards ⛔

Uh oh!

danieldeidda commented Oct 8, 2026

Uh oh!

Uh oh!

Uh oh!

KrisThielemans left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

KrisThielemans left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

danieldeidda commented Oct 16, 2026

Uh oh!

KrisThielemans commented Oct 16, 2026

Uh oh!

KrisThielemans left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

codacy-production Bot commented Sep 9, 2026 •

edited

Loading

danieldeidda commented Sep 14, 2026 •

edited

Loading

casperdcl commented Oct 7, 2026 •

edited

Loading

codacy-production Bot commented Oct 8, 2026 •

edited

Loading

danieldeidda commented Oct 19, 2026 •

edited

Loading