diff options
| -rw-r--r-- | cuda/3d/par3d_bp.cu | 67 |
1 files changed, 37 insertions, 30 deletions
diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu index 1dc75ce..748086e 100644 --- a/cuda/3d/par3d_bp.cu +++ b/cuda/3d/par3d_bp.cu @@ -35,10 +35,6 @@ along with the ASTRA Toolbox. If not, see <http://www.gnu.org/licenses/>. #include <cuda.h> -typedef texture<float, 3, cudaReadModeElementType> texture3D; - -static texture3D gT_par3DProjTexture; - namespace astraCUDA3d { static const unsigned int g_volBlockZ = 6; @@ -58,26 +54,31 @@ __constant__ DevPar3DParams gC_C[g_MaxAngles]; __constant__ float gC_scale[g_MaxAngles]; -static bool bindProjDataTexture(const cudaArray* array) +static bool bindProjDataTexture(cudaArray* array, cudaTextureObject_t& texObj) { - cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>(); - - gT_par3DProjTexture.addressMode[0] = cudaAddressModeBorder; - gT_par3DProjTexture.addressMode[1] = cudaAddressModeBorder; - gT_par3DProjTexture.addressMode[2] = cudaAddressModeBorder; - gT_par3DProjTexture.filterMode = cudaFilterModeLinear; - gT_par3DProjTexture.normalized = false; - - cudaBindTextureToArray(gT_par3DProjTexture, array, channelDesc); - - // TODO: error value? - - return true; + cudaChannelFormatDesc channelDesc = + cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); + + cudaResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = cudaResourceTypeArray; + resDesc.res.array.array = array; + + cudaTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = cudaAddressModeBorder; + texDesc.addressMode[1] = cudaAddressModeBorder; + texDesc.addressMode[2] = cudaAddressModeBorder; + texDesc.filterMode = cudaFilterModeLinear; + texDesc.readMode = cudaReadModeElementType; + texDesc.normalizedCoords = 0; + + return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "par3d_bp texture"); } template<unsigned int ZSIZE> -__global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale) +__global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, cudaTextureObject_t tex, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale) { float* volData = (float*)D_volData; @@ -125,7 +126,7 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAn for (int idx = 0; idx < ZSIZE; ++idx) { - float fVal = tex3D(gT_par3DProjTexture, fU, fAngle, fV); + float fVal = tex3D<float>(tex, fU, fAngle, fV); Z[idx] += fVal * fS; fU += fCu.z; @@ -144,7 +145,7 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAn } // supersampling version -__global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, int iRaysPerVoxelDim, float fOutputScale) +__global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, cudaTextureObject_t tex, int startAngle, int angleOffset, const SDimensions3D dims, int iRaysPerVoxelDim, float fOutputScale) { float* volData = (float*)D_volData; @@ -206,7 +207,7 @@ __global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int star const float fU = fCu.w + fXs * fCu.x + fYs * fCu.y + fZs * fCu.z; const float fV = fCv.w + fXs * fCv.x + fYs * fCv.y + fZs * fCv.z; - fVal += tex3D(gT_par3DProjTexture, fU, fAngle, fV) * fS; + fVal += tex3D<float>(tex, fU, fAngle, fV) * fS; fZs += fSubStep; } fYs += fSubStep; @@ -259,18 +260,21 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData, const SDimensions3D& dims, const SPar3DProjection* angles, const SProjectorParams3D& params) { - bindProjDataTexture(D_projArray); + cudaTextureObject_t D_texObj; + bindProjDataTexture(D_projArray, D_texObj); float fOutputScale = params.fOutputScale * params.fVolScaleX * params.fVolScaleY * params.fVolScaleZ; + bool ok = true; + for (unsigned int th = 0; th < dims.iProjAngles; th += g_MaxAngles) { unsigned int angleCount = g_MaxAngles; if (th + angleCount > dims.iProjAngles) angleCount = dims.iProjAngles - th; - bool ok = transferConstants(angles, angleCount, params); + ok = transferConstants(angles, angleCount, params); if (!ok) - return false; + break; dim3 dimBlock(g_volBlockX, g_volBlockY); @@ -283,23 +287,26 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData, // 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) { if (dims.iVolZ == 1) { - dev_par3D_BP<1><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale); + dev_par3D_BP<1><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, i, th, dims, fOutputScale); } else { - dev_par3D_BP<g_volBlockZ><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, fOutputScale); + dev_par3D_BP<g_volBlockZ><<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, 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); + dev_par3D_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), D_texObj, i, th, dims, params.iRaysPerVoxelDim, fOutputScale); } // TODO: Consider not synchronizing here, if possible. - if (!checkCuda(cudaThreadSynchronize(), "cone_bp")) - return false; + ok = checkCuda(cudaThreadSynchronize(), "cone_bp"); + if (!ok) + break; angles = angles + angleCount; // printf("%f\n", toc(t)); } + cudaDestroyTextureObject(D_texObj); + return true; } |