Image reconstruction algorithm problem

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all.

This is part of the pseudo code I am implementing in CUDA as part of an image reconstruction algorithm:

for each xbin(0->detectorXDim/2-1):

     for each ybin(0->detectorYDim-1):

          rayInit=(xbin*xBinSize+0.5,ybin*xBinSize+0.5,-detectordistance)

          rayEnd=beamFocusCoord

          slopeVector=rayEnd-rayInit

          //knowing that r=rayInit+t*slopeVector;

          //x=rayInit[0]+t*slopeVector[0]

          //y=rayInit[1]+t*slopeVector[1]

          //z=rayInit[2]+t*slopeVector[2]

          //to find ray xx intersections:

          for each xinteger(xbin+1->detectorXDim/2):

                    solve t for x=xinteger*xBinSize;

                    find corresponding y and z

                    add to intersections array

          //find ray yy intersections(analogous to xx intersections)

          //find ray zz intersections(analogous to xx intersections)

So far, this is what I have come up with:

__global__ void sysmat(int xfocus,int yfocus, int zfocus, int xbin,int xbinsize,int ybin,int ybinsize, int zbin, int projecoes){

	int tx=threadIdx.x, ty=threadIdx.y,tz=threadIdx.z, bx=blockIdx.x, by=blockIdx.y,i,x,y,z;

	int idx=ty+by*blocksize;

	int idy=tx+bx*blocksize;

	int slopeVectorx=xfocus-idx*xbinsize+0.5;

	int slopeVectory=yfocus-idy*ybinsize+0.5;

	int slopeVectorz=zfocus-zdetector;

	__syncthreads();

	

	//points where the ray intersects x axis

	int xint=idx+1;

        int yint=idy+1;

	int*intersectionsx[(detectorXDim/2-xint)+(detectorYDim-yint)+(zfocus)];

	int*intersectionsy[(detectorXDim/2-xint)+(detectorYDim-yint)+(zfocus)];

	int*intersectionsz[(detectorXDim/2-xint)+(detectorYDim-yint)+(zfocus)];

	for(xint=xint; xint<detectorXDim/2;xint++){

		x=xint*xbinsize;

		t=(x-idx)/slopeVectorx;

		y=idy+t*slopeVectory;

		z=z+t*slopeVectorz;

	

		intersectionsx[xint-1]=x;

		intersectionsy[xint-1]=y;

		intersectionsz[xint-1]=z;

		__syncthreads();

	}

...

}

So, this is just a piece of the code. I know that there might be some errors(you can point them if they are blatantly wrong) but what I am more concerned is this:

Each thread(which corresponds to a detector bin) needs three arrays so it can save the points where the ray(which passes through this thread/bin) intersects multiples of the x,y and z axis. Each array’s length depend on the place of the thread/bin(it’s index) in the detector and on the beamFocusCoord(which are fixed). In order to do this I wrote this piece of code, which I am almost certain can not be done:

int*intersectionsx[(detectorXDim/2-xint)+(detectorXDim-yint)+(zfocus)];

	int*intersectionsy[(detectorXDim/2-xint)+(detectorXDim-yint)+(zfocus)];

	int*intersectionsz[(detectorXDim/2-xint)+(detectorXDim-yint)+(zfocus)];

In the end I wanted to be sure that this is wrong, and how can I have one thread be responsible for three arrays, which length depend on the thread’s index

2

You can’t store floats in integer variables…

int slopeVectorx=xfocus-idxxbinsize+0.5;
int slopeVectory=yfocus-idyybinsize+0.5;

3

Thank you for pointing that out. Any thoughts on the arrays problem?

4

Anyone?

Briefly, my question is: can I have a thread allocate one(or more) vectors(which length depends on the threadIdx) in its local memory and have that thread work on that vector? Basically, as if each thread was working like a CPU