Update version to 1.3

1 files changed, 646 insertions, 0 deletions

diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu
new file mode 100644
index 0000000..ad0604c
--- /dev/null
+++ b/cuda/3d/fdk.cu
@@ -0,0 +1,646 @@
+/*
+-----------------------------------------------------------------------
+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 "cone_fp.h"
+#include "testutil.h"
+#endif
+
+#include "dims3d.h"
+#include "../2d/fft.h"
+
+typedef texture<float, 3, cudaReadModeElementType> 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 = 2048;
+
+__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<float>();
+
+	gT_coneProjTexture.addressMode[0] = cudaAddressModeClamp;
+	gT_coneProjTexture.addressMode[1] = cudaAddressModeClamp;
+	gT_coneProjTexture.addressMode[2] = cudaAddressModeClamp;
+	gT_coneProjTexture.filterMode = cudaFilterModeLinear;
+	gT_coneProjTexture.normalized = false;
+
+	cudaBindTextureToArray(gT_coneProjTexture, array, channelDesc);
+
+	// TODO: error value?
+
+	return true;
+}
+
+
+__global__ void devBP_FDK(void* D_volData, unsigned int volPitch, int startAngle, float fSrcOrigin, float fDetOrigin, float fSrcZ, float fDetZ, float fInvDetUSize, float fInvDetVSize, 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 - fSrcZ;
+
+	const float fU_base = 0.5f*dims.iProjU - 0.5f + 1.5f;
+	const float fV_base = 0.5f*dims.iProjV - 0.5f + 1.5f + (fDetZ-fSrcZ);
+
+	// Note re. fZ/rV_base: the computations below are all relative to the
+	// optical axis, so we do the Z-adjustments beforehand.
+
+	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 cos_theta = gC_angle_cos[angle];
+			const float sin_theta = gC_angle_sin[angle];
+
+			const float fR = fSrcOrigin;
+			const float fD = fR - fX * sin_theta + fY * cos_theta;
+			float fWeight = fR / fD;
+			fWeight *= fWeight;
+
+			const float fScaleFactor = (fR + fDetOrigin) / fD;
+			const float fU = fU_base + (fX*cos_theta+fY*sin_theta) * fScaleFactor * fInvDetUSize;
+			const float fV = fV_base + fZ * fScaleFactor * fInvDetVSize;
+
+			fVal += tex3D(gT_coneProjTexture, fU, fAngle, fV);
+
+		}
+
+		volData[(Z*dims.iVolY+Y)*volPitch+X] += fVal;
+//		projData[(angle*dims.iProjV+detectorV)*projPitch+detectorU] = 10.0f;
+//		if (threadIdx.x == 0 && threadIdx.y == 0) { printf("%d,%d,%d [%d / %d] -> %f\n", angle, detectorU, detectorV, (angle*dims.iProjV+detectorV)*projPitch+detectorU, projPitch, projData[(angle*dims.iProjV+detectorV)*projPitch+detectorU]); }
+	}
+
+}
+
+
+bool FDK_BP(cudaPitchedPtr D_volumeData,
+            cudaPitchedPtr D_projData,
+            float fSrcOrigin, float fDetOrigin,
+            float fSrcZ, float fDetZ, float fDetUSize, float fDetVSize,
+            const SDimensions3D& dims, const float* angles)
+{
+	// transfer projections to array
+
+	cudaArray* cuArray = allocateProjectionArray(dims);
+	transferProjectionsToArray(D_projData, cuArray, dims);
+
+	bindProjDataTexture(cuArray);
+
+	float* angle_sin = new float[dims.iProjAngles];
+	float* angle_cos = new float[dims.iProjAngles];
+
+	for (unsigned int i = 0; i < dims.iProjAngles; ++i) {
+		angle_sin[i] = sinf(angles[i]);
+		angle_cos[i] = cosf(angles[i]);
+	}
+	cudaError_t e1 = cudaMemcpyToSymbol(gC_angle_sin, angle_sin, dims.iProjAngles*sizeof(float), 0, cudaMemcpyHostToDevice);
+	cudaError_t e2 = cudaMemcpyToSymbol(gC_angle_cos, angle_cos, dims.iProjAngles*sizeof(float), 0, cudaMemcpyHostToDevice);
+	assert(e1 == cudaSuccess);
+	assert(e2 == cudaSuccess);
+
+	delete[] angle_sin;
+	delete[] angle_cos;
+
+	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) {
+		devBP_FDK<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, fSrcOrigin, fDetOrigin, fSrcZ, fDetZ, 1.0f / fDetUSize, 1.0f / fDetVSize, dims);
+	}
+
+	cudaTextForceKernelsCompletion();
+
+	cudaFreeArray(cuArray);
+
+	// printf("%f\n", toc(t));
+
+	return true;
+}
+
+__global__ void devFDK_preweight(void* D_projData, unsigned int projPitch, unsigned int startAngle, unsigned int endAngle, float fSrcOrigin, float fDetOrigin, float fSrcZ, float fDetZ, 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 + fDetZ - fSrcZ;
+
+	for (int detectorV = startDetectorV; detectorV < endDetectorV; ++detectorV)
+	{
+		const float fRayLength = sqrtf(fT + fV * fV);
+
+		const float fWeight = fCentralRayLength / fRayLength;
+
+		projData[(detectorV*dims.iProjAngles+angle)*projPitch+detectorU] *= fWeight;
+
+		fV += 1.0f;
+	}
+}
+
+__global__ void devFDK_ParkerWeight(void* D_projData, unsigned int projPitch, unsigned int startAngle, unsigned int endAngle, float fSrcOrigin, float fDetOrigin, float fSrcZ, float fDetZ, 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);
+	const float fBeta = -gC_angle[angle];
+
+	// 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_Filter(cudaPitchedPtr D_projData,
+                cufftComplex * D_filter,
+                float fSrcOrigin, float fDetOrigin,
+                float fSrcZ, float fDetZ,
+                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<<<dimGrid, dimBlock>>>(D_projData.ptr, projPitch, 0, dims.iProjAngles, fSrcOrigin, fDetOrigin, fSrcZ, fDetZ, fDetUSize, fDetVSize, dims);
+
+	cudaTextForceKernelsCompletion();
+
+	if (bShortScan) {
+		// We do short-scan Parker weighting
+
+		cudaError_t e1 = cudaMemcpyToSymbol(gC_angle, angles,
+		                                    dims.iProjAngles*sizeof(float), 0,
+		                                    cudaMemcpyHostToDevice);
+		assert(!e1);
+
+		// TODO: detector pixel size!
+		float fCentralFanAngle = atanf((dims.iProjU*0.5f) /
+		                               (fSrcOrigin + fDetOrigin));
+
+		devFDK_ParkerWeight<<<dimGrid, dimBlock>>>(D_projData.ptr, projPitch, 0, dims.iProjAngles, fSrcOrigin, fDetOrigin, fSrcZ, fDetZ, fDetUSize, fCentralFanAngle, dims);
+
+	}
+
+	cudaTextForceKernelsCompletion();
+
+
+	// The filtering is a regular ramp filter per detector line.
+
+	int iPaddedDetCount = calcNextPowerOfTwo(2 * dims.iProjU);
+	int iHalfFFTSize = calcFFTFourSize(iPaddedDetCount);
+
+	
+
+	// 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, 0,
+		                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,
+		                 0, 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,
+         float fSrcOrigin, float fDetOrigin,
+         float fSrcZ, float fDetZ, float fDetUSize, float fDetVSize,
+         const SDimensions3D& dims, const float* angles, bool bShortScan)
+{
+	bool ok;
+	// Generate filter
+	// TODO: Check errors
+	cufftComplex * D_filter;
+	int iPaddedDetCount = calcNextPowerOfTwo(2 * dims.iProjU);
+	int iHalfFFTSize = calcFFTFourSize(iPaddedDetCount);
+
+	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, fSrcOrigin, fDetOrigin,
+	                fSrcZ, fDetZ, fDetUSize, fDetVSize,
+	                bShortScan, dims, angles);
+
+	// Clean up filter
+	freeComplexOnDevice(D_filter);
+
+
+	if (!ok)
+		return false;
+
+	// Perform BP
+
+	ok = FDK_BP(D_volumeData, D_projData, fSrcOrigin, fDetOrigin, fSrcZ, fDetZ,
+	            fDetUSize, fDetVSize, dims, angles);
+
+	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