two dimension shared memory access problem

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

I am working on image construction for CT. My problem is that I am trying to simulate a a grid(detector) as 2 dimensional array. The problem that I have; the computed value in the grid is filled in pairs behavior. I mean by pairs that program makes the value in zdetector[0][0] equal to any zdetector[threadIdx.x][threadIdx.y] in the same block, which is incorrect because every value in the array has a different value. I added down here a snippet of the code that make this computing:

const int IdX = threadIdx.x+blockIdx.xBLOCKDIM;
const int IdZ = threadIdx.y+blockIdx.yBLOCKDIM;

const int  IdXlast = 16+blockIdx.x*(BLOCKDIM);
const int IdZlast = 16 + blockIdx.y*(BLOCKDIM);
const float xF = phantomParams.srcobjdist * sinf(0.0f);
const float yF = phantomParams.srcdetdist * cosf(0.0f);

__shared__ float xDetector[detectorSizeX];

__shared__ float zDetector[detectorSizeX][detectorSizeZ];
//calculate the absolute voxel coordinates

#define ABSOLUTEIMAGEX (phantomParams.voxelsphantomParams.voxelpitch0.5f)
#define ABSOLUTEIMAGEZ (phantomParams.voxelsphantomParams.voxelpitch0.5f)
#define ABSOLUTEIMAGEY (phantomParams.voxelsphantomParams.voxelpitch0.5f)
#define ABSOLUTEDETX (((phantomParams.dimx * 0.5f) - phantomParams.detoffsetu)* phantomParams.pixelpitch)
#define ABSOLUTEDETZ (((phantomParams.dimz0.5f )- phantomParams.detoffsetv) phantomParams.pixelpitch)
#define yDet0 (-phantomParams.detObjdist)

/*const float absoluteImageX = phantomParams.voxels*phantomParams.voxelpitch*0.5f;//They are the same absoluteImageY,Z 
const float absoluteDetX = ((phantomParams.dimx * 0.5f) - phantomParams.detoffsetu)* phantomParams.pixelpitch;
const float absoluteDetZ = ((phantomParams.dimz*0.5f )- phantomParams.detoffsetv)* phantomParams.pixelpitch;*/

//intersection computing

const float xDet0 = IdX * phantomParams.pixelpitch - ABSOLUTEDETX;
const float xDet = xDet0* cosf(0.0f)  + yDet0*sinf(0.0f);// This will be chnaged to angles
const float yDet = yDet0* cosf(0.0f) - sinf(0.0f)*xDet0;
//This for extra column 
const float xDet01 = IdXlast * phantomParams.pixelpitch - ABSOLUTEDETX;
const float xDet1 = xDet01 * cosf(0.0f)  +yDet0*sinf(0.0f);
const float yDet1 = yDet0* cosf(0.0f) - sinf(0.0f)*xDet01;
const float zDet0 = IdZ* phantomParams.pixelpitch - ABSOLUTEDETZ;
//This for the extra Z
const float zDet01 = IdZlast*phantomParams.pixelpitch - ABSOLUTEDETZ;


zDetector[threadIdx.x][threadIdx.y] = -yDet*(0.0-zDet0)/(yF - yDet)+zDet0;

I hope that anybody can help

2

You might want to edit your post and put that code in a code block. Also, it will help if you’re more explicit with your problem… it’s hard to understand what you’re asking as it is worded now. Try to simplify it and mention how and with what you’re calling your kernel, etc.