summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2016-03-02 14:06:19 +0100
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2016-04-18 11:54:18 +0200
commitb474576d36554f9322b57fedeeae493d88491f31 (patch)
treeeb1b69009e37cffa388067483f916f642c009e40
parentc366f2b07ce16c4ccdafc7cc4199fdac2d3ffef2 (diff)
downloadastra-b474576d36554f9322b57fedeeae493d88491f31.tar.gz
astra-b474576d36554f9322b57fedeeae493d88491f31.tar.bz2
astra-b474576d36554f9322b57fedeeae493d88491f31.tar.xz
astra-b474576d36554f9322b57fedeeae493d88491f31.zip
Add FDKWeighting option to standard cone_bp
-rw-r--r--cuda/3d/cone_bp.cu12
-rw-r--r--cuda/3d/dims3d.h4
2 files changed, 12 insertions, 4 deletions
diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu
index f077f0d..6bd9d16 100644
--- a/cuda/3d/cone_bp.cu
+++ b/cuda/3d/cone_bp.cu
@@ -77,6 +77,7 @@ bool bindProjDataTexture(const cudaArray* array)
//__launch_bounds__(32*16, 4)
+template<bool FDKWEIGHT>
__global__ void dev_cone_BP(void* D_volData, unsigned int volPitch, int startAngle,
int angleOffset, const astraCUDA3d::SDimensions3D dims,
float fOutputScale)
@@ -134,7 +135,10 @@ __global__ void dev_cone_BP(void* D_volData, unsigned int volPitch, int startAng
fU = fUNum * fr;
fV = fVNum * fr;
float fVal = tex3D(gT_coneProjTexture, fU, fAngle, fV);
- Z[idx] += fVal; // fr*fr*fVal;
+ if (FDKWEIGHT)
+ Z[idx] += fr*fr*fVal;
+ else
+ Z[idx] += fVal;
fUNum += fCu.z;
fVNum += fCv.z;
@@ -297,8 +301,10 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
for (unsigned int i = 0; i < angleCount; i += g_anglesPerBlock) {
// printf("Calling BP: %d, %dx%d, %dx%d to %p\n", i, dimBlock.x, dimBlock.y, dimGrid.x, dimGrid.y, (void*)D_volumeData.ptr);
- if (params.iRaysPerVoxelDim == 1)
- dev_cone_BP<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
+ if (params.bFDKWeighting)
+ dev_cone_BP<true><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
+ else if (params.iRaysPerVoxelDim == 1)
+ dev_cone_BP<false><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale);
else
dev_cone_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, params.iRaysPerVoxelDim, fOutputScale);
}
diff --git a/cuda/3d/dims3d.h b/cuda/3d/dims3d.h
index a15c67a..eb7045f 100644
--- a/cuda/3d/dims3d.h
+++ b/cuda/3d/dims3d.h
@@ -58,7 +58,8 @@ struct SProjectorParams3D {
iRaysPerDetDim(1), iRaysPerVoxelDim(1),
fOutputScale(1.0f),
fVolScaleX(1.0f), fVolScaleY(1.0f), fVolScaleZ(1.0f),
- ker(ker3d_default)
+ ker(ker3d_default),
+ bFDKWeighting(false)
{ }
unsigned int iRaysPerDetDim;
@@ -68,6 +69,7 @@ struct SProjectorParams3D {
float fVolScaleY;
float fVolScaleZ;
Cuda3DProjectionKernel ker;
+ bool bFDKWeighting;
};
}