diff options
Diffstat (limited to 'cuda')
| -rw-r--r-- | cuda/2d/darthelper.cu | 13 | ||||
| -rw-r--r-- | cuda/2d/par_fp.cu | 2 | ||||
| -rw-r--r-- | cuda/3d/cone_fp.cu | 2 | ||||
| -rw-r--r-- | cuda/3d/par3d_fp.cu | 2 |
4 files changed, 10 insertions, 9 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); diff --git a/cuda/2d/par_fp.cu b/cuda/2d/par_fp.cu index d0ca7ff..bb8b909 100644 --- a/cuda/2d/par_fp.cu +++ b/cuda/2d/par_fp.cu @@ -487,7 +487,7 @@ bool FP_simple_internal(float* D_volumeData, unsigned int volumePitch, unsigned int blockEnd = 0; bool blockVertical = false; for (unsigned int a = 0; a <= dims.iProjAngles; ++a) { - bool vertical; + bool vertical = false; // TODO: Having <= instead of < below causes a 5% speedup. // Maybe we should detect corner cases and put them in the optimal // group of angles. diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu index bda71ba..b36d2bc 100644 --- a/cuda/3d/cone_fp.cu +++ b/cuda/3d/cone_fp.cu @@ -340,7 +340,7 @@ bool ConeFP_Array_internal(cudaPitchedPtr D_projData, // tic(t); for (unsigned int a = 0; a <= angleCount; ++a) { - int dir; + int dir = -1; if (a != angleCount) { float dX = fabsf(angles[a].fSrcX - (angles[a].fDetSX + dims.iProjU*angles[a].fDetUX*0.5f + dims.iProjV*angles[a].fDetVX*0.5f)); float dY = fabsf(angles[a].fSrcY - (angles[a].fDetSY + dims.iProjU*angles[a].fDetUY*0.5f + dims.iProjV*angles[a].fDetVY*0.5f)); diff --git a/cuda/3d/par3d_fp.cu b/cuda/3d/par3d_fp.cu index 8d44540..b14c494 100644 --- a/cuda/3d/par3d_fp.cu +++ b/cuda/3d/par3d_fp.cu @@ -440,7 +440,7 @@ bool Par3DFP_Array_internal(cudaPitchedPtr D_projData, // tic(t); for (unsigned int a = 0; a <= angleCount; ++a) { - int dir; + int dir = -1; if (a != dims.iProjAngles) { float dX = fabsf(angles[a].fRayX); float dY = fabsf(angles[a].fRayY); |