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

加速计算 CUDA
1

首先给大家看看我的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;
}

2

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

3

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

4

[

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

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

5

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

6

[

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

7

看了您的后面的部分:

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

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

8

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

9

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

10

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

11

[

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

12

对了。您的意图是:
根据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的点的坐标吗？

希望我没有理解错。。

13

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

14

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

15

[

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

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 编辑 ]

16

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

17

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

18

[

我给你举个特例:
假设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 编辑 ]

19

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

20

[

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