diff options
| author | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2021-11-16 10:21:35 +0100 |
|---|---|---|
| committer | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2021-11-16 10:21:35 +0100 |
| commit | 231dd3e5e28319aa16155efd9ec7fdc69834666b (patch) | |
| tree | 5b061e199bc3b7b5e39e48c8fbfa79eb5c5b906d /cuda/3d/par3d_bp.cu | |
| parent | 295ad96290c78a57ef982100014b1242a89d915c (diff) | |
| download | astra-231dd3e5e28319aa16155efd9ec7fdc69834666b.tar.gz astra-231dd3e5e28319aa16155efd9ec7fdc69834666b.tar.bz2 astra-231dd3e5e28319aa16155efd9ec7fdc69834666b.tar.xz astra-231dd3e5e28319aa16155efd9ec7fdc69834666b.zip | |
Add specialization to BP3D kernels for single slice volumes
Diffstat (limited to 'cuda/3d/par3d_bp.cu')
| -rw-r--r-- | cuda/3d/par3d_bp.cu | 14 |
1 files changed, 9 insertions, 5 deletions
diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu index 004fc99..d356b9f 100644 --- a/cuda/3d/par3d_bp.cu +++ b/cuda/3d/par3d_bp.cu @@ -41,8 +41,7 @@ static texture3D gT_par3DProjTexture; namespace astraCUDA3d { -#define ZSIZE 6 -static const unsigned int g_volBlockZ = ZSIZE; +static const unsigned int g_volBlockZ = 6; static const unsigned int g_anglesPerBlock = 32; static const unsigned int g_volBlockX = 16; @@ -77,6 +76,7 @@ static bool bindProjDataTexture(const cudaArray* array) } +template<unsigned int ZSIZE> __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale) { float* volData = (float*)D_volData; @@ -281,9 +281,13 @@ bool Par3DBP_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_par3D_BP<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale); - else + if (params.iRaysPerVoxelDim == 1) { + if (dims.iVolZ == 1) { + dev_par3D_BP<1><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale); + } else { + dev_par3D_BP<g_volBlockZ><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale); + } + } else dev_par3D_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, params.iRaysPerVoxelDim, fOutputScale); } |