summaryrefslogtreecommitdiffstats
path: root/cuda/2d/darthelper.cu
diff options
context:
space:
mode:
authorWim van Aarle <wimvanaarle@gmail.com>2015-05-26 15:43:08 +0200
committerWim van Aarle <wimvanaarle@gmail.com>2015-05-26 15:43:08 +0200
commit3117c7a61a83ca28ba00211d1eb65eaf13a7307a (patch)
tree4851ae878cf499deb2ae2f77575b3c48ecaf5d04 /cuda/2d/darthelper.cu
parent9e884e760b71be354c08892f6bce2efe723917b3 (diff)
parent57668bcdd086145eb14598e32fd5d0dca711865c (diff)
downloadastra-3117c7a61a83ca28ba00211d1eb65eaf13a7307a.tar.gz
astra-3117c7a61a83ca28ba00211d1eb65eaf13a7307a.tar.bz2
astra-3117c7a61a83ca28ba00211d1eb65eaf13a7307a.tar.xz
astra-3117c7a61a83ca28ba00211d1eb65eaf13a7307a.zip
Merge remote-tracking branch 'upstream/master' into parallel_vec
Conflicts: astra_vc11.vcxproj.filters python/astra/functions.py src/ParallelBeamBlobKernelProjector2D.cpp src/ProjectionGeometry2D.cpp
Diffstat (limited to 'cuda/2d/darthelper.cu')
-rw-r--r--cuda/2d/darthelper.cu13
1 files changed, 7 insertions, 6 deletions
diff --git a/cuda/2d/darthelper.cu b/cuda/2d/darthelper.cu
index 28ca557..1d10d49 100644
--- a/cuda/2d/darthelper.cu
+++ b/cuda/2d/darthelper.cu
@@ -57,7 +57,7 @@ void roiSelect(float* out, float radius, unsigned int width, unsigned int height
// We abuse dims here...
SDimensions dims;
dims.iVolWidth = width;
- dims.iVolHeight = width;
+ dims.iVolHeight = height;
allocateVolumeData(D_data, pitch, dims);
copyVolumeToDevice(out, width, dims, D_data, pitch);
@@ -245,7 +245,7 @@ void dartMask(float* mask, const float* segmentation, unsigned int conn, unsigne
// We abuse dims here...
SDimensions dims;
dims.iVolWidth = width;
- dims.iVolHeight = width;
+ dims.iVolHeight = height;
allocateVolumeData(D_segmentationData, pitch, dims);
copyVolumeToDevice(segmentation, width, dims, D_segmentationData, pitch);
@@ -278,7 +278,7 @@ __global__ void devDartSmoothingRadius(float* out, const float* in, float b, uns
unsigned int x = threadIdx.x + 16*blockIdx.x;
unsigned int y = threadIdx.y + 16*blockIdx.y;
- // Sacrifice the border pixels to simplify the implementation.
+ // Sacrifice the border pixels to simplify the implementation.
if (x > radius-1 && x < width - radius && y > radius-1 && y < height - radius)
{
float* d = (float*)in;
@@ -286,9 +286,10 @@ __global__ void devDartSmoothingRadius(float* out, const float* in, float b, uns
unsigned int o2 = y*pitch+x;
int r = radius;
+ float count = 4*r*(r+1);
float res = -d[o2];
- for (int row = -r; row < r; row++)
+ for (int row = -r; row <= r; row++)
{
unsigned int o1 = (y+row)*pitch+x;
for (int col = -r; col <= r; col++)
@@ -297,7 +298,7 @@ __global__ void devDartSmoothingRadius(float* out, const float* in, float b, uns
}
}
- res *= b / 4*r*(r+1);
+ res *= b / count;
res += (1.0f-b) * d[o2];
m[o2] = res;
@@ -333,7 +334,7 @@ void dartSmoothing(float* out, const float* in, float b, unsigned int radius, un
// We abuse dims here...
SDimensions dims;
dims.iVolWidth = width;
- dims.iVolHeight = width;
+ dims.iVolHeight = height;
allocateVolumeData(D_inData, pitch, dims);
copyVolumeToDevice(in, width, dims, D_inData, pitch);