/* ----------------------------------------------------------------------- Copyright: 2010-2021, imec Vision Lab, University of Antwerp 2014-2021, CWI, Amsterdam Contact: astra@astra-toolbox.com Website: http://www.astra-toolbox.com/ This file is part of the ASTRA Toolbox. The ASTRA Toolbox is free software: you can redistribute it and/or modify it under the terms of the GNU General Public License as published by the Free Software Foundation, either version 3 of the License, or (at your option) any later version. The ASTRA Toolbox is distributed in the hope that it will be useful, but WITHOUT ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License for more details. You should have received a copy of the GNU General Public License along with the ASTRA Toolbox. If not, see . ----------------------------------------------------------------------- */ #include "astra/cuda/2d/util.h" #include "astra/cuda/2d/arith.h" #include #include #include #include #include namespace astraCUDA { static const unsigned g_MaxAngles = 2560; __constant__ float gC_angle[g_MaxAngles]; __constant__ float gC_angle_offset[g_MaxAngles]; __constant__ float gC_angle_detsize[g_MaxAngles]; // optimization parameters static const unsigned int g_anglesPerBlock = 16; static const unsigned int g_detBlockSize = 32; static const unsigned int g_blockSlices = 64; // projection for angles that are roughly horizontal // (detector roughly vertical) __global__ void FPhorizontal_simple(float* D_projData, unsigned int projPitch, cudaTextureObject_t tex, unsigned int startSlice, unsigned int startAngle, unsigned int endAngle, const SDimensions dims, float outputScale) { const int relDet = threadIdx.x; const int relAngle = threadIdx.y; int angle = startAngle + blockIdx.x * g_anglesPerBlock + relAngle; if (angle >= endAngle) return; const float theta = gC_angle[angle]; const float cos_theta = __cosf(theta); const float sin_theta = __sinf(theta); // compute start detector for this block/angle: const int detRegion = blockIdx.y; const int detector = detRegion * g_detBlockSize + relDet; // Now project the part of the ray to angle,detector through // slices startSlice to startSlice+g_blockSlices-1 if (detector < 0 || detector >= dims.iProjDets) return; const float fDetStep = -gC_angle_detsize[angle] / sin_theta; float fSliceStep = cos_theta / sin_theta; float fDistCorr; if (sin_theta > 0.0f) fDistCorr = outputScale / sin_theta; else fDistCorr = -outputScale / sin_theta; float fVal = 0.0f; // project detector on slice float fP = (detector - 0.5f*dims.iProjDets + 0.5f - gC_angle_offset[angle]) * fDetStep + (startSlice - 0.5f*dims.iVolWidth + 0.5f) * fSliceStep + 0.5f*dims.iVolHeight - 0.5f + 0.5f; float fS = startSlice + 0.5f; int endSlice = startSlice + g_blockSlices; if (endSlice > dims.iVolWidth) endSlice = dims.iVolWidth; if (dims.iRaysPerDet > 1) { fP += (-0.5f*dims.iRaysPerDet + 0.5f)/dims.iRaysPerDet * fDetStep; const float fSubDetStep = fDetStep / dims.iRaysPerDet; fDistCorr /= dims.iRaysPerDet; fSliceStep -= dims.iRaysPerDet * fSubDetStep; for (int slice = startSlice; slice < endSlice; ++slice) { for (int iSubT = 0; iSubT < dims.iRaysPerDet; ++iSubT) { fVal += tex2D(tex, fS, fP); fP += fSubDetStep; } fP += fSliceStep; fS += 1.0f; } } else { for (int slice = startSlice; slice < endSlice; ++slice) { fVal += tex2D(tex, fS, fP); fP += fSliceStep; fS += 1.0f; } } D_projData[angle*projPitch+detector] += fVal * fDistCorr; } // projection for angles that are roughly vertical // (detector roughly horizontal) __global__ void FPvertical_simple(float* D_projData, unsigned int projPitch, cudaTextureObject_t tex, unsigned int startSlice, unsigned int startAngle, unsigned int endAngle, const SDimensions dims, float outputScale) { const int relDet = threadIdx.x; const int relAngle = threadIdx.y; int angle = startAngle + blockIdx.x * g_anglesPerBlock + relAngle; if (angle >= endAngle) return; const float theta = gC_angle[angle]; const float cos_theta = __cosf(theta); const float sin_theta = __sinf(theta); // compute start detector for this block/angle: const int detRegion = blockIdx.y; const int detector = detRegion * g_detBlockSize + relDet; // Now project the part of the ray to angle,detector through // slices startSlice to startSlice+g_blockSlices-1 if (detector < 0 || detector >= dims.iProjDets) return; const float fDetStep = gC_angle_detsize[angle] / cos_theta; float fSliceStep = sin_theta / cos_theta; float fDistCorr; if (cos_theta < 0.0f) fDistCorr = -outputScale / cos_theta; else fDistCorr = outputScale / cos_theta; float fVal = 0.0f; float fP = (detector - 0.5f*dims.iProjDets + 0.5f - gC_angle_offset[angle]) * fDetStep + (startSlice - 0.5f*dims.iVolHeight + 0.5f) * fSliceStep + 0.5f*dims.iVolWidth - 0.5f + 0.5f; float fS = startSlice+0.5f; int endSlice = startSlice + g_blockSlices; if (endSlice > dims.iVolHeight) endSlice = dims.iVolHeight; if (dims.iRaysPerDet > 1) { fP += (-0.5f*dims.iRaysPerDet + 0.5f)/dims.iRaysPerDet * fDetStep; const float fSubDetStep = fDetStep / dims.iRaysPerDet; fDistCorr /= dims.iRaysPerDet; fSliceStep -= dims.iRaysPerDet * fSubDetStep; for (int slice = startSlice; slice < endSlice; ++slice) { for (int iSubT = 0; iSubT < dims.iRaysPerDet; ++iSubT) { fVal += tex2D(tex, fP, fS); fP += fSubDetStep; } fP += fSliceStep; fS += 1.0f; } } else { for (int slice = startSlice; slice < endSlice; ++slice) { fVal += tex2D(tex, fP, fS); fP += fSliceStep; fS += 1.0f; } } D_projData[angle*projPitch+detector] += fVal * fDistCorr; } // Coordinates of center of detector pixel number t: // x = (t - 0.5*nDets + 0.5 - fOffset) * fSize * cos(fAngle) // y = - (t - 0.5*nDets + 0.5 - fOffset) * fSize * sin(fAngle) static void convertAndUploadAngles(const SParProjection *projs, unsigned int nth, unsigned int ndets) { float *angles = new float[nth]; float *offsets = new float[nth]; float *detsizes = new float[nth]; for (int i = 0; i < nth; ++i) getParParameters(projs[i], ndets, angles[i], detsizes[i], offsets[i]); cudaMemcpyToSymbol(gC_angle, angles, nth*sizeof(float), 0, cudaMemcpyHostToDevice); cudaMemcpyToSymbol(gC_angle_offset, offsets, nth*sizeof(float), 0, cudaMemcpyHostToDevice); cudaMemcpyToSymbol(gC_angle_detsize, detsizes, nth*sizeof(float), 0, cudaMemcpyHostToDevice); delete [] angles; delete [] offsets; delete [] detsizes; } bool FP_simple_internal(float* D_volumeData, unsigned int volumePitch, float* D_projData, unsigned int projPitch, const SDimensions& dims, const SParProjection* angles, float outputScale) { assert(dims.iProjAngles <= g_MaxAngles); assert(angles); cudaArray* D_dataArray; cudaTextureObject_t D_texObj; if (!createTextureObject2D(D_volumeData, D_dataArray, D_texObj, volumePitch, dims.iVolWidth, dims.iVolHeight)) return false; convertAndUploadAngles(angles, dims.iProjAngles, dims.iProjDets); dim3 dimBlock(g_detBlockSize, g_anglesPerBlock); // detector block size, angles std::list streams; // Run over all angles, grouping them into groups of the same // orientation (roughly horizontal vs. roughly vertical). // Start a stream of grids for each such group. // TODO: Check if it's worth it to store this info instead // of recomputing it every FP. unsigned int blockStart = 0; unsigned int blockEnd = 0; bool blockVertical = false; for (unsigned int a = 0; a <= dims.iProjAngles; ++a) { 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. if (a != dims.iProjAngles) vertical = (fabsf(angles[a].fRayX) <= fabsf(angles[a].fRayY)); if (a == dims.iProjAngles || vertical != blockVertical) { // block done blockEnd = a; if (blockStart != blockEnd) { dim3 dimGrid((blockEnd-blockStart+g_anglesPerBlock-1)/g_anglesPerBlock, (dims.iProjDets+g_detBlockSize-1)/g_detBlockSize); // angle blocks, detector blocks // TODO: consider limiting number of handle (chaotic) geoms // with many alternating directions cudaStream_t stream; cudaStreamCreate(&stream); streams.push_back(stream); //printf("angle block: %d to %d, %d\n", blockStart, blockEnd, blockVertical); if (!blockVertical) for (unsigned int i = 0; i < dims.iVolWidth; i += g_blockSlices) FPhorizontal_simple<<>>(D_projData, projPitch, D_texObj, i, blockStart, blockEnd, dims, outputScale); else for (unsigned int i = 0; i < dims.iVolHeight; i += g_blockSlices) FPvertical_simple<<>>(D_projData, projPitch, D_texObj, i, blockStart, blockEnd, dims, outputScale); } blockVertical = vertical; blockStart = a; } } bool ok = true; for (std::list::iterator iter = streams.begin(); iter != streams.end(); ++iter) { ok &= checkCuda(cudaStreamSynchronize(*iter), "par_fp"); cudaStreamDestroy(*iter); } cudaFreeArray(D_dataArray); cudaDestroyTextureObject(D_texObj); return ok; } bool FP_simple(float* D_volumeData, unsigned int volumePitch, float* D_projData, unsigned int projPitch, const SDimensions& dims, const SParProjection* angles, float outputScale) { for (unsigned int iAngle = 0; iAngle < dims.iProjAngles; iAngle += g_MaxAngles) { SDimensions subdims = dims; unsigned int iEndAngle = iAngle + g_MaxAngles; if (iEndAngle >= dims.iProjAngles) iEndAngle = dims.iProjAngles; subdims.iProjAngles = iEndAngle - iAngle; bool ret; ret = FP_simple_internal(D_volumeData, volumePitch, D_projData + iAngle * projPitch, projPitch, subdims, angles + iAngle, outputScale); if (!ret) return false; } return true; } bool FP(float* D_volumeData, unsigned int volumePitch, float* D_projData, unsigned int projPitch, const SDimensions& dims, const SParProjection* angles, float outputScale) { return FP_simple(D_volumeData, volumePitch, D_projData, projPitch, dims, angles, outputScale); } }