急急!CPU函数转到GPU运行输出的数据不一样

首先给大家看看我的CPU写的:
for (int i=0; i<Depth;i++)
for (int j=0; j<Depth;j++)
for (int k=0; k<Depth;k++)
{
int index = iDepthDepth+j*Depth+k;
if (voxelStatus[index]==0)
{
for(int jz=0; jz<BoxWidth; jz++)
for(int jy=0; jy<BoxWidth; jy++)
for(int jx=0; jx<BoxWidth; jx++)
{
int x = k-(BoxWidth/2)+jx;
int y = j-(BoxWidth/2)+jy;
int z = i-(BoxWidth/2)+jz;

int boxIndex = zDepthDepth+y*Depth+x;

if(BoxVoxelStatus[boxIndex] == 0)
{

BoxVoxelStatus[boxIndex] = 1;///存在问题 其它线程并不一定知道它已经标记为0;
}
}
}
}

下面是我转换成GPU后的:
这样分配的线程:
dim3 dimGrid(Depth,Depth);
dim3 dimBlock(Depth,1);

index=(blockIdx.y*gridDim.x+blockIdx.x)*blockDim.x+threadIdx.x;

cudaVoxelStatus[index] = (s+6)/7-1 ;
//查找每个边界体素的包围盒
if(cudaVoxelStatus[index] == 0)
{
for(int jz=0; jz<BoxWidth; jz++)
for(int jy=0; jy<BoxWidth; jy++)
for(int jx=0; jx<BoxWidth; jx++)
{
int x = threadIdx.x-(BoxWidth/2)+jx;
int y = blockIdx.x-(BoxWidth/2)+jy;
int z = blockIdx.y-(BoxWidth/2)+jz;

//int boxIndex = z*(Depth+1)(Depth+1)+y(Depth+1)+x;
int boxIndex = zDepthDepth+y*Depth+x;

if(cudaBoxVoxelStatus[boxIndex] == 0)
{

cudaBoxVoxelStatus[boxIndex] = 1;
}

没仔细看,但是估计线程没有同步造成的~

顶一下,的确是没有同步的问题,但是做过尝试性同步没有解决问题,还等高手来解~

[

似乎没有发现任何问题。。。除了比CPU版本的多了一句话。
对了。你发下调用前后的上下文?
请高手前来解答!

[ 本帖最后由 悠闲的小猫 于 2011-3-27 15:10 编辑 ]

恩 针对您的回复我把这个kernel的逻辑说一下:首先这个s是前面一个语句的计算和,这个和是没问题的,因为cudaVoxelStatus[index] 这个值的输出是正确的。整个的意思是先判断cudaVoxelStatus[index] 是否为0(实际是标注的一个小立方体),如果它为0,则在以这个小立方体为中心立方体在其周围Boxwidth范围内找到所有的立方体并把它们标注为1(初始为0)即cudaBoxVoxelStatus[boxIndex]=1。这是这个kernel的整个含义啊。请您给看看
[

[

那就奇怪了。。还是没有发现任何问题。。。。太奇怪了。。。

看了您的后面的部分:

[
调用完BoxIndexKernel, 再用BoxIndexKernel_CleanUp<<<1,1>>>()。

[ 本帖最后由 悠闲的小猫 于 2011-3-27 18:27 编辑 ]

恩 我已经把判断完cudaBoxVoxelStatus[boxIndex] = 1后,接下来要算起对应空间的x y z值的kernel发给您短消息了,麻烦您看看,它们两个是连着的啊 是不是那里出错了啊!!谢谢
[

[
额。。我看到了。在我发上个回贴的同时你发了这个帖子。呵呵。您看下。

[
这里的我之所以是上面的 是因为 我是先沿着xz平面剖分 然后 再沿着y方向增加的。threadIdx.x; blockIdx.x;blockIdx.y; 如果您不介意 可以加您的QQ 给看看我的程序吗 谢谢先

[

QQ不在线。。如果这里是故意的。那就完全没有任何问题了。。。太奇怪了。。。

对了。您的意图是:
根据voxelStatus的(x0,y0,z0)处的值,来设置BoxVoxelStatus(x0-W/2,y0-W/2,z0-W/2)到BoxVoxelStatus(x0+W/2,y0+W/2,z0+W/2)的值,最后再取出BoxVoxelStatus里这些值为1的点的坐标吗?

希望我没有理解错。。

[
恩 还是谢谢了啊 不知道还有没有人来给解答啊

[
恩 恩 是啊是啊 就是这个逻辑啊 输出的是BoxVoxelStatus为1的对应的 x y z值啊 就是这些值和cpu计算出的不一样啊

[

看了看。。真的完全没问题。。太奇怪了。。。
对了。你看看你这两处:

int x = threadIdx.x-(BoxWidth/2)+jx;
int y = blockIdx.x-(BoxWidth/2)+jy;
int z = blockIdx.y-(BoxWidth/2)+jz;

cudaBoxIndex_X[tmp] = blockIdx.y;//iy
cudaBoxIndex_Y[tmp] = blockIdx.x;//iz
cudaBoxIndex_Z[tmp] = threadIdx.x;//ix
你确定你的确是故意这样而且完全认为正确吗?

[ 本帖最后由 悠闲的小猫 于 2011-3-27 18:46 编辑 ]

[
cuda中 应该是按threadIdx.x blockIdx.x blockIdx.y每个递增的吧 jx jy jz 也是对应的先线程 然后再块的

[
cudaBoxVoxelStatus 是按这个下标找的啊z y x : int boxIndex = zDepthDepth+y*Depth+x;

[

我给你举个特例:
假设BoxWidth=2, 同时假设在threadIdx.x为2,blockIdx.x为3,blockIdx.y为4的时候,
才可以满足voxelStatus[index]==0, 那么:

你的第一次kernel标记了
BoxVoxelStatus中对应如下的点:
(1,2,3) (2,2,3) (1,3,3) (2,3,3)
(1,2,4) (2,2,4) (1,3,4) (2,3,4)
一共8个。

然后在你的第二个kernel里面,
你却会得到什么呢?根据你的写法, 是:
cudaBoxIndex_X,cudaBoxIndex_Y,cudaBoxIndex_Z)
偏移a 3,2,1
偏移b 3,2,2
偏移c 3,3,1
偏移d 3,3,2
偏移e 4,2,1
偏移f 4,2,2
偏移g 4,3,1
偏移h 4,3,2
(a-h是0-7里面的8个数的某种可能的排列)

所以您看。。您第二个kernel取出的结果和您第一个kernel得到的结果,是不是有点不对?

跟据您的写法,您无法得到本回帖的第一段中提出的和第二段中您所确定的目的。

[ 本帖最后由 悠闲的小猫 于 2011-3-27 19:53 编辑 ]

[
但是 我是这样int boxIndex = zDepthDepth+y*Depth+x; 所以在第一段中也是先存的z 即 您举例中的blockIdx.y 。所以在第二个kernel中 应该是blockIdx.y 在前的啊,这个可以由之前的cpu值对应起来的 。另外,我设计的是以voxelStatus==0;的为中心立方体,所以boxwidth按奇数算的。

[

好吧。。。我没有看懂。。我得出的结论只是从您的代码和grid shape(Depth,Depth,1), block shape (Depth,1,1)中得出的。请原谅我的没有理解。
还请高人指点!