-
Notifications
You must be signed in to change notification settings - Fork 95
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Merge pull request #1408 from Imraj-Singh/cuda_stir
Add CudaRelativeDifferencePrior (currently limited to 3x3x3 neighbourhood)
- Loading branch information
Showing
27 changed files
with
1,038 additions
and
129 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
46 changes: 46 additions & 0 deletions
46
recon_test_pack/CUDA/OSMAPOSL_test_sim_Parallelproj_CUDARDP.par
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,46 @@ | ||
| OSMAPOSLParameters := | ||
| ; test file for OSMAPOSL with a quadratic prior (and ray tracing projection matrix) | ||
| objective function type:= PoissonLogLikelihoodWithLinearModelForMeanAndProjData | ||
| PoissonLogLikelihoodWithLinearModelForMeanAndProjData Parameters:= | ||
|
|
||
| input file := my_prompts${suffix}.hs | ||
| zero end planes of segment 0:= 0 ; segment 0 is has direct and indirect planes | ||
| ; if disabled, defaults to maximum segment number in the file | ||
| ;maximum absolute segment number to process := 3 | ||
|
|
||
| projector pair type := parallelproj | ||
| Projector Pair Using Parallelproj Parameters:= | ||
| End Projector Pair Using Parallelproj Parameters:= | ||
|
|
||
| Bin Normalisation type := from projdata | ||
| Bin Normalisation From ProjData := | ||
| normalisation projdata filename:= my_acfs${suffix}.hs | ||
| End Bin Normalisation From ProjData:= | ||
|
|
||
| additive sinogram := my_additive_sinogram${suffix}.hs | ||
|
|
||
| xy output image size (in pixels) := 91 | ||
| zoom := .5 | ||
|
|
||
| use subset sensitivities:=1 | ||
| subset sensitivity filenames:= my_test_sim${suffix}_sens_PP_s${num_subsets}_%d.hv | ||
| recompute_sensitivity:=1 | ||
|
|
||
| prior type := Relative Difference Prior | ||
| Relative Difference Prior Parameters:= | ||
| penalisation factor := 0.5 | ||
| ; next defaults to 0, set to 1 for 2D inverse Euclidean weights, 0 for 3D | ||
| only 2D:= 0 | ||
| END Relative Difference Prior Parameters:= | ||
|
|
||
| end PoissonLogLikelihoodWithLinearModelForMeanAndProjData Parameters:= | ||
|
|
||
| output filename prefix := my_test_sim${suffix}_image_OSMAPOSL_PP_CUDARDP | ||
| number of subsets:= ${num_subsets} | ||
| start at subset:= 0 | ||
| number of subiterations:= 28 | ||
| save estimates at subiteration intervals:= ${num_subsets} | ||
|
|
||
| ;report objective function values interval := 1 | ||
|
|
||
| END := |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,69 @@ | ||
| /* | ||
| Copyright (C) 2024, University College London | ||
| This file is part of STIR. | ||
| SPDX-License-Identifier: Apache-2.0 | ||
| See STIR/LICENSE.txt for details | ||
| */ | ||
|
|
||
| #ifndef __stir_cuda_utilities_H__ | ||
| #define __stir_cuda_utilities_H__ | ||
|
|
||
| /*! | ||
| \file | ||
| \ingroup CUDA | ||
| \brief some utilities for STIR and CUDA | ||
| \author Kris Thielemans | ||
| */ | ||
| #include "stir/Array.h" | ||
| #include "stir/info.h" | ||
| #include <vector> | ||
|
|
||
| START_NAMESPACE_STIR | ||
|
|
||
| template <int num_dimensions, typename elemT> | ||
| inline void | ||
| 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(); | ||
| } | ||
| else | ||
| { | ||
| info("array_to_device non-contiguous", 100); | ||
| // Allocate host memory to get contiguous vector, copy array to it and copy from device to host | ||
| std::vector<elemT> tmp_data(stir_array.size_all()); | ||
| std::copy(stir_array.begin_all(), stir_array.end_all(), tmp_data.begin()); | ||
| cudaMemcpy(dev_data, tmp_data.data(), stir_array.size_all() * sizeof(elemT), cudaMemcpyHostToDevice); | ||
| } | ||
| } | ||
|
|
||
| template <int num_dimensions, typename elemT> | ||
| inline void | ||
| array_to_host(Array<num_dimensions, elemT>& stir_array, const elemT* dev_data) | ||
| { | ||
| if (stir_array.is_contiguous()) | ||
| { | ||
| info("array_to_host contiguous", 100); | ||
| cudaMemcpy(stir_array.get_full_data_ptr(), dev_data, stir_array.size_all() * sizeof(elemT), cudaMemcpyDeviceToHost); | ||
| stir_array.release_full_data_ptr(); | ||
| } | ||
| else | ||
| { | ||
| info("array_to_host non-contiguous", 100); | ||
| // Allocate host memory for the result and copy from device to host | ||
| std::vector<elemT> tmp_data(stir_array.size_all()); | ||
| cudaMemcpy(tmp_data.data(), dev_data, stir_array.size_all() * sizeof(elemT), cudaMemcpyDeviceToHost); | ||
| // Copy the data to the stir_array | ||
| std::copy(tmp_data.begin(), tmp_data.end(), stir_array.begin_all()); | ||
| } | ||
| } | ||
|
|
||
| END_NAMESPACE_STIR | ||
|
|
||
| #endif |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.