Skip to content

Commit 7cef2b3

Browse files
AnderBiguriAnderBiguri
authored
Merge pull request #700 from rozenn1/tv_proximal_2d
2d handling for TV proximal algs
2 parents ce640c8 + 28881a9 commit 7cef2b3

8 files changed

Lines changed: 18 additions & 18 deletions

File tree

‎Common/CUDA/GD_TV.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -307,12 +307,12 @@ do { \
307307
size_t mem_free=mem_GPU_global-mem_auxiliary;
308308

309309
splits=(unsigned int)(ceil(((float)(3*mem_size_image)/(float)(deviceCount))/mem_free));
310-
// Now, there is an overhead here, as each splits should have 2 slices more, to accoutn for overlap of images.
310+
// Now, there is an overhead here, as each splits should have 2 slices more, to account for overlap of images.
311311
// lets make sure these 2 slices fit, if they do not, add 1 to splits.
312312
slices_per_split=(image_size[2]+deviceCount*splits-1)/(deviceCount*splits);
313313
mem_img_each_GPU=(mem_slice_image*(slices_per_split+buffer_length*2));
314314

315-
// if the new stuff does not fit in the GPU, it measn we are in the edge case where adding that extra slice will overflow memory
315+
// if the new stuff does not fit in the GPU, it means we are in the edge case where adding that extra slice will overflow memory
316316
if (mem_GPU_global< 3*mem_img_each_GPU+mem_auxiliary){
317317
// one more split should do the job, as its an edge case.
318318
splits++;

‎Common/CUDA/tv_proximal.cu‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -229,12 +229,12 @@ void cpy_from_host(float* device_array,float* host_array,
229229
size_t mem_free=mem_GPU_global;
230230

231231
splits=(unsigned int)(ceil(((float)(5*mem_size_image)/(float)(deviceCount))/mem_free));
232-
// Now, there is an overhead here, as each splits should have 2 slices more, to accoutn for overlap of images.
232+
// Now, there is an overhead here, as each splits should have 2 slices more, to account for overlap of images.
233233
// lets make sure these 2 slices fit, if they do not, add 1 to splits.
234234
slices_per_split=(image_size[2]+deviceCount*splits-1)/(deviceCount*splits);
235235
mem_img_each_GPU=(mem_slice_image*(slices_per_split+buffer_length*2));
236236

237-
// if the new stuff does not fit in the GPU, it measn we are in the edge case where adding that extra slice will overflow memory
237+
// if the new stuff does not fit in the GPU, it means we are in the edge case where adding that extra slice will overflow memory
238238
if (mem_GPU_global< 5*mem_img_each_GPU){
239239
// one more split should do the job, as its an edge case.
240240
splits++;
@@ -677,12 +677,12 @@ void cpy_from_host(float* device_array,float* host_array,
677677

678678
// Initial and last cases are special. These define the boundary condition. In our case, we are using Neumann boundary condition
679679
// so we need to copy the edge slice into the buffer
680-
if(is_first_chunk){
680+
if(is_first_chunk && image_size[2] > 1){
681681
for (unsigned int j=0;j<buffer_length;j++){
682-
cudaMemcpyAsync(device_array+pixels_per_slice*j, host_array+pixels_per_slice*(buffer_length-j), pixels_per_slice*sizeof(float), cudaMemcpyHostToDevice,stream);
683-
}
682+
cudaMemcpyAsync(device_array+pixels_per_slice*j, host_array+pixels_per_slice*(buffer_length-j), pixels_per_slice*sizeof(float), cudaMemcpyHostToDevice,stream);
683+
}
684684
}
685-
if(is_last_chunk){
685+
if(is_last_chunk && image_size[2] > 1){
686686

687687
for (unsigned int j=0;j<buffer_length;j++){
688688
cudaMemcpyAsync(device_array+bytes_device+pixels_per_slice*j, host_array+pixels_per_slice*(image_size[2]-j-2), pixels_per_slice*sizeof(float), cudaMemcpyHostToDevice,stream);

‎MATLAB/Demos/d16_2Dtomography.m‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -62,7 +62,7 @@
6262
angles=linspace(0,2*pi,100);
6363
phantom=single(phantom('Modified Shepp-Logan',geo.nVoxel(1)));
6464
projections=Ax(phantom,geo,angles);
65-
%% recosntruct
65+
%% reconstruct
6666

6767
imgOSSART=OS_SART(projections,geo,angles,40);
6868
imgASDPOCS=ASD_POCS(projections,geo,angles,40);
@@ -107,7 +107,7 @@
107107
phantom=single(phantom('Modified Shepp-Logan',geo.nVoxel(1)));
108108
% phantom=cat(3,phantom,phantom);
109109
projections=Ax(phantom,geo,angles,'interpolated');
110-
%% recosntruct
110+
%% reconstruct
111111

112112
imgOSSART=OS_SART(projections,geo,angles,40);
113113
% in 2D, the TV norm minimization happens in CPU, this causes the

‎MATLAB/Demos/d24_Algorithms06.m‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
%
33
%
44
% This demo presents the TIGRE capability of using image priors for
5-
% CT recosntruction. The most known algorithm for this is PICCS
5+
% CT reconstruction. The most known algorithm for this is PICCS
66
%
77
%--------------------------------------------------------------------------
88
%--------------------------------------------------------------------------

‎Python/demos/d16_2Dtomography.py‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -64,7 +64,7 @@
6464
head = sample_loader.load_head_phantom(geo.nVoxel)
6565
projections = tigre.Ax(head, geo, angles)
6666
tigre.plotSinogram(projections, 0)
67-
#%% recosntruct
67+
#%% reconstruct
6868

6969
imgOSSART = algs.ossart(projections, geo, angles, 40)
7070
imgCGLS = algs.cgls(projections, geo, angles, 40)
@@ -105,7 +105,7 @@
105105
angles = np.linspace(0, 2 * np.pi, 100)
106106
head = sample_loader.load_head_phantom(geo.nVoxel)
107107
projections = tigre.Ax(head, geo, angles)
108-
#%% recosntruct
108+
#%% reconstruct
109109

110110
imgOSSART = algs.ossart(projections, geo, angles, 40)
111111
imgCGLS = algs.cgls(projections, geo, angles, 40)

‎Python/tigre/algorithms/ista_algorithms.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@ class FISTA(IterativeReconAlg):
1919
Parameters
2020
----------
2121
:param proj: (np.ndarray, dtype=np.float32)
22-
Input data, shape = (geo.nDector, nangles)
22+
Input data, shape = (nangles, geo.nDetector)
2323
2424
:param geo: (tigre.geometry)
2525
Geometry of detector and image (see examples/Demo code)

‎Python/tigre/algorithms/iterative_recon_alg.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,7 @@ class IterativeReconAlg(object):
3636
Parameters
3737
----------
3838
:param proj: (np.ndarray, dtype=np.float32)
39-
Input data, shape = (geo.nDector, nangles)
39+
Input data, shape = (nangles, geo.nDetector)
4040
4141
:param geo: (tigre.geometry)
4242
Geometry of detector and image (see examples/Demo code)

‎Python/tigre/algorithms/pocs_algorithms.py‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ class initASD_POCS(IterativeReconAlg):
2020
Parameters
2121
----------
2222
:param proj: (np.ndarray, dtype=np.float32)
23-
Input data, shape = (geo.nDector, nangles)
23+
Input data, shape = (nangles, geo.nDetector)
2424
2525
:param geo: (tigre.geometry)
2626
Geometry of detector and image (see examples/Demo code)
@@ -192,7 +192,7 @@ class initPICCS(IterativeReconAlg):
192192
Parameters
193193
----------
194194
:param proj: (np.ndarray, dtype=np.float32)
195-
Input data, shape = (geo.nDector, nangles)
195+
Input data, shape = (nangles, geo.nDetector)
196196
197197
:param geo: (tigre.geometry)
198198
Geometry of detector and image (see examples/Demo code)
@@ -439,7 +439,7 @@ class initPCSD(IterativeReconAlg):
439439
Parameters
440440
----------
441441
:param proj: (np.ndarray, dtype=np.float32)
442-
Input data, shape = (geo.nDector, nangles)
442+
Input data, shape = (nangles, geo.nDetector)
443443
444444
:param geo: (tigre.geometry)
445445
Geometry of detector and image (see examples/Demo code)

0 commit comments

Comments
 (0)