From 3042b1369a96eef4798ea4280dd7aa1a8be2fcca Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Mon, 30 Mar 2015 17:17:54 +0200 Subject: Initialize variables to avoid warning These variables are never used when uninitialized, but Visual Studio complains about them. --- cuda/2d/par_fp.cu | 2 +- cuda/3d/cone_fp.cu | 2 +- cuda/3d/par3d_fp.cu | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) (limited to 'cuda') 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); -- cgit v1.2.3 From 4bb0a8cfc636582daa8ad62fc2f957239556be81 Mon Sep 17 00:00:00 2001 From: Valerii Sokolov Date: Thu, 9 Apr 2015 13:14:01 +0200 Subject: Fixed a few CUDA 2D DART bugs. * Mixed width and height led to incorrect work on rectangular images. * Incorrect weight calculation in `devDartSmoothingRadius` (#47). --- cuda/2d/darthelper.cu | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) (limited to 'cuda') 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); -- cgit v1.2.3