summaryrefslogtreecommitdiffstats
path: root/cuda/3d/par3d_bp.cu
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <WillemJan.Palenstijn@uantwerpen.be>2013-07-01 22:34:11 +0000
committerwpalenst <WillemJan.Palenstijn@uantwerpen.be>2013-07-01 22:34:11 +0000
commitb2fc6c70434674d74551c3a6c01ffb3233499312 (patch)
treeb17f080ebc504ab85ebb7c3d89f917fd87ce9e00 /cuda/3d/par3d_bp.cu
downloadastra-b2fc6c70434674d74551c3a6c01ffb3233499312.tar.gz
astra-b2fc6c70434674d74551c3a6c01ffb3233499312.tar.bz2
astra-b2fc6c70434674d74551c3a6c01ffb3233499312.tar.xz
astra-b2fc6c70434674d74551c3a6c01ffb3233499312.zip
Update version to 1.3
Diffstat (limited to 'cuda/3d/par3d_bp.cu')
-rw-r--r--cuda/3d/par3d_bp.cu464
1 files changed, 464 insertions, 0 deletions
diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu
new file mode 100644
index 0000000..872b1eb
--- /dev/null
+++ b/cuda/3d/par3d_bp.cu
@@ -0,0 +1,464 @@
+/*
+-----------------------------------------------------------------------
+Copyright 2012 iMinds-Vision Lab, University of Antwerp
+
+Contact: astra@ua.ac.be
+Website: http://astra.ua.ac.be
+
+
+This file is part of the
+All Scale Tomographic Reconstruction Antwerp Toolbox ("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 <http://www.gnu.org/licenses/>.
+
+-----------------------------------------------------------------------
+$Id$
+*/
+
+#include <cstdio>
+#include <cassert>
+#include <iostream>
+#include <list>
+
+#include <cuda.h>
+#include "util3d.h"
+
+#ifdef STANDALONE
+#include "par3d_fp.h"
+#include "testutil.h"
+#endif
+
+#include "dims3d.h"
+
+typedef texture<float, 3, cudaReadModeElementType> texture3D;
+
+static texture3D gT_par3DProjTexture;
+
+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 g_MaxAngles = 1024;
+
+__constant__ float gC_Cux[g_MaxAngles];
+__constant__ float gC_Cuy[g_MaxAngles];
+__constant__ float gC_Cuz[g_MaxAngles];
+__constant__ float gC_Cuc[g_MaxAngles];
+__constant__ float gC_Cvx[g_MaxAngles];
+__constant__ float gC_Cvy[g_MaxAngles];
+__constant__ float gC_Cvz[g_MaxAngles];
+__constant__ float gC_Cvc[g_MaxAngles];
+
+
+static bool bindProjDataTexture(const cudaArray* array)
+{
+ cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
+
+ gT_par3DProjTexture.addressMode[0] = cudaAddressModeClamp;
+ gT_par3DProjTexture.addressMode[1] = cudaAddressModeClamp;
+ gT_par3DProjTexture.addressMode[2] = cudaAddressModeClamp;
+ gT_par3DProjTexture.filterMode = cudaFilterModeLinear;
+ gT_par3DProjTexture.normalized = false;
+
+ cudaBindTextureToArray(gT_par3DProjTexture, array, channelDesc);
+
+ // TODO: error value?
+
+ return true;
+}
+
+
+__global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAngle, const SDimensions3D dims)
+{
+ float* volData = (float*)D_volData;
+
+ int endAngle = startAngle + g_anglesPerBlock;
+ if (endAngle > dims.iProjAngles)
+ endAngle = dims.iProjAngles;
+
+ // threadIdx: x = rel x
+ // y = rel y
+
+ // blockIdx: x = x + y
+ // y = z
+
+
+ // TO TRY: precompute part of detector intersection formulas in shared mem?
+ // TO TRY: inner loop over z, gather ray values in shared mem
+
+ const int X = blockIdx.x % ((dims.iVolX+g_volBlockX-1)/g_volBlockX) * g_volBlockX + threadIdx.x;
+ const int Y = blockIdx.x / ((dims.iVolX+g_volBlockX-1)/g_volBlockX) * g_volBlockY + threadIdx.y;
+
+ if (X >= dims.iVolX)
+ return;
+ if (Y >= dims.iVolY)
+ return;
+
+ const int startZ = blockIdx.y * g_volBlockZ;
+ int endZ = startZ + g_volBlockZ;
+ if (endZ > dims.iVolZ)
+ endZ = dims.iVolZ;
+
+ float fX = X - 0.5f*dims.iVolX + 0.5f;
+ float fY = Y - 0.5f*dims.iVolY + 0.5f;
+ float fZ = startZ - 0.5f*dims.iVolZ + 0.5f;
+
+ for (int Z = startZ; Z < endZ; ++Z, fZ += 1.0f)
+ {
+
+ float fVal = 0.0f;
+ float fAngle = startAngle + 0.5f;
+
+ for (int angle = startAngle; angle < endAngle; ++angle, fAngle += 1.0f)
+ {
+
+ const float fCux = gC_Cux[angle];
+ const float fCuy = gC_Cuy[angle];
+ const float fCuz = gC_Cuz[angle];
+ const float fCuc = gC_Cuc[angle];
+ const float fCvx = gC_Cvx[angle];
+ const float fCvy = gC_Cvy[angle];
+ const float fCvz = gC_Cvz[angle];
+ const float fCvc = gC_Cvc[angle];
+
+ const float fUNum = fCuc + fX * fCux + fY * fCuy + fZ * fCuz;
+ const float fVNum = fCvc + fX * fCvx + fY * fCvy + fZ * fCvz;
+
+ const float fU = fUNum + 1.0f;
+ const float fV = fVNum + 1.0f;
+
+ fVal += tex3D(gT_par3DProjTexture, fU, fAngle, fV); // TODO: check order
+
+ }
+
+ volData[(Z*dims.iVolY+Y)*volPitch+X] += fVal;
+ }
+
+}
+
+// supersampling version
+__global__ void dev_par3D_BP_SS(void* D_volData, unsigned int volPitch, int startAngle, const SDimensions3D dims)
+{
+ float* volData = (float*)D_volData;
+
+ int endAngle = startAngle + g_anglesPerBlock;
+ if (endAngle > dims.iProjAngles)
+ endAngle = dims.iProjAngles;
+
+ // threadIdx: x = rel x
+ // y = rel y
+
+ // blockIdx: x = x + y
+ // y = z
+
+
+ // TO TRY: precompute part of detector intersection formulas in shared mem?
+ // TO TRY: inner loop over z, gather ray values in shared mem
+
+ const int X = blockIdx.x % ((dims.iVolX+g_volBlockX-1)/g_volBlockX) * g_volBlockX + threadIdx.x;
+ const int Y = blockIdx.x / ((dims.iVolX+g_volBlockX-1)/g_volBlockX) * g_volBlockY + threadIdx.y;
+
+ if (X >= dims.iVolX)
+ return;
+ if (Y >= dims.iVolY)
+ return;
+
+ const int startZ = blockIdx.y * g_volBlockZ;
+ int endZ = startZ + g_volBlockZ;
+ if (endZ > dims.iVolZ)
+ endZ = dims.iVolZ;
+
+ float fX = X - 0.5f*dims.iVolX + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim;
+ float fY = Y - 0.5f*dims.iVolY + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim;
+ float fZ = startZ - 0.5f*dims.iVolZ + 0.5f - 0.5f + 0.5f/dims.iRaysPerVoxelDim;
+
+ const float fSubStep = 1.0f/dims.iRaysPerVoxelDim;
+
+ for (int Z = startZ; Z < endZ; ++Z, fZ += 1.0f)
+ {
+
+ float fVal = 0.0f;
+ float fAngle = startAngle + 0.5f;
+
+ for (int angle = startAngle; angle < endAngle; ++angle, fAngle += 1.0f)
+ {
+ const float fCux = gC_Cux[angle];
+ const float fCuy = gC_Cuy[angle];
+ const float fCuz = gC_Cuz[angle];
+ const float fCuc = gC_Cuc[angle];
+ const float fCvx = gC_Cvx[angle];
+ const float fCvy = gC_Cvy[angle];
+ const float fCvz = gC_Cvz[angle];
+ const float fCvc = gC_Cvc[angle];
+
+ float fXs = fX;
+ for (int iSubX = 0; iSubX < dims.iRaysPerVoxelDim; ++iSubX) {
+ float fYs = fY;
+ for (int iSubY = 0; iSubY < dims.iRaysPerVoxelDim; ++iSubY) {
+ float fZs = fZ;
+ for (int iSubZ = 0; iSubZ < dims.iRaysPerVoxelDim; ++iSubZ) {
+
+ const float fUNum = fCuc + fXs * fCux + fYs * fCuy + fZs * fCuz;
+ const float fVNum = fCvc + fXs * fCvx + fYs * fCvy + fZs * fCvz;
+
+ const float fU = fUNum + 1.0f;
+ const float fV = fVNum + 1.0f;
+
+ fVal += tex3D(gT_par3DProjTexture, fU, fAngle, fV); // TODO: check order
+ fZs += fSubStep;
+ }
+ fYs += fSubStep;
+ }
+ fXs += fSubStep;
+ }
+
+ }
+
+ volData[(Z*dims.iVolY+Y)*volPitch+X] += fVal / (dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim*dims.iRaysPerVoxelDim);
+ }
+
+}
+
+bool Par3DBP_Array(cudaPitchedPtr D_volumeData,
+ cudaArray *D_projArray,
+ const SDimensions3D& dims, const SPar3DProjection* angles)
+{
+ bindProjDataTexture(D_projArray);
+
+
+ // transfer angles to constant memory
+ float* tmp = new float[dims.iProjAngles];
+
+#define TRANSFER_TO_CONSTANT(expr,name) do { for (unsigned int i = 0; i < dims.iProjAngles; ++i) tmp[i] = (expr) ; cudaMemcpyToSymbol(gC_##name, tmp, dims.iProjAngles*sizeof(float), 0, cudaMemcpyHostToDevice); } while (0)
+
+#define DENOM (angles[i].fRayX*angles[i].fDetUY*angles[i].fDetVZ - angles[i].fRayX*angles[i].fDetUZ*angles[i].fDetVY - angles[i].fRayY*angles[i].fDetUX*angles[i].fDetVZ + angles[i].fRayY*angles[i].fDetUZ*angles[i].fDetVX + angles[i].fRayZ*angles[i].fDetUX*angles[i].fDetVY - angles[i].fRayZ*angles[i].fDetUY*angles[i].fDetVX)
+
+
+ TRANSFER_TO_CONSTANT( ( - (angles[i].fRayY*angles[i].fDetVZ - angles[i].fRayZ*angles[i].fDetVY)) / DENOM , Cux );
+ TRANSFER_TO_CONSTANT( ( (angles[i].fRayX*angles[i].fDetVZ - angles[i].fRayZ*angles[i].fDetVX)) / DENOM , Cuy );
+ TRANSFER_TO_CONSTANT( (- (angles[i].fRayX*angles[i].fDetVY - angles[i].fRayY*angles[i].fDetVX) ) / DENOM , Cuz );
+ TRANSFER_TO_CONSTANT( (-(angles[i].fDetSY*angles[i].fDetVZ - angles[i].fDetSZ*angles[i].fDetVY)*angles[i].fRayX + (angles[i].fRayY*angles[i].fDetVZ - angles[i].fRayZ*angles[i].fDetVY)*angles[i].fDetSX - (angles[i].fRayY*angles[i].fDetSZ - angles[i].fRayZ*angles[i].fDetSY)*angles[i].fDetVX) / DENOM , Cuc );
+
+ TRANSFER_TO_CONSTANT( ((angles[i].fRayY*angles[i].fDetUZ - angles[i].fRayZ*angles[i].fDetUY) ) / DENOM , Cvx );
+ TRANSFER_TO_CONSTANT( (- (angles[i].fRayX*angles[i].fDetUZ - angles[i].fRayZ*angles[i].fDetUX) ) / DENOM , Cvy );
+ TRANSFER_TO_CONSTANT( ((angles[i].fRayX*angles[i].fDetUY - angles[i].fRayY*angles[i].fDetUX) ) / DENOM , Cvz );
+ TRANSFER_TO_CONSTANT( ((angles[i].fDetSY*angles[i].fDetUZ - angles[i].fDetSZ*angles[i].fDetUY)*angles[i].fRayX - (angles[i].fRayY*angles[i].fDetUZ - angles[i].fRayZ*angles[i].fDetUY)*angles[i].fDetSX + (angles[i].fRayY*angles[i].fDetSZ - angles[i].fRayZ*angles[i].fDetSY)*angles[i].fDetUX ) / DENOM , Cvc );
+
+#undef TRANSFER_TO_CONSTANT
+#undef DENOM
+
+ delete[] tmp;
+
+ dim3 dimBlock(g_volBlockX, g_volBlockY);
+
+ dim3 dimGrid(((dims.iVolX+g_volBlockX-1)/g_volBlockX)*((dims.iVolY+g_volBlockY-1)/g_volBlockY), (dims.iVolZ+g_volBlockZ-1)/g_volBlockZ);
+
+ // timeval t;
+ // tic(t);
+
+ for (unsigned int i = 0; i < dims.iProjAngles; 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 (dims.iRaysPerVoxelDim == 1)
+ dev_par3D_BP<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, dims);
+ else
+ dev_par3D_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, dims);
+ }
+
+ cudaTextForceKernelsCompletion();
+
+ // printf("%f\n", toc(t));
+
+ return true;
+}
+
+bool Par3DBP(cudaPitchedPtr D_volumeData,
+ cudaPitchedPtr D_projData,
+ const SDimensions3D& dims, const SPar3DProjection* angles)
+{
+ // transfer projections to array
+
+ cudaArray* cuArray = allocateProjectionArray(dims);
+ transferProjectionsToArray(D_projData, cuArray, dims);
+
+ bool ret = Par3DBP_Array(D_volumeData, cuArray, dims, angles);
+
+ cudaFreeArray(cuArray);
+
+ return ret;
+}
+
+
+}
+
+#ifdef STANDALONE
+int main()
+{
+ SDimensions3D dims;
+ dims.iVolX = 256;
+ dims.iVolY = 256;
+ dims.iVolZ = 256;
+ dims.iProjAngles = 180;
+ dims.iProjU = 512;
+ dims.iProjV = 512;
+ 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);
+
+ 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
+ }
+
+
+ SPar3DProjection angle[180];
+ angle[0].fRayX = 1;
+ angle[0].fRayY = 0;
+ angle[0].fRayZ = 0;
+
+ angle[0].fDetSX = 512;
+ angle[0].fDetSY = -256;
+ angle[0].fDetSZ = -256;
+
+ 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(Ray, 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::Par3DFP(volData, projData, dims, angle, 1.0f);
+#if 1
+ float* bufs = new float[180*512];
+
+ for (int i = 0; i < 512; ++i) {
+ cudaMemcpy(bufs, ((float*)projData.ptr)+180*512*i, 180*512*sizeof(float), cudaMemcpyDeviceToHost);
+
+ printf("%d %d %d\n", projData.pitch, projData.xsize, projData.ysize);
+
+ char fname[20];
+ sprintf(fname, "sino%03d.png", i);
+ saveImage(fname, 180, 512, bufs, 0, 512);
+ }
+
+ float* bufp = new float[512*512];
+
+ for (int i = 0; i < 180; ++i) {
+ for (int j = 0; j < 512; ++j) {
+ cudaMemcpy(bufp+512*j, ((float*)projData.ptr)+180*512*j+512*i, 512*sizeof(float), cudaMemcpyDeviceToHost);
+ }
+
+ char fname[20];
+ sprintf(fname, "proj%03d.png", i);
+ saveImage(fname, 512, 512, bufp, 0, 512);
+ }
+#endif
+ for (unsigned int i = 0; i < 256*256; ++i)
+ slice[i] = 0.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);
+ }
+
+ astraCUDA3d::Par3DBP(volData, projData, dims, angle);
+#if 1
+ float* buf = new float[256*256];
+
+ for (int i = 0; i < 256; ++i) {
+ cudaMemcpy(buf, ((float*)volData.ptr)+256*256*i, 256*256*sizeof(float), cudaMemcpyDeviceToHost);
+
+ printf("%d %d %d\n", volData.pitch, volData.xsize, volData.ysize);
+
+ char fname[20];
+ sprintf(fname, "vol%03d.png", i);
+ saveImage(fname, 256, 256, buf, 0, 60000);
+ }
+#endif
+
+}
+#endif