diff options
| author | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2016-03-02 14:06:19 +0100 |
|---|---|---|
| committer | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2016-04-18 11:54:18 +0200 |
| commit | b474576d36554f9322b57fedeeae493d88491f31 (patch) | |
| tree | eb1b69009e37cffa388067483f916f642c009e40 /cuda | |
| parent | c366f2b07ce16c4ccdafc7cc4199fdac2d3ffef2 (diff) | |
| download | astra-b474576d36554f9322b57fedeeae493d88491f31.tar.gz astra-b474576d36554f9322b57fedeeae493d88491f31.tar.bz2 astra-b474576d36554f9322b57fedeeae493d88491f31.tar.xz astra-b474576d36554f9322b57fedeeae493d88491f31.zip | |
Add FDKWeighting option to standard cone_bp
Diffstat (limited to 'cuda')
| -rw-r--r-- | cuda/3d/cone_bp.cu | 12 | ||||
| -rw-r--r-- | cuda/3d/dims3d.h | 4 |
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; }; } |