summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorValerii Sokolov <valerii.sokolov@uantwerpen.be>2015-04-09 13:14:01 +0200
committerValerii Sokolov <valerii.sokolov@uantwerpen.be>2015-04-09 13:14:01 +0200
commit4bb0a8cfc636582daa8ad62fc2f957239556be81 (patch)
tree75410351523c6685db5ba23e228d3080cc0b657e
parent3042b1369a96eef4798ea4280dd7aa1a8be2fcca (diff)
downloadastra-4bb0a8cfc636582daa8ad62fc2f957239556be81.tar.gz
astra-4bb0a8cfc636582daa8ad62fc2f957239556be81.tar.bz2
astra-4bb0a8cfc636582daa8ad62fc2f957239556be81.tar.xz
astra-4bb0a8cfc636582daa8ad62fc2f957239556be81.zip
Fixed a few CUDA 2D DART bugs.
* Mixed width and height led to incorrect work on rectangular images. * Incorrect weight calculation in `devDartSmoothingRadius` (#47).
-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);