/* ----------------------------------------------------------------------- Copyright: 2010-2015, iMinds-Vision Lab, University of Antwerp 2014-2015, CWI, Amsterdam Contact: astra@uantwerpen.be Website: http://sf.net/projects/astra-toolbox 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 . ----------------------------------------------------------------------- $Id$ */ #include #include #include #include #include #include "util3d.h" #ifdef STANDALONE #include "cone_fp.h" #include "testutil.h" #endif #include "dims3d.h" #include "arith3d.h" #include "cone_bp.h" #include "../2d/fft.h" #include "../../include/astra/Logging.h" typedef texture texture3D; static texture3D gT_coneProjTexture; namespace astraCUDA3d { static const unsigned int g_volBlockZ = 16; static const unsigned int g_anglesPerBlock = 64; static const unsigned int g_volBlockX = 32; static const unsigned int g_volBlockY = 16; static const unsigned int g_anglesPerWeightBlock = 16; static const unsigned int g_detBlockU = 32; static const unsigned int g_detBlockV = 32; static const unsigned g_MaxAngles = 12000; __constant__ float gC_angle_sin[g_MaxAngles]; __constant__ float gC_angle_cos[g_MaxAngles]; __constant__ float gC_angle[g_MaxAngles]; // per-detector u/v shifts? static bool bindProjDataTexture(const cudaArray* array) { cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); gT_coneProjTexture.addressMode[0] = cudaAddressModeBorder; gT_coneProjTexture.addressMode[1] = cudaAddressModeBorder; gT_coneProjTexture.addressMode[2] = cudaAddressModeBorder; gT_coneProjTexture.filterMode = cudaFilterModeLinear; gT_coneProjTexture.normalized = false; cudaBindTextureToArray(gT_coneProjTexture, array, channelDesc); // TODO: error value? return true; } __global__ void devFDK_preweight(void* D_projData, unsigned int projPitch, unsigned int startAngle, unsigned int endAngle, float fSrcOrigin, float fDetOrigin, float fZShift, float fDetUSize, float fDetVSize, const SDimensions3D dims) { float* projData = (float*)D_projData; int angle = startAngle + blockIdx.y * g_anglesPerWeightBlock + threadIdx.y; if (angle >= endAngle) return; const int detectorU = (blockIdx.x%((dims.iProjU+g_detBlockU-1)/g_detBlockU)) * g_detBlockU + threadIdx.x; const int startDetectorV = (blockIdx.x/((dims.iProjU+g_detBlockU-1)/g_detBlockU)) * g_detBlockV; int endDetectorV = startDetectorV + g_detBlockV; if (endDetectorV > dims.iProjV) endDetectorV = dims.iProjV; // We need the length of the central ray and the length of the ray(s) to // our detector pixel(s). const float fCentralRayLength = fSrcOrigin + fDetOrigin; const float fU = (detectorU - 0.5f*dims.iProjU + 0.5f) * fDetUSize; const float fT = fCentralRayLength * fCentralRayLength + fU * fU; float fV = (startDetectorV - 0.5f*dims.iProjV + 0.5f) * fDetVSize + fZShift; //const float fW = fCentralRayLength; //const float fW = fCentralRayLength * (M_PI / 2.0f) / (float)dims.iProjAngles; const float fW1 = fSrcOrigin * fDetUSize * fDetVSize; const float fW = fCentralRayLength * fW1 * fW1 * (M_PI / 2.0f) / (float)dims.iProjAngles; for (int detectorV = startDetectorV; detectorV < endDetectorV; ++detectorV) { const float fRayLength = sqrtf(fT + fV * fV); const float fWeight = fW / fRayLength; projData[(detectorV*dims.iProjAngles+angle)*projPitch+detectorU] *= fWeight; fV += fDetVSize; } } __global__ void devFDK_ParkerWeight(void* D_projData, unsigned int projPitch, unsigned int startAngle, unsigned int endAngle, float fSrcOrigin, float fDetOrigin, float fDetUSize, float fCentralFanAngle, const SDimensions3D dims) { float* projData = (float*)D_projData; int angle = startAngle + blockIdx.y * g_anglesPerWeightBlock + threadIdx.y; if (angle >= endAngle) return; const int detectorU = (blockIdx.x%((dims.iProjU+g_detBlockU-1)/g_detBlockU)) * g_detBlockU + threadIdx.x; const int startDetectorV = (blockIdx.x/((dims.iProjU+g_detBlockU-1)/g_detBlockU)) * g_detBlockV; int endDetectorV = startDetectorV + g_detBlockV; if (endDetectorV > dims.iProjV) endDetectorV = dims.iProjV; // We need the length of the central ray and the length of the projection // of our ray onto the central slice const float fCentralRayLength = fSrcOrigin + fDetOrigin; // TODO: Detector pixel size const float fU = (detectorU - 0.5f*dims.iProjU + 0.5f) * fDetUSize; //const float fGamma = atanf(fU / fCentralRayLength); //const float fBeta = gC_angle[angle]; const float fGamma = atanf(fU / fCentralRayLength); float fBeta = -gC_angle[angle]; if (fBeta < 0.0f) fBeta += 2*M_PI; if (fBeta >= 2*M_PI) fBeta -= 2*M_PI; // compute the weight depending on the location in the central fan's radon // space float fWeight; if (fBeta <= 0.0f) { fWeight = 0.0f; } else if (fBeta <= 2.0f*(fCentralFanAngle + fGamma)) { fWeight = sinf((M_PI / 4.0f) * fBeta / (fCentralFanAngle + fGamma)); fWeight *= fWeight; } else if (fBeta <= M_PI + 2*fGamma) { fWeight = 1.0f; } else if (fBeta <= M_PI + 2*fCentralFanAngle) { fWeight = sinf((M_PI / 4.0f) * (M_PI + 2.0f*fCentralFanAngle - fBeta) / (fCentralFanAngle - fGamma)); fWeight *= fWeight; } else { fWeight = 0.0f; } for (int detectorV = startDetectorV; detectorV < endDetectorV; ++detectorV) { projData[(detectorV*dims.iProjAngles+angle)*projPitch+detectorU] *= fWeight; } } // Perform the FDK pre-weighting and filtering bool FDK_PreWeight(cudaPitchedPtr D_projData, float fSrcOrigin, float fDetOrigin, float fZShift, float fDetUSize, float fDetVSize, bool bShortScan, const SDimensions3D& dims, const float* angles) { // The pre-weighting factor for a ray is the cosine of the angle between // the central line and the ray. dim3 dimBlock(g_detBlockU, g_anglesPerWeightBlock); dim3 dimGrid( ((dims.iProjU+g_detBlockU-1)/g_detBlockU)*((dims.iProjV+g_detBlockV-1)/g_detBlockV), (dims.iProjAngles+g_anglesPerWeightBlock-1)/g_anglesPerWeightBlock); int projPitch = D_projData.pitch/sizeof(float); devFDK_preweight<<>>(D_projData.ptr, projPitch, 0, dims.iProjAngles, fSrcOrigin, fDetOrigin, fZShift, fDetUSize, fDetVSize, dims); cudaTextForceKernelsCompletion(); if (bShortScan && dims.iProjAngles > 1) { ASTRA_DEBUG("Doing Parker weighting"); // We do short-scan Parker weighting // First, determine (in a very basic way) the interval that's // been scanned. We assume angles[0] is one of the endpoints of the // range. float fdA = angles[1] - angles[0]; while (fdA < -M_PI) fdA += 2*M_PI; while (fdA >= M_PI) fdA -= 2*M_PI; float fAngleBase; if (fdA >= 0.0f) { // going up from angles[0] fAngleBase = angles[dims.iProjAngles - 1]; } else { // going down from angles[0] fAngleBase = angles[0]; } // We pick the highest end of the range, and then // move all angles so they fall in (-2pi,0] float *fRelAngles = new float[dims.iProjAngles]; for (unsigned int i = 0; i < dims.iProjAngles; ++i) { float f = angles[i] - fAngleBase; while (f > 0) f -= 2*M_PI; while (f <= -2*M_PI) f += 2*M_PI; fRelAngles[i] = f; } cudaError_t e1 = cudaMemcpyToSymbol(gC_angle, fRelAngles, dims.iProjAngles*sizeof(float), 0, cudaMemcpyHostToDevice); assert(!e1); delete[] fRelAngles; float fCentralFanAngle = atanf(fDetUSize * (dims.iProjU*0.5f) / (fSrcOrigin + fDetOrigin)); devFDK_ParkerWeight<<>>(D_projData.ptr, projPitch, 0, dims.iProjAngles, fSrcOrigin, fDetOrigin, fDetUSize, fCentralFanAngle, dims); } cudaTextForceKernelsCompletion(); return true; } bool FDK_Filter(cudaPitchedPtr D_projData, cufftComplex * D_filter, const SDimensions3D& dims) { // The filtering is a regular ramp filter per detector line. int iPaddedDetCount = calcNextPowerOfTwo(2 * dims.iProjU); int iHalfFFTSize = calcFFTFourSize(iPaddedDetCount); int projPitch = D_projData.pitch/sizeof(float); // We process one sinogram at a time. float* D_sinoData = (float*)D_projData.ptr; cufftComplex * D_sinoFFT = NULL; allocateComplexOnDevice(dims.iProjAngles, iHalfFFTSize, &D_sinoFFT); bool ok = true; for (int v = 0; v < dims.iProjV; ++v) { ok = runCudaFFT(dims.iProjAngles, D_sinoData, projPitch, dims.iProjU, iPaddedDetCount, iHalfFFTSize, D_sinoFFT); if (!ok) break; applyFilter(dims.iProjAngles, iHalfFFTSize, D_sinoFFT, D_filter); ok = runCudaIFFT(dims.iProjAngles, D_sinoFFT, D_sinoData, projPitch, dims.iProjU, iPaddedDetCount, iHalfFFTSize); if (!ok) break; D_sinoData += (dims.iProjAngles * projPitch); } freeComplexOnDevice(D_sinoFFT); return ok; } bool FDK(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, const SConeProjection* angles, const SDimensions3D& dims, SProjectorParams3D params, bool bShortScan) { bool ok; // Generate filter // TODO: Check errors cufftComplex * D_filter; int iPaddedDetCount = calcNextPowerOfTwo(2 * dims.iProjU); int iHalfFFTSize = calcFFTFourSize(iPaddedDetCount); // NB: We don't support arbitrary cone_vec geometries here. // Only those that are vertical sub-geometries // (cf. CompositeGeometryManager) of regular cone geometries. assert(dims.iProjAngles > 0); const SConeProjection& p0 = angles[0]; // assuming U is in the XY plane, V is parallel to Z axis float fDetCX = p0.fDetSX + 0.5*dims.iProjU*p0.fDetUX; float fDetCY = p0.fDetSY + 0.5*dims.iProjU*p0.fDetUY; float fDetCZ = p0.fDetSZ + 0.5*dims.iProjV*p0.fDetVZ; float fSrcOrigin = sqrt(p0.fSrcX*p0.fSrcX + p0.fSrcY*p0.fSrcY); float fDetOrigin = sqrt(fDetCX*fDetCX + fDetCY*fDetCY); float fDetUSize = sqrt(p0.fDetUX*p0.fDetUX + p0.fDetUY*p0.fDetUY); float fDetVSize = abs(p0.fDetVZ); float fZShift = fDetCZ - p0.fSrcZ; float *pfAngles = new float[dims.iProjAngles]; for (unsigned int i = 0; i < dims.iProjAngles; ++i) { // FIXME: Sign/order pfAngles[i] = -atan2(angles[i].fSrcX, angles[i].fSrcY) + M_PI; } #if 1 ok = FDK_PreWeight(D_projData, fSrcOrigin, fDetOrigin, fZShift, fDetUSize, fDetVSize, bShortScan, dims, pfAngles); #else ok = true; #endif delete[] pfAngles; if (!ok) return false; #if 1 cufftComplex *pHostFilter = new cufftComplex[dims.iProjAngles * iHalfFFTSize]; memset(pHostFilter, 0, sizeof(cufftComplex) * dims.iProjAngles * iHalfFFTSize); genFilter(FILTER_RAMLAK, 1.0f, dims.iProjAngles, pHostFilter, iPaddedDetCount, iHalfFFTSize); allocateComplexOnDevice(dims.iProjAngles, iHalfFFTSize, &D_filter); uploadComplexArrayToDevice(dims.iProjAngles, iHalfFFTSize, pHostFilter, D_filter); delete [] pHostFilter; // Perform filtering ok = FDK_Filter(D_projData, D_filter, dims); // Clean up filter freeComplexOnDevice(D_filter); #endif if (!ok) return false; // Perform BP params.bFDKWeighting = true; //ok = FDK_BP(D_volumeData, D_projData, fSrcOrigin, fDetOrigin, 0.0f, 0.0f, fDetUSize, fDetVSize, dims, pfAngles); ok = ConeBP(D_volumeData, D_projData, dims, angles, params); if (!ok) return false; return true; } } #ifdef STANDALONE void dumpVolume(const char* filespec, const cudaPitchedPtr& data, const SDimensions3D& dims, float fMin, float fMax) { float* buf = new float[dims.iVolX*dims.iVolY]; unsigned int pitch = data.pitch / sizeof(float); for (int i = 0; i < dims.iVolZ; ++i) { cudaMemcpy2D(buf, dims.iVolX*sizeof(float), ((float*)data.ptr)+pitch*dims.iVolY*i, data.pitch, dims.iVolX*sizeof(float), dims.iVolY, cudaMemcpyDeviceToHost); char fname[512]; sprintf(fname, filespec, dims.iVolZ-i-1); saveImage(fname, dims.iVolY, dims.iVolX, buf, fMin, fMax); } } void dumpSinograms(const char* filespec, const cudaPitchedPtr& data, const SDimensions3D& dims, float fMin, float fMax) { float* bufs = new float[dims.iProjAngles*dims.iProjU]; unsigned int pitch = data.pitch / sizeof(float); for (int i = 0; i < dims.iProjV; ++i) { cudaMemcpy2D(bufs, dims.iProjU*sizeof(float), ((float*)data.ptr)+pitch*dims.iProjAngles*i, data.pitch, dims.iProjU*sizeof(float), dims.iProjAngles, cudaMemcpyDeviceToHost); char fname[512]; sprintf(fname, filespec, i); saveImage(fname, dims.iProjAngles, dims.iProjU, bufs, fMin, fMax); } } void dumpProjections(const char* filespec, const cudaPitchedPtr& data, const SDimensions3D& dims, float fMin, float fMax) { float* bufp = new float[dims.iProjV*dims.iProjU]; unsigned int pitch = data.pitch / sizeof(float); for (int i = 0; i < dims.iProjAngles; ++i) { for (int j = 0; j < dims.iProjV; ++j) { cudaMemcpy(bufp+dims.iProjU*j, ((float*)data.ptr)+pitch*dims.iProjAngles*j+pitch*i, dims.iProjU*sizeof(float), cudaMemcpyDeviceToHost); } char fname[512]; sprintf(fname, filespec, i); saveImage(fname, dims.iProjV, dims.iProjU, bufp, fMin, fMax); } } int main() { #if 0 SDimensions3D dims; dims.iVolX = 512; dims.iVolY = 512; dims.iVolZ = 512; dims.iProjAngles = 180; dims.iProjU = 1024; dims.iProjV = 1024; dims.iRaysPerDet = 1; cudaExtent extentV; extentV.width = dims.iVolX*sizeof(float); extentV.height = dims.iVolY; extentV.depth = dims.iVolZ; cudaPitchedPtr volData; // pitch, ptr, xsize, ysize cudaMalloc3D(&volData, extentV); cudaExtent extentP; extentP.width = dims.iProjU*sizeof(float); extentP.height = dims.iProjAngles; extentP.depth = dims.iProjV; cudaPitchedPtr projData; // pitch, ptr, xsize, ysize cudaMalloc3D(&projData, extentP); cudaMemset3D(projData, 0, extentP); #if 0 float* slice = new float[256*256]; cudaPitchedPtr ptr; ptr.ptr = slice; ptr.pitch = 256*sizeof(float); ptr.xsize = 256*sizeof(float); ptr.ysize = 256; for (unsigned int i = 0; i < 256*256; ++i) slice[i] = 1.0f; for (unsigned int i = 0; i < 256; ++i) { cudaExtent extentS; extentS.width = dims.iVolX*sizeof(float); extentS.height = dims.iVolY; extentS.depth = 1; cudaPos sp = { 0, 0, 0 }; cudaPos dp = { 0, 0, i }; cudaMemcpy3DParms p; p.srcArray = 0; p.srcPos = sp; p.srcPtr = ptr; p.dstArray = 0; p.dstPos = dp; p.dstPtr = volData; p.extent = extentS; p.kind = cudaMemcpyHostToDevice; cudaMemcpy3D(&p); #if 0 if (i == 128) { for (unsigned int j = 0; j < 256*256; ++j) slice[j] = 0.0f; } #endif } #endif SConeProjection angle[180]; angle[0].fSrcX = -1536; angle[0].fSrcY = 0; angle[0].fSrcZ = 0; angle[0].fDetSX = 1024; angle[0].fDetSY = -512; angle[0].fDetSZ = 512; angle[0].fDetUX = 0; angle[0].fDetUY = 1; angle[0].fDetUZ = 0; angle[0].fDetVX = 0; angle[0].fDetVY = 0; angle[0].fDetVZ = -1; #define ROTATE0(name,i,alpha) do { angle[i].f##name##X = angle[0].f##name##X * cos(alpha) - angle[0].f##name##Y * sin(alpha); angle[i].f##name##Y = angle[0].f##name##X * sin(alpha) + angle[0].f##name##Y * cos(alpha); } while(0) for (int i = 1; i < 180; ++i) { angle[i] = angle[0]; ROTATE0(Src, i, i*2*M_PI/180); ROTATE0(DetS, i, i*2*M_PI/180); ROTATE0(DetU, i, i*2*M_PI/180); ROTATE0(DetV, i, i*2*M_PI/180); } #undef ROTATE0 astraCUDA3d::ConeFP(volData, projData, dims, angle, 1.0f); //dumpSinograms("sino%03d.png", projData, dims, 0, 512); //dumpProjections("proj%03d.png", projData, dims, 0, 512); astraCUDA3d::zeroVolumeData(volData, dims); float* angles = new float[dims.iProjAngles]; for (int i = 0; i < 180; ++i) angles[i] = i*2*M_PI/180; astraCUDA3d::FDK(volData, projData, 1536, 512, 0, 0, dims, angles); dumpVolume("vol%03d.png", volData, dims, -20, 100); #else SDimensions3D dims; dims.iVolX = 1000; dims.iVolY = 999; dims.iVolZ = 500; dims.iProjAngles = 376; dims.iProjU = 1024; dims.iProjV = 524; dims.iRaysPerDet = 1; float* angles = new float[dims.iProjAngles]; for (int i = 0; i < dims.iProjAngles; ++i) angles[i] = -i*(M_PI)/360; cudaPitchedPtr volData = astraCUDA3d::allocateVolumeData(dims); cudaPitchedPtr projData = astraCUDA3d::allocateProjectionData(dims); astraCUDA3d::zeroProjectionData(projData, dims); astraCUDA3d::zeroVolumeData(volData, dims); timeval t; tic(t); for (int i = 0; i < dims.iProjAngles; ++i) { char fname[256]; sprintf(fname, "/home/wpalenst/tmp/Elke/proj%04d.png", i); unsigned int w,h; float* bufp = loadImage(fname, w,h); int pitch = projData.pitch / sizeof(float); for (int j = 0; j < dims.iProjV; ++j) { cudaMemcpy(((float*)projData.ptr)+dims.iProjAngles*pitch*j+pitch*i, bufp+dims.iProjU*j, dims.iProjU*sizeof(float), cudaMemcpyHostToDevice); } delete[] bufp; } printf("Load time: %f\n", toc(t)); //dumpSinograms("sino%03d.png", projData, dims, -8.0f, 256.0f); //astraCUDA3d::FDK(volData, projData, 7350, 62355, 0, 10, dims, angles); //astraCUDA3d::FDK(volData, projData, 7350, -380, 0, 10, dims, angles); tic(t); astraCUDA3d::FDK(volData, projData, 7383.29867, 0, 0, 10, dims, angles); printf("FDK time: %f\n", toc(t)); tic(t); dumpVolume("vol%03d.png", volData, dims, -65.9f, 200.0f); //dumpVolume("vol%03d.png", volData, dims, 0.0f, 256.0f); printf("Save time: %f\n", toc(t)); #endif } #endif