列出帖子

该操作将允许你查看该会员所有的帖子,注意你只能看到你有权限看到的板块的帖子。


显示所有帖子 - 奈奈

页: [1] 2 3 ... 5
1
CUDA / Re: 使用cuda 非32位数据的共享内存使用方法
« 于: 十一月 13, 2020, 10:08:47 am »

感谢奈奈同学工作之余, 百忙之中的回复.

楼主的代码(double)实际上在profiler下, 并不会报告bank conflict了. 该行为从计算能力5.0+就可以观察到(包括楼主的pascal的显卡), 但没有在计算能力7.5(图灵)+上测试过.

你可以和楼主用一张6.1的卡在profiler下观察到这点.

我们(你, 我, W)当初的笔记和讨论实际上并不非常适用于楼主今日的话题. 楼主主要是用于连续顺序访问的. 这种情况的确不会, 但是我不知道原因, 手册也没有做出解释.

我个人的猜测是shared memory和相关的单元在处理的时候, 使用了一个很小的缓冲区, 将32个8B里面的前面的16个8B组合了一下, 然后后面的16B个8B组合了一下, 有点像按照half warp(本例)拆分了访问一样. 或者对于一次性的非4B的访问(指真正的访存指令时候的超过8B, 例如用double和float2都可以激活这种效果), 例如8B的访问的时候, 按照half warp进行, 而对于标准的4B访问, 按照warp进行.

我建议奈奈同学分别使用4B, 8B, 16B访问来测试profiler的bank conflict报告情况. (即: float, float2, float4, 或者float, double, double2), 这3种类型(以及他们对应的整数类型)均会导致编译器生成LDS, LDS.64, LDS.128这种指令. 丹丹可以猜测和验证是否4B是按照warp来的, 8B按照half warp来的, 而16B按照1/4 warp来的. 无论验证或者否定, 都将会是一个有趣的发现, 不是吗?

看到你给出的链接里面的手抄笔记的熟悉的字体, 感慨万千, 虽然当年参与讨论的人已经有一个不在了, 但是该网站保存了足够美好的回忆.

S写于祖国的边疆
好久不见。
如果如你所说“4B是按照warp来的, 8B按照half warp来的, 而16B按照1/4 warp来访问”那么楼主的例子的确不会出现conflict,因为16个线程内的确没有不同线程访问同一个bank。是我轻率了,我还以为现在的卡都是warp访问。看来理解得还不够透彻才导致了浅薄的推导。虽然我的卡是7.5的,但我也会抽空验证一下你的猜测。 
另外,再次感叹日月如梭,看下时间都过了三年,向你们请教的日子让我成长很多,包括现在重新学习CUDA,回看当时的记录仍旧有新的感慨和领悟。现在自己也会遇到问题,会查会问,无果时会自己推论、验证、思考反复循环,过程艰难又快乐。依旧感谢你们当初的指导。

2
CUDA / Re: 使用cuda 非32位数据的共享内存使用方法
« 于: 十一月 12, 2020, 06:38:39 pm »
共享内存的数据类型通常为float,因为某种特殊需求我需要读取double的数据
__shared__ double a[256];
软硬件环境:计算力6.1的pascal 架构显卡 cuda 10.1
问题:
1. 对于一个block大小为256的块,每个线程直接访问对应数组的元素(如下),会出现bank conflict吗?
     a[threadIdx.x]=1.0;
2.如果会出现bank conflict,应该怎么处理?
3.如果不会,按照手册里说maxwell架构之后的,共享内存的bank size 都是4字节,他是通过什么方式避免了bank conflict?
4.不同架构下,处理__shared__ double a[256];这种数据访问的处理方式有什么不同?
个人愚见:来源https://blog.csdn.net/wd1603926823/article/details/78326570 若有不对请指出,相互学习
1、现在的卡一般都是warp访问,32个banks,每个bank内32bit,所以我假设你的卡也是如此。那么会出现conflict,因为数据是64bit占2个banks。也就是a[0]放在bank0和bank1简称b0、b1;a[1]--b2、b3;依次类推,a[15]在b30、b31;a[16]在b0、b1,其实类推到此时就可以看到线程0和线程16即不同的线程访问了同一个bank0,拟继续类推会发现更严重的conflict。你的线程0、16、32、64、128即5个不同线程都访问了bank0。按照conflict的定义:不同线程访问同一个bank会出现conflict,所以此处存在conflict。
2、怎么处理,我刚刚推了一下,可以通过补位的方式,至于补多少,你可以看完这个网址上的介绍自己推。假设每16个数补4个无意义的数extend,那么也就是a[0]放在bank0和bank1简称b0、b1;a[1]--b2、b3;依次类推,a[15]在b30、b31;extend[0]在b0,b1;extend[1]在b2、b3以此类推;a[16]在b8、b9;a[17]在b10、b11...可以看到至少目前没有发生conflict,后面你自己慢慢推吧。其实你自己像这个网址中这样用一张草稿推更简单明了。
3、4、前面的推明白了,这两点也知道了

3
我不是很懂你这里所做的计算,不过还是给出一些看法。
(1)我觉得最好不要在一个核函数内使用线程块之间同步的技巧。我觉得这样做容易出错,而且代码难以理解。
(2)你说你试了用一个block的方式。我觉得如果你可以同时处理很多图片的话,这样的方案应该很好。这样,你可以用一个block处理一个图片,每个图片的处理相互独立。我觉得这样做应该是最高效的。
谢谢,的确在GPU端对不同blocks的threads进行同步这是不对的,会有问题。我试用了1)在CPU端多次启动一个核函数来隐形做所有blocks所有threads的同步 2)只设置一个block启动一次核函数,这个block内的所有线程合作完成计算。这两种方法测试了很多图片正确,但效率对我来说太慢了,可能我要改算法思路,这种思路下两种编写方式已经到瓶颈了,我也想不出更优的方案了。

4
我今天还试了一种方式,总共1个block,在GPU端做同步不需cpu端控制迭代。结果与CPU完全一致,缺点就是真的很慢
程序代码: [选择]
__global__ void dtgpu_5st(size_t width,size_t height,uchar* dev_dtimg)
{
int id = threadIdx.x;//0~1023

///////////////////临时距离变换结果初始化为0
__shared__ uchar localdtimg[pixelsnum];
for(int inde=id;inde<pixelsnum;inde += blockDim.x)
{
localdtimg[inde]=0;//临时距离变换结果初始化
}

//////////////////计算全图距离值为1,默认不是黑图
for(int index=id;index<pixelsnum;index += blockDim.x)
{
int x=index % width;
int y=index / width;
//二维纹理不需要user处理图像边界情况
uchar center= tex2D(texbw,x,y);//本线程中心点
if(center!=0)
{
uchar up,down,left,right,upleft,upright,downright,downleft;
up= tex2D(texbw,x,y-1);
down= tex2D(texbw,x,y+1);
left= tex2D(texbw,x-1,y);
right= tex2D(texbw,x+1,y);
upleft= tex2D(texbw,x-1,y-1);
upright= tex2D(texbw,x+1,y-1);
downright= tex2D(texbw,x-1,y+1);
downleft= tex2D(texbw,x+1,y+1);
int neighbor8=up*down*left*right*upleft*upright*downright*downleft;
if(neighbor8==0)
{
localdtimg[index]=1;//第一次迭代即石头最外面一圈
}
}
}

__shared__ bool nexttime;
__shared__ size_t dtvalue;
if(id==0)
{
nexttime=true;//下一次是否需要迭代的标志
dtvalue = 1;// dtvalue=1,2,3,4,...max
}
__syncthreads();//第一次迭代即石头最外面一圈完成
//以上都没有时序问题

//////////////////计算全图距离值为2,3,4,5,...,maxdtvalue
while(nexttime)
{
//开始此次迭代,dtvalue=2开始
if(id==0)
{
nexttime=false;//下一次是否需要迭代的标志
dtvalue += 1;
}
//把大图上一次的结果临时保存,以防时序导致错乱
__shared__ uchar lastdttmp[pixelsnum];
for(int inde=id;inde<pixelsnum;inde += blockDim.x)
{
lastdttmp[inde]=localdtimg[inde];
}
__syncthreads();

//计算大图这一次的结果
for(int index=id;index<pixelsnum;index += blockDim.x)
{
//计算大图的一小部分
int x=index % width;
int y=index / width;
if((x>0 && x<width-1)&&(y>0 && y<height-1))
{
uchar lastdis=lastdttmp[index];//上一轮迭代时此线程计算的距离值
uchar center= tex2D(texbw,x,y);//本线程中心点
if(center!=0 && lastdis==0)
{
//有活儿干的线程
uchar up,down,left,right,upleft,upright,downright,downleft;
up=lastdttmp[index-width];
down=lastdttmp[index+width];
left=lastdttmp[index-1];
right=lastdttmp[index+1];
upleft=lastdttmp[index-width-1];
upright=lastdttmp[index-width+1];
downright=lastdttmp[index+width+1];
downleft=lastdttmp[index+width-1];

uchar min=255;
if(up!=0 && up<min)
{
min=up;
}
if(down!=0 && down<min)
{
min=down;
}
if(left!=0 && left<min)
{
min=left;
}
if(right!=0 && right<min)
{
min=right;
}
if(upleft!=0 && upleft<min)
{
min=upleft;
}
if(upright!=0 && upright<min)
{
min=upright;
}
if(downright!=0 && downright<min)
{
min=downright;
}
if(downleft!=0 && downleft<min)
{
min=downleft;
}
nexttime=true;//所有干活线程去写这个全局变量吗??怎么控制只有1个线程去写
localdtimg[index]=min+1;
}
}
//此次迭代只计算了全图的一部分,继续计算全图的其它部分
}
__syncthreads();
//此次迭代结束,准备开始下一次迭代
}
//所有迭代完毕,距离变换计算完毕


//将临时距离变换结果写回全局距离变换结果中
for(int index=id;index<pixelsnum;index += blockDim.x)
{
dev_dtimg[index]=localdtimg[index];
}
//__syncthreads();//不用多此一举,核函数结束也有一次同步

//核函数结束了
}

5
我现在使用的是CPU端控制每次迭代的核函数,那么每次核函数运行完毕,回到CPU端时相当于对所有线程进行了同步操作,然后再开始下一次迭代:
程序代码: [选择]
//距离变换结果使用全局内存 读写
uchar *dev_dtimg;
cudaMalloc((void**)&dev_dtimg, imgsize);

//下一次是否需要迭代的标志
int *dev_next;
cudaMalloc((void**)&dev_next,sizeof(int));

//计算最外围的1
dtgpu_1st<<<blocks,threads>>>(cols,rows,dev_dtimg);
cudaDeviceSynchronize();

//计算2到n,CPU控制核函数计算完毕即同步
int maxdtvalue=0;
int nextresult[1]={1};
while(nextresult[0]==1)
{
nextresult[0]=0;//是否需要继续迭代的标志
cudaMemcpy(dev_next, nextresult, sizeof(int), cudaMemcpyHostToDevice);
dtgpu_2<<<blocks,threads>>>(cols,rows,dev_dtimg,dev_next);
cudaMemcpy(nextresult,dev_next, sizeof(int), cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();

maxdtvalue++;//迭代次数就是最大距离值
}
//gpu 端最终结果
cudaMemcpy(dtimg.data,dev_dtimg,imgsize,cudaMemcpyDeviceToHost);

//CPU端结果对比
cv::Mat dist_min=cv::Mat::zeros(bwimg.rows,bwimg.cols,CV_8UC1);
int maxvalue=0;
minDT(bwimg,dist_min,maxvalue);
bool cpugpumatch=true;
for(int r=1;r!=rows-1;r++)
{
for(int c=1;c!=cols-1;c++)
{
int dt_cpu=dist_min.ptr<uchar>(r)[c];
int dt_gpu=dtimg.ptr<uchar>(r)[c];
if(dt_cpu!=dt_gpu)
{
cpugpumatch=false;
break;
}
}
}
if(cpugpumatch)
{
std::cout<<"image: "<<i<<" --cpu and gpu results match ! maxdtvalue: "<<maxdtvalue<<std::endl;
}
else
{
std::cout<<"image: "<<i<<"--Error:cpu and gpu results dismatch !!!!! maxdtvalue: "<<maxdtvalue<<std::endl;
}

程序代码: [选择]
__global__ void dtgpu_2(size_t width,size_t height,uchar* dev_dtimg,int *next)
{
int x = threadIdx.x + blockIdx.x * blockDim.x;//这个是width的坐标(横坐标)
int y = threadIdx.y + blockIdx.y * blockDim.y;//这个是height的坐标(纵坐标)
int offset = x + y * blockDim.x * gridDim.x;

uchar lastdis=dev_dtimg[offset];//上一轮迭代时此线程计算的距离值
uchar center,up,down,left,right,upleft,upright,downright,downleft;
if((x>0 && x<width-1)&&(y>0 && y<height-1))
{
up=dev_dtimg[offset-width];
down=dev_dtimg[offset+width];
left=dev_dtimg[offset-1];
right=dev_dtimg[offset+1];
upleft=dev_dtimg[offset-width-1];
upright=dev_dtimg[offset-width+1];
downright=dev_dtimg[offset+width+1];
downleft=dev_dtimg[offset+width-1];

center= tex2D(texbw,x,y);//本线程中心点
if(center!=0 && lastdis==0)
{
//有活儿干的线程
uchar min=255;
if(up!=0 && up<min)
{
min=up;
}
if(down!=0 && down<min)
{
min=down;
}
if(left!=0 && left<min)
{
min=left;
}
if(right!=0 && right<min)
{
min=right;
}
if(upleft!=0 && upleft<min)
{
min=upleft;
}
if(upright!=0 && upright<min)
{
min=upright;
}
if(downright!=0 && downright<min)
{
min=downright;
}
if(downleft!=0 && downleft<min)
{
min=downleft;
}
next[0]=1;//所有计算线程去写这个全局变量吗??怎么控制只有1个线程去写
lastdis=min+1;
}
}
//为下一次迭代做准备
//将临时dt值写到全局距离变换结果中
dev_dtimg[offset]=lastdis;
}

目前测试了很多图片 CPU与GPU结果一致,但我很担心效率

6
当我将核函数改成下面这样,居然运行结果和CPU版本一致。明明对比上一个版本,我没做什么实质性改动,上一个核函数每次迭代循环是直接从全局变量中得到8邻域并开始计算,我改成了先将全局变量中得到的8邻域保存到这个线程的私有寄存器中,然后对私有寄存器中保存的8个数开始计算。

//第3个版本 将上一些迭代时的8邻域结果从全局保存到每个线程的私有寄存器中,然后每个线程访问寄存器中的数据进行计算
__global__ void dtgpu_min(size_t width,size_t height,uchar* dev_dtimg)
{
   //存储每个block内每次迭代时,每个thread计算的临时dt值
   //每次迭代完时,所有blocks内的所有threads一起写给最终的dev_dtimg
   uchar tmp_dt=0;//初始化为0

   int x = threadIdx.x + blockIdx.x * blockDim.x;//这个是width的坐标(横坐标)
   int y = threadIdx.y + blockIdx.y * blockDim.y;//这个是height的坐标(纵坐标)
   int offset = x + y * blockDim.x * gridDim.x;

   /////第一次迭代 默认肯定有石头 不会是黑图
   uchar center,up,down,left,right,upleft,upright,downright,downleft;
   //二维纹理不需要user处理图像边界情况
   center= tex2D(texbw,x,y);//本线程中心点
   up= tex2D(texbw,x,y-1);
   down= tex2D(texbw,x,y+1);
   left= tex2D(texbw,x-1,y);
   right= tex2D(texbw,x+1,y);
   upleft= tex2D(texbw,x-1,y-1);
   upright= tex2D(texbw,x+1,y-1);
   downright= tex2D(texbw,x-1,y+1);
   downleft= tex2D(texbw,x+1,y+1);
   if(center!=0)
   {
      int neighbor8=up*down*left*right*upleft*upright*downright*downleft;
      if(neighbor8==0)
      {
         tmp_dt=1;
         dev_dtimg[offset]=1;//第一次迭代即石头最外面一圈
          __threadfence();
      }
   }

   bool has255=true;//默认一张图肯定未处理之前是有255的
   while(has255)
   {
      has255=false;//假设此次迭代后就没有255了
      //////第n次迭代
      //所有线程 将周围8邻域存到私有寄存器内
      if((x>0 && x<width-1)&&(y>0 && y<height-1))
      {
         up=dev_dtimg[offset-width];
         down=dev_dtimg[offset+width];
         left=dev_dtimg[offset-1];
         right=dev_dtimg[offset+1];
         upleft=dev_dtimg[offset-width-1];
         upright=dev_dtimg[offset-width+1];
         downright=dev_dtimg[offset+width+1];
         downleft=dev_dtimg[offset+width-1];

         uchar lastdis=tmp_dt;//上一轮迭代时此线程计算的距离值
         if(center!=0 && lastdis==0)
         {
            //有活儿干的线程
            uchar min=255;
            if(up!=0 && up<min)
            {
               min=up;
            }
            if(down!=0 && down<min)
            {
               min=down;
            }
            if(left!=0 && left<min)
            {
               min=left;
            }
            if(right!=0 && right<min)
            {
               min=right;
            }
            if(upleft!=0 && upleft<min)
            {
               min=upleft;
            }
            if(upright!=0 && upright<min)
            {
               min=upright;
            }
            if(downright!=0 && downright<min)
            {
               min=downright;
            }
            if(downleft!=0 && downleft<min)
            {
               min=downleft;
            }
            has255=true;//迭代过程中发现还是有255,说明下次还要迭代
            tmp_dt=min+1;
         }
      }
      //为下一次迭代做准备
      //将临时dt值写到全局距离变换结果中
      dev_dtimg[offset]=tmp_dt;
      __threadfence();
      //开始下一次迭代
   }
}

GPU version...
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 2 2 2 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 3 3 3 3 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 3 3 4 4 3 3 3 3 2 2 2 2 2 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 2 2 2 2 2 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 5 5 5 5 4 4 4 4 4 3 3 3 3 3 3 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 4 4 5 5 5 6 6 6 6 5 5 5 5 5 4 4 4 4 4 4 3 3 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 5 5 5 5 6 6 6 7 7 6 6 6 6 6 5 5 5 5 5 5 4 4 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 6 6 6 6 7 7 7 7 7 7 7 7 6 6 6 6 6 6 5 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 4 5 5 6 6 7 7 7 7 8 8 8 8 8 8 7 7 7 7 6 6 5 5 4 3 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 2 2 3 3 4 4 5 5 5 6 6 7 7 8 8 8 8 9 9 9 9 8 8 7 6 6 6 5 5 4 4 3 2 2 1 1 0 0 0 0 0 0 0
0 0 0 1 1 2 3 3 4 4 5 5 6 6 6 7 7 8 8 8 8 8 8 8 8 8 8 8 7 6 5 5 5 4 4 3 3 2 1 1 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 4 4 5 5 6 6 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 6 5 4 4 4 3 3 2 2 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 5 5 6 6 7 7 7 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 5 4 3 3 3 2 2 1 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 4 5 5 6 6 6 6 6 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 4 3 2 2 2 1 1 0 0 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 3 4 4 5 5 6 6 5 5 5 4 4 4 4 4 4 4 4 4 4 4 4 4 5 4 3 2 1 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 1 1 2 2 3 3 4 4 5 5 5 5 4 4 4 3 3 3 3 3 3 3 3 3 3 3 4 4 4 3 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 5 4 4 4 3 3 3 2 2 2 2 2 2 2 2 2 3 3 4 4 3 2 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 4 4 4 3 3 3 2 2 2 1 1 1 1 1 1 1 2 2 3 4 4 3 2 2 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 2 2 2 1 1 1 0 0 0 0 0 1 1 2 3 4 4 3 3 2 2 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 2 3 4 4 4 3 3 2 2 2 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 1 1 2 3 4 5 4 4 3 3 3 3 2 2 1 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 1 1 1 1 1 1 2 2 2 3 4 5 5 4 4 4 4 3 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 2 2 3 3 3 4 5 5 5 5 5 4 4 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 3 3 3 3 4 4 4 5 6 6 6 5 5 4 3 2 1 1 1 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 3 4 4 4 4 4 4 5 5 5 6 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 4 5 5 5 5 5 5 6 6 6 7 7 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 5 6 6 6 6 6 6 7 7 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 6 6 7 7 7 7 7 7 7 6 5 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 5 5 6 6 7 7 6 6 6 6 6 5 4 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 4 4 5 5 6 6 6 6 5 5 5 5 5 4 3 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 5 5 5 5 5 5 4 4 4 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 3 3 3 3 3 2 2 2 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 2 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0


CPU version...
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 2 2 2 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 3 3 3 3 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 3 3 4 4 3 3 3 3 2 2 2 2 2 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 2 2 2 2 2 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 5 5 5 5 4 4 4 4 4 3 3 3 3 3 3 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 4 4 5 5 5 6 6 6 6 5 5 5 5 5 4 4 4 4 4 4 3 3 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 5 5 5 5 6 6 6 7 7 6 6 6 6 6 5 5 5 5 5 5 4 4 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 6 6 6 6 7 7 7 7 7 7 7 7 6 6 6 6 6 6 5 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 4 5 5 6 6 7 7 7 7 8 8 8 8 8 8 7 7 7 7 6 6 5 5 4 3 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 2 2 3 3 4 4 5 5 5 6 6 7 7 8 8 8 8 9 9 9 9 8 8 7 6 6 6 5 5 4 4 3 2 2 1 1 0 0 0 0 0 0 0
0 0 0 1 1 2 3 3 4 4 5 5 6 6 6 7 7 8 8 8 8 8 8 8 8 8 8 8 7 6 5 5 5 4 4 3 3 2 1 1 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 4 4 5 5 6 6 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 6 5 4 4 4 3 3 2 2 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 5 5 6 6 7 7 7 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 5 4 3 3 3 2 2 1 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 4 5 5 6 6 6 6 6 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 4 3 2 2 2 1 1 0 0 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 3 4 4 5 5 6 6 5 5 5 4 4 4 4 4 4 4 4 4 4 4 4 4 5 4 3 2 1 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 1 1 2 2 3 3 4 4 5 5 5 5 4 4 4 3 3 3 3 3 3 3 3 3 3 3 4 4 4 3 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 5 4 4 4 3 3 3 2 2 2 2 2 2 2 2 2 3 3 4 4 3 2 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 4 4 4 3 3 3 2 2 2 1 1 1 1 1 1 1 2 2 3 4 4 3 2 2 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 2 2 2 1 1 1 0 0 0 0 0 1 1 2 3 4 4 3 3 2 2 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 2 3 4 4 4 3 3 2 2 2 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 1 1 2 3 4 5 4 4 3 3 3 3 2 2 1 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 1 1 1 1 1 1 2 2 2 3 4 5 5 4 4 4 4 3 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 2 2 3 3 3 4 5 5 5 5 5 4 4 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 3 3 3 3 4 4 4 5 6 6 6 5 5 4 3 2 1 1 1 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 3 4 4 4 4 4 4 5 5 5 6 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 4 5 5 5 5 5 5 6 6 6 7 7 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 5 6 6 6 6 6 6 7 7 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 6 6 7 7 7 7 7 7 7 6 5 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 5 5 6 6 7 7 6 6 6 6 6 5 4 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 4 4 5 5 6 6 6 6 5 5 5 5 5 4 3 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 5 5 5 5 5 5 4 4 4 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 3 3 3 3 3 2 2 2 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 2 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0

难道是将全局变量的8个数保存到私有寄存器中这个过程的时间冲掉了我之前说的有的线程跑得太快而导致数据出入问题??
我再测几组图试试,也许只是偶然现象,按道理应该和上一版一样与CPU版结果有几个数不一致的啊

这种方式是不行的,因为每次迭代结束没有对所有线程同步,有的线程跑得快,有的慢比如线程A去读全局dev_img中A右边位置B,本来B=3,而B跑得比较快,可能已经完成了B自身的写操作变成4,那么A读B时就读到了4!

7
CUDA / Re: CUDA在对分水岭算法改编出现问题
« 于: 十月 26, 2020, 03:15:09 pm »
楼主 你的GPU版的分水岭最终实现出来了吗?

8
当我将核函数改成下面这样,居然运行结果和CPU版本一致。明明对比上一个版本,我没做什么实质性改动,上一个核函数每次迭代循环是直接从全局变量中得到8邻域并开始计算,我改成了先将全局变量中得到的8邻域保存到这个线程的私有寄存器中,然后对私有寄存器中保存的8个数开始计算。

//第3个版本 将上一些迭代时的8邻域结果从全局保存到每个线程的私有寄存器中,然后每个线程访问寄存器中的数据进行计算
__global__ void dtgpu_min(size_t width,size_t height,uchar* dev_dtimg)
{
   //存储每个block内每次迭代时,每个thread计算的临时dt值
   //每次迭代完时,所有blocks内的所有threads一起写给最终的dev_dtimg
   uchar tmp_dt=0;//初始化为0

   int x = threadIdx.x + blockIdx.x * blockDim.x;//这个是width的坐标(横坐标)
   int y = threadIdx.y + blockIdx.y * blockDim.y;//这个是height的坐标(纵坐标)
   int offset = x + y * blockDim.x * gridDim.x;

   /////第一次迭代 默认肯定有石头 不会是黑图
   uchar center,up,down,left,right,upleft,upright,downright,downleft;
   //二维纹理不需要user处理图像边界情况
   center= tex2D(texbw,x,y);//本线程中心点
   up= tex2D(texbw,x,y-1);
   down= tex2D(texbw,x,y+1);
   left= tex2D(texbw,x-1,y);
   right= tex2D(texbw,x+1,y);
   upleft= tex2D(texbw,x-1,y-1);
   upright= tex2D(texbw,x+1,y-1);
   downright= tex2D(texbw,x-1,y+1);
   downleft= tex2D(texbw,x+1,y+1);
   if(center!=0)
   {
      int neighbor8=up*down*left*right*upleft*upright*downright*downleft;
      if(neighbor8==0)
      {
         tmp_dt=1;
         dev_dtimg[offset]=1;//第一次迭代即石头最外面一圈
          __threadfence();
      }
   }

   bool has255=true;//默认一张图肯定未处理之前是有255的
   while(has255)
   {
      has255=false;//假设此次迭代后就没有255了
      //////第n次迭代
      //所有线程 将周围8邻域存到私有寄存器内
      if((x>0 && x<width-1)&&(y>0 && y<height-1))
      {
         up=dev_dtimg[offset-width];
         down=dev_dtimg[offset+width];
         left=dev_dtimg[offset-1];
         right=dev_dtimg[offset+1];
         upleft=dev_dtimg[offset-width-1];
         upright=dev_dtimg[offset-width+1];
         downright=dev_dtimg[offset+width+1];
         downleft=dev_dtimg[offset+width-1];

         uchar lastdis=tmp_dt;//上一轮迭代时此线程计算的距离值
         if(center!=0 && lastdis==0)
         {
            //有活儿干的线程
            uchar min=255;
            if(up!=0 && up<min)
            {
               min=up;
            }
            if(down!=0 && down<min)
            {
               min=down;
            }
            if(left!=0 && left<min)
            {
               min=left;
            }
            if(right!=0 && right<min)
            {
               min=right;
            }
            if(upleft!=0 && upleft<min)
            {
               min=upleft;
            }
            if(upright!=0 && upright<min)
            {
               min=upright;
            }
            if(downright!=0 && downright<min)
            {
               min=downright;
            }
            if(downleft!=0 && downleft<min)
            {
               min=downleft;
            }
            has255=true;//迭代过程中发现还是有255,说明下次还要迭代
            tmp_dt=min+1;
         }
      }
      //为下一次迭代做准备
      //将临时dt值写到全局距离变换结果中
      dev_dtimg[offset]=tmp_dt;
      __threadfence();
      //开始下一次迭代
   }
}

GPU version...
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 2 2 2 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 3 3 3 3 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 3 3 4 4 3 3 3 3 2 2 2 2 2 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 2 2 2 2 2 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 5 5 5 5 4 4 4 4 4 3 3 3 3 3 3 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 4 4 5 5 5 6 6 6 6 5 5 5 5 5 4 4 4 4 4 4 3 3 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 5 5 5 5 6 6 6 7 7 6 6 6 6 6 5 5 5 5 5 5 4 4 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 6 6 6 6 7 7 7 7 7 7 7 7 6 6 6 6 6 6 5 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 4 5 5 6 6 7 7 7 7 8 8 8 8 8 8 7 7 7 7 6 6 5 5 4 3 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 2 2 3 3 4 4 5 5 5 6 6 7 7 8 8 8 8 9 9 9 9 8 8 7 6 6 6 5 5 4 4 3 2 2 1 1 0 0 0 0 0 0 0
0 0 0 1 1 2 3 3 4 4 5 5 6 6 6 7 7 8 8 8 8 8 8 8 8 8 8 8 7 6 5 5 5 4 4 3 3 2 1 1 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 4 4 5 5 6 6 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 6 5 4 4 4 3 3 2 2 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 5 5 6 6 7 7 7 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 5 4 3 3 3 2 2 1 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 4 5 5 6 6 6 6 6 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 4 3 2 2 2 1 1 0 0 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 3 4 4 5 5 6 6 5 5 5 4 4 4 4 4 4 4 4 4 4 4 4 4 5 4 3 2 1 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 1 1 2 2 3 3 4 4 5 5 5 5 4 4 4 3 3 3 3 3 3 3 3 3 3 3 4 4 4 3 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 5 4 4 4 3 3 3 2 2 2 2 2 2 2 2 2 3 3 4 4 3 2 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 4 4 4 3 3 3 2 2 2 1 1 1 1 1 1 1 2 2 3 4 4 3 2 2 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 2 2 2 1 1 1 0 0 0 0 0 1 1 2 3 4 4 3 3 2 2 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 2 3 4 4 4 3 3 2 2 2 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 1 1 2 3 4 5 4 4 3 3 3 3 2 2 1 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 1 1 1 1 1 1 2 2 2 3 4 5 5 4 4 4 4 3 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 2 2 3 3 3 4 5 5 5 5 5 4 4 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 3 3 3 3 4 4 4 5 6 6 6 5 5 4 3 2 1 1 1 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 3 4 4 4 4 4 4 5 5 5 6 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 4 5 5 5 5 5 5 6 6 6 7 7 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 5 6 6 6 6 6 6 7 7 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 6 6 7 7 7 7 7 7 7 6 5 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 5 5 6 6 7 7 6 6 6 6 6 5 4 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 4 4 5 5 6 6 6 6 5 5 5 5 5 4 3 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 5 5 5 5 5 5 4 4 4 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 3 3 3 3 3 2 2 2 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 2 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0


CPU version...
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 2 2 2 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 3 3 3 3 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 3 3 4 4 3 3 3 3 2 2 2 2 2 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 2 2 2 2 2 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 5 5 5 5 4 4 4 4 4 3 3 3 3 3 3 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 4 4 5 5 5 6 6 6 6 5 5 5 5 5 4 4 4 4 4 4 3 3 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 5 5 5 5 6 6 6 7 7 6 6 6 6 6 5 5 5 5 5 5 4 4 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 6 6 6 6 7 7 7 7 7 7 7 7 6 6 6 6 6 6 5 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 4 5 5 6 6 7 7 7 7 8 8 8 8 8 8 7 7 7 7 6 6 5 5 4 3 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 2 2 3 3 4 4 5 5 5 6 6 7 7 8 8 8 8 9 9 9 9 8 8 7 6 6 6 5 5 4 4 3 2 2 1 1 0 0 0 0 0 0 0
0 0 0 1 1 2 3 3 4 4 5 5 6 6 6 7 7 8 8 8 8 8 8 8 8 8 8 8 7 6 5 5 5 4 4 3 3 2 1 1 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 4 4 5 5 6 6 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 6 5 4 4 4 3 3 2 2 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 5 5 6 6 7 7 7 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 5 4 3 3 3 2 2 1 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 4 5 5 6 6 6 6 6 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 4 3 2 2 2 1 1 0 0 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 3 4 4 5 5 6 6 5 5 5 4 4 4 4 4 4 4 4 4 4 4 4 4 5 4 3 2 1 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 1 1 2 2 3 3 4 4 5 5 5 5 4 4 4 3 3 3 3 3 3 3 3 3 3 3 4 4 4 3 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 5 4 4 4 3 3 3 2 2 2 2 2 2 2 2 2 3 3 4 4 3 2 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 4 4 4 3 3 3 2 2 2 1 1 1 1 1 1 1 2 2 3 4 4 3 2 2 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 2 2 2 1 1 1 0 0 0 0 0 1 1 2 3 4 4 3 3 2 2 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 2 3 4 4 4 3 3 2 2 2 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 1 1 2 3 4 5 4 4 3 3 3 3 2 2 1 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 1 1 1 1 1 1 2 2 2 3 4 5 5 4 4 4 4 3 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 2 2 3 3 3 4 5 5 5 5 5 4 4 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 3 3 3 3 4 4 4 5 6 6 6 5 5 4 3 2 1 1 1 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 3 4 4 4 4 4 4 5 5 5 6 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 4 5 5 5 5 5 5 6 6 6 7 7 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 5 6 6 6 6 6 6 7 7 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 6 6 7 7 7 7 7 7 7 6 5 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 5 5 6 6 7 7 6 6 6 6 6 5 4 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 4 4 5 5 6 6 6 6 5 5 5 5 5 4 3 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 5 5 5 5 5 5 4 4 4 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 3 3 3 3 3 2 2 2 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 2 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0

难道是将全局变量的8个数保存到私有寄存器中这个过程的时间冲掉了我之前说的有的线程跑得太快而导致数据出入问题??
我再测几组图试试,也许只是偶然现象,按道理应该和上一版一样与CPU版结果有几个数不一致的啊

9
我在实现距离变换(非0点到最近的0的距离)的GPU版,对应CPU版结果如下:
CPU version...
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 2 2 2 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 3 3 3 3 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 3 3 4 4 3 3 3 3 2 2 2 2 2 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 2 2 2 2 2 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 5 5 5 5 4 4 4 4 4 3 3 3 3 3 3 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 4 4 5 5 5 6 6 6 6 5 5 5 5 5 4 4 4 4 4 4 3 3 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 5 5 5 5 6 6 6 7 7 6 6 6 6 6 5 5 5 5 5 5 4 4 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 6 6 6 6 7 7 7 7 7 7 7 7 6 6 6 6 6 6 5 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 4 5 5 6 6 7 7 7 7 8 8 8 8 8 8 7 7 7 7 6 6 5 5 4 3 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 2 2 3 3 4 4 5 5 5 6 6 7 7 8 8 8 8 9 9 9 9 8 8 7 6 6 6 5 5 4 4 3 2 2 1 1 0 0 0 0 0 0 0
0 0 0 1 1 2 3 3 4 4 5 5 6 6 6 7 7 8 8 8 8 8 8 8 8 8 8 8 7 6 5 5 5 4 4 3 3 2 1 1 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 4 4 5 5 6 6 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 6 5 4 4 4 3 3 2 2 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 5 5 6 6 7 7 7 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 5 4 3 3 3 2 2 1 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 4 5 5 6 6 6 6 6 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 4 3 2 2 2 1 1 0 0 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 3 4 4 5 5 6 6 5 5 5 4 4 4 4 4 4 4 4 4 4 4 4 4 5 4 3 2 1 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 1 1 2 2 3 3 4 4 5 5 5 5 4 4 4 3 3 3 3 3 3 3 3 3 3 3 4 4 4 3 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 5 4 4 4 3 3 3 2 2 2 2 2 2 2 2 2 3 3 4 4 3 2 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 4 4 4 3 3 3 2 2 2 1 1 1 1 1 1 1 2 2 3 4 4 3 2 2 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 2 2 2 1 1 1 0 0 0 0 0 1 1 2 3 4 4 3 3 2 2 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 2 3 4 4 4 3 3 2 2 2 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 1 1 2 3 4 5 4 4 3 3 3 3 2 2 1 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 1 1 1 1 1 1 2 2 2 3 4 5 5 4 4 4 4 3 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 2 2 3 3 3 4 5 5 5 5 5 4 4 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 3 3 3 3 4 4 4 5 6 6 6 5 5 4 3 2 1 1 1 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 3 4 4 4 4 4 4 5 5 5 6 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 4 5 5 5 5 5 5 6 6 6 7 7 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 5 6 6 6 6 6 6 7 7 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 6 6 7 7 7 7 7 7 7 6 5 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 5 5 6 6 7 7 6 6 6 6 6 5 4 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 4 4 5 5 6 6 6 6 5 5 5 5 5 4 3 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 5 5 5 5 5 5 4 4 4 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 3 3 3 3 3 2 2 2 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 2 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0

但是GPU版,最外面几层的计算我打印出来还正确:
第一次未进入while循环:
GPU version...
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 0 0 0 0 0 0 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0
0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
0 0 0 0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0
0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0
0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0 0 0
0 0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0
0 0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0
0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 1 1 1 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 1 1 1 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 1 1 1 0 0 0 1 1 1 0 0 0 0 0 0 0 0 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 0 0 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
这些1实际就代表了这个点的线程此次参数了主要计算

第一次后开始进入while循环:GPU version...
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 2 2 2 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 0 0 0 0 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 0 0 0 0 0 0 0 0 0 0 2 2 2 2 2 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 2 2 2 2 2 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 2 1 1 0 0 0 0 0 0 0
0 0 0 1 1 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 1 0 0 0 0 0 0 0 0
0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 2 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 2 1 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 2 2 1 1 0 0 0 0 0 0 0 0 0 0
0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 2 2 2 2 2 2 2 2 2 0 0 0 0 0 2 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 2 2 2 1 1 1 1 1 1 1 2 2 0 0 0 0 2 2 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 0 0 0 0 0 2 2 2 1 1 1 0 0 0 0 0 1 1 2 0 0 0 0 0 2 2 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 2 2 2 0 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 2 0 0 0 0 0 0 2 2 2 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 1 1 2 0 0 0 0 0 0 0 0 0 2 2 1 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 1 1 1 1 1 1 2 2 2 0 0 0 0 0 0 0 0 0 0 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 2 2 0 0 0 0 0 0 0 0 0 0 0 0 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 1 1 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 0 0 0 0 0 0 0 0 0 0 2 2 2 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 0 2 2 2 2 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
这些2实际就是此次循环时参数计算的线程即active的线程

但GPU最终结果如下:
max dt value:1 !max dt value:2 !GPU version...
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 2 2 2 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 3 3 3 3 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 3 3 4 4 3 3 3 3 2 2 2 2 2 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 2 2 2 2 2 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 5 5 5 5 4 4 4 4 4 3 3 3 3 3 3 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 4 4 5 5 5 6 6 6 6 5 5 5 5 5 4 4 4 4 4 4 3 3 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 5 5 5 5 6 6 6 7 7 6 6 6 6 6 5 5 5 5 5 5 4 4 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 6 6 6 6 7 8 7 7 7 7 7 7 6 6 6 6 6 6 5 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 4 5 5 6 6 7 7 7 7 8 9 8 8 8 8 7 7 7 7 6 6 5 5 4 3 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 1 2 2 3 3 4 4 5 5 5 6 6 7 7 9 9 9 9 9 9 9 9 8 8 7 6 6 6 5 5 4 4 3 2 2 1 1 0 0 0 0 0 0 0
0 0 0 1 1 2 3 3 4 4 5 5 6 6 6 7 7 8 8 8 8 8 8 8 8 8 8 8 7 6 5 5 5 4 4 3 3 2 1 1 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 4 4 5 5 6 6 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 6 5 4 4 4 3 3 2 2 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 5 5 6 6 7 7 7 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 5 4 3 3 3 2 2 1 1 0 0 0 0 0 0 0 0 0
0 0 1 2 2 3 3 4 4 5 5 6 6 6 6 6 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 4 3 2 2 2 1 1 0 0 0 0 0 0 0 0 0 0
0 0 1 1 2 2 3 3 4 4 5 5 6 6 5 5 5 4 4 4 4 4 4 4 4 4 4 4 4 4 5 4 3 2 1 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 1 1 2 2 3 3 4 4 5 5 5 5 4 4 4 3 3 3 3 3 3 3 3 3 3 3 4 4 4 3 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 1 1 2 2 3 3 4 4 5 4 4 4 3 3 3 2 2 2 2 2 2 2 2 2 3 3 4 4 3 2 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 1 1 2 2 3 3 4 4 4 3 3 3 2 2 2 1 1 1 1 1 1 1 2 2 3 4 4 3 2 2 1 1 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 2 2 2 1 1 1 0 0 0 0 0 1 1 2 3 4 4 3 3 2 2 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 2 3 4 4 4 3 3 2 2 2 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 0 0 0 0 0 0 0 0 1 1 1 2 3 4 5 4 4 3 3 3 3 2 2 1 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 1 1 1 1 1 1 2 2 2 3 4 5 5 4 4 4 4 3 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 2 2 2 3 3 3 4 5 5 5 5 5 4 4 3 2 2 1 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 3 3 3 3 4 4 4 5 6 6 6 5 5 4 3 2 1 1 1 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 3 4 4 4 4 4 4 5 5 5 6 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 3 4 4 4 5 5 5 5 5 5 6 6 6 7 7 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 4 4 5 5 5 6 6 6 6 6 6 7 7 7 6 6 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 4 4 4 5 5 6 6 7 7 7 7 7 7 7 6 5 5 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 2 3 3 3 4 4 5 5 6 6 7 7 6 6 6 6 6 5 4 4 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 2 2 2 3 3 4 4 5 5 6 6 6 6 5 5 5 5 5 4 3 3 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 3 3 4 4 5 5 5 5 5 5 4 4 4 4 4 3 2 2 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 4 4 4 4 4 4 4 4 3 3 3 3 3 2 1 1 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 3 3 3 3 3 3 3 3 3 3 2 2 2 2 2 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 2 2 2 3 2 2 2 2 2 2 2 2 1 1 1 1 1 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 2 2 2 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0

可以看到GPU的最后几个线程结果与CPU版有出入。GPU代码如下:
texture<uchar, 2,cudaReadModeElementType>  texbw;
#define THREADS_SIZE 34

__device__ uint count=0;
__global__ void dtgpu_min(size_t width,size_t height,uchar* dev_dtimg)
{
   uchar tmp_dt=0;//初始化为0
   int x = threadIdx.x + blockIdx.x * blockDim.x;//这个是width的坐标(横坐标)
   int y = threadIdx.y + blockIdx.y * blockDim.y;//这个是height的坐标(纵坐标)
   int offset = x + y * blockDim.x * gridDim.x;

   /////第一次迭代 默认肯定有石头 不会是黑图
   uchar center,up,down,left,right,upleft,upright,downright,downleft;
   //二维纹理不需要user处理图像边界情况
   center= tex2D(texbw,x,y);//本线程中心点
   up= tex2D(texbw,x,y-1);
   down= tex2D(texbw,x,y+1);
   left= tex2D(texbw,x-1,y);
   right= tex2D(texbw,x+1,y);
   upleft= tex2D(texbw,x-1,y-1);
   upright= tex2D(texbw,x+1,y-1);
   downright= tex2D(texbw,x-1,y+1);
   downleft= tex2D(texbw,x+1,y+1);
   if(center!=0)
   {
      int neighbor8=up*down*left*right*upleft*upright*downright*downleft;
      if(neighbor8==0)
      {
         tmp_dt=1;
         dev_dtimg[offset]=1;//第一次迭代即石头最外面一圈
          __threadfence();
      }
   }
   if(offset==0)
   {
      atomicAdd(&count, 1);//让一个线程去改变全图最大距离值
      printf("max dt value:%d !\n",int(count));
   }

   bool has255=true;//默认一张图肯定未处理之前是有255的
   while(has255)
   {
      has255=false;//假设此次迭代后就没有255了
      //////第n次迭代
      uchar lastdis=tmp_dt;//上一轮迭代时此线程计算的距离值
      if((center!=0 && lastdis==0)&&(x>0 && x<width-1)&&(y>0 && y<height-1))
      {
         //有活儿干的线程
         uchar min=255;
         up=dev_dtimg[offset-width];
         if(up!=0 && up<min)
         {
            min=up;
         }
         down=dev_dtimg[offset+width];
         if(down!=0 && down<min)
         {
            min=down;
         }
         left=dev_dtimg[offset-1];
         if(left!=0 && left<min)
         {
            min=left;
         }
         right=dev_dtimg[offset+1];
         if(right!=0 && right<min)
         {
            min=right;
         }
         upleft=dev_dtimg[offset-width-1];
         if(upleft!=0 && upleft<min)
         {
            min=upleft;
         }
         upright=dev_dtimg[offset-width+1];
         if(upright!=0 && upright<min)
         {
            min=upright;
         }
         downright=dev_dtimg[offset+width+1];
         if(downright!=0 && downright<min)
         {
            min=downright;
         }
         downleft=dev_dtimg[offset+width-1];
         if(downleft!=0 && downleft<min)
         {
            min=downleft;
         }
         has255=true;//迭代过程中发现还是有255,说明下次还要迭代
         tmp_dt=min+1;
      }
      //为下一次迭代做准备//将临时dt值写到全局距离变换结果中
      dev_dtimg[offset]=tmp_dt;
      __threadfence();
      //改变全局最大距离值
      if(offset==0)
      {
         int maxdis=atomicAdd(&count, 1);//让一个线程去改变全图最大距离值
         printf("max dt value:%d !\n",int(count));
      }
      //开始下一次迭代
   }
}

int main()
{
   cv::Mat testimg = cv::imread("/media/root/Ubuntu43/xrt/imgs/watershed-min/0.bmp",-1);
   cv::cvtColor(testimg,testimg,cv::COLOR_BGR2GRAY);
   //二值图像
   cv::Mat bwimg;
   cv::threshold(testimg,bwimg,100,255,cv::THRESH_BINARY);
   int rows=bwimg.rows;
   int cols=bwimg.cols;
   int imgsize=sizeof(uchar)*rows*cols;
   //距离变换结果
   cv::Mat dtimg=cv::Mat::zeros(rows,cols,CV_8UC1);
   //GPU版本--不扩展图像,2D纹理自动会处理图像边界问题
   dim3 threads(THREADS_SIZE, 1);
   dim3 blocks(2,rows);//68/64=2

   //二值图bwimg使用2D纹理内存 texbw 只读
   cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar>();
   cudaArray *dev_Src;
   cudaMallocArray(&dev_Src,&desc,cols,rows);
   cudaMemcpyToArray(dev_Src,0,0,bwimg.data,imgsize,cudaMemcpyHostToDevice);
   texbw.filterMode =cudaFilterModePoint;
   texbw.addressMode[0] = cudaAddressModeWrap;
   texbw.addressMode[1] = cudaAddressModeWrap;
   cudaBindTextureToArray(&texbw,dev_Src,&desc);

   //距离变换结果使用全局内存 读写
   uchar *dev_dtimg;
   cudaMalloc((void**)&dev_dtimg, imgsize);
   dtgpu_min<<<blocks,threads>>>(cols,rows,dev_dtimg);

   cudaMemcpy(dtimg.data,dev_dtimg,imgsize,cudaMemcpyDeviceToHost);
   cudaThreadSynchronize();
   
   cudaUnbindTexture(texbw);// 解绑和释放内存
   cudaFree(dev_Src);
   cudaFree(dev_dtimg);

   return 0;
}
(1)我想了一下为什么全局距离最大值应该输出9,而我输出2,是因为我这种写法有以下问题:比如第一次进入while,offset0这个线程给全局count加1,在这个动作期间,active的线程们可能已经计算到1或者2或者3了,别的有活儿干的线程并不会等待offset0完成将count=0加1变成1这个动作,是这样吗?
(2)为什么GPU结果最后几个线程,即输出结果为8和9那里与CPU版有点出入也是因为不是所有的6计算完毕再所有线程一起进入下一次while,可能某个线程T本该写7,而他周围3x3邻域内已有4个线程写了7,一个线程写了9,还有两个线程Tm、Tn正在写6,而这个线程T比Tm、Tn更快,所以它环顾自己的邻域发现最小的是7,所以输出自己的结果为8。也就是线程T没有等到Tm或Tn写完6。是这样吗?怎么让所有block的所有线程都等待呢,难道真只有核函数结束这一种办法?如果是这样,那我这个思路就得大改了。
(3)我想输出最大值9作为全局变量,参考了您说的手册中有一个例子是对一个数组(长度m*n)求和,每个block完成m个数的求和,然后每个block派出自己的thread0将本组求的一个临时和写到全局。总共n个block。通过全局变量count的计数atomicInc,其实就相当于标志每个thread0的快慢,最慢的一个计数完肯定是n-1,所以只要判断哪个thread0计数完是n-1,那么其代表的block就是最慢的。然后由这个最慢的block完成对大小为n的临时和的最终求和。我本来想利用这个例子,但好像不适合,因为我每次while内active的线程数不固定,我无法找到最慢的thread或block。

10
CUDA / Re: 求最大值的问题
« 于: 十月 22, 2020, 10:59:42 am »
请问第(3)点中所说的手册是指《CUDA C编程手册》还是《CUDA C专家手册》或《CUDA C编程指南》?我准备下载阅读。
已找到。是《CUDA C programming guide》!感叹[名词2]对这本书很熟悉,连里面的例子都记得这么清楚!向您学习。我正在学习这本书。

11
CUDA / Re: 求最大值的问题
« 于: 十月 22, 2020, 10:39:08 am »
这东西有3个常见做法, 其中两种可以在单kernel内完成, 另外一种需要启动两次kernel. 大致如下:

(1)无kernel代码改动, 简单的将kernel启动两次即可. 第一次上100个block,每个里面依然是统计4096个数中的最大值. 这样你就得到了100个数, 其中最终的最大值必然是这100数之一. 然后第二次继续以1个block, 启动kernel, 从100数中继续选择出一个最大的.

优点: 无代码改动. 缺点: 你需要启动两次kernel, 很多人会绕不过弯.

(2)改写kernel, 只需要启动一次100个block的kernel, 每个block选出局部的一个最大值, 然后将本block的最大值用, atomicMax之类的原子操作处理后, 即可得到全体中的最大值.

优点: 一步到位, 容易理解. 缺点: 目前只有整数支持原子操作求最大值, 浮点数需要用其他方式变通实现原子操作最大值.

(3)改写kernel, 上100个block, 每个block结束前查看自己是否是最后结束的block, 如果是, 进行一次扫尾操作, 对最终的100个局部最大值完成求全局最大值.

优点: 手册上有参考实现, 可以直接抄. 缺点: 可能是最难理解的.

本论坛总是推荐方式(1).

请问第(3)点中所说的手册是指《CUDA C编程手册》还是《CUDA C专家手册》或《CUDA C编程指南》?我准备下载阅读。

12
OpenCL /
« 于: 四月 18, 2018, 05:31:33 pm »
这样做可以是可以, 但是没有必要上2个context, 使用1个context, 里面2个设备即可.

对于这种实质上是一个 ...

很谢谢,我消化一下,会好好考虑你的建议。

13
OpenCL /
« 于: 四月 18, 2018, 03:09:13 pm »
mutex肯定是只能有一个的,这个必须的。 你有多个kernel对象,也可以上多个, 每个对象对应好1个就可以 ...

我修改后是应用当前平台下的2个devices,故为每个device创建了一个context、kernel_imgProc、共享buffers、互斥锁。(每个device一份)。我将cpu_thread=6即前6个host线程放在devices[0]上,后6个host线程放在devices[1]上。
单例PrepareOpenCL.cpp中:
程序代码: [选择]
cl_context_properties context_props[]={CL_CONTEXT_PLATFORM,cl_context_properties(platformInUse),0};
context[0]=clCreateContext(context_props,1,&device[0],NULL,NULL,&status);
context[1]=clCreateContext(context_props,1,&device[1],NULL,NULL,&status);

std::ifstream srcFile("/home/wangdan/ore_granule/multiThreadsFluoreTest12Channels_GPU/objDetectFluoreBmp.cl");
std::string srcProg(std::istreambuf_iterator(srcFile),(std::istreambuf_iterator()));
const char * src = srcProg.c_str();
size_t length = srcProg.length();

cl_program program[number_devices];
program[0]=clCreateProgramWithSource(context[0],1,&src,&length,&status);
program[1]=clCreateProgramWithSource(context[1],1,&src,&length,&status);
status=clBuildProgram(program[0],1,&device[0],NULL,NULL,NULL);
status=clBuildProgram(program[1],1,&device[1],NULL,NULL,NULL);

rgbArray_buffer[0]=clCreateBuffer(context[0],CL_MEM_READ_ONLY,rgbsize*sizeof(uchar),0,&status);
rgbArray_buffer[1]=clCreateBuffer(context[1],CL_MEM_READ_ONLY,rgbsize*sizeof(uchar),0,&status);
Bjimg=imread("/home/wangdan/ore_granule/multiThreadsFluoreTest12Channels_GPU/blackcor_rotate.bmp",0);
if(!Bjimg.data)
{
cout<<"error:BJ-image does not exist!"< }
Bjimg_buffer[0]=clCreateBuffer(context[0],CL_MEM_READ_ONLY,484*364*sizeof(uchar),0,&status);
Bjimg_buffer[1]=clCreateBuffer(context[1],CL_MEM_READ_ONLY,484*364*sizeof(uchar),0,&status);

for(int i=0;i {
if(i<(cpu_thread))
{
queue[i]=clCreateCommandQueue(context[0],device[0],CL_QUEUE_PROFILING_ENABLE,NULL);
status=clEnqueueWriteBuffer(queue[i], rgbArray_buffer[0], CL_FALSE, 0, rgbsize* sizeof(uchar),rgbarray, 0, NULL, NULL);
srcdata_buffer[i] = clCreateBuffer(context[0], CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,NULL);
srcdata_back_buffer[i] = clCreateBuffer(context[0], CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,NULL);
status = clEnqueueWriteBuffer(queue[i], Bjimg_buffer[0], CL_FALSE, 0, 484*364* sizeof(uchar), Bjimg.data, 0, NULL, NULL);
haveStone_buffer[i]=clCreateBuffer(context[0], CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, 2*sizeof(int), NULL,&status);
}
else
{
queue[i]=clCreateCommandQueue(context[1],device[1],CL_QUEUE_PROFILING_ENABLE,NULL);
status=clEnqueueWriteBuffer(queue[i], rgbArray_buffer[1], CL_FALSE, 0, rgbsize* sizeof(uchar),rgbarray, 0, NULL, NULL);
srcdata_buffer[i] = clCreateBuffer(context[1], CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,NULL);
srcdata_back_buffer[i] = clCreateBuffer(context[1], CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, srcdatasize, NULL,NULL);
status = clEnqueueWriteBuffer(queue[i], Bjimg_buffer[1], CL_FALSE, 0, 484*364* sizeof(uchar), Bjimg.data, 0, NULL, NULL);
haveStone_buffer[i]=clCreateBuffer(context[1], CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, 2*sizeof(int), NULL,&status);
}
}

kernel_imgProc[0]=clCreateKernel(program[0],"imgProcess",&status);
status |= clSetKernelArg(kernel_imgProc[0], 0, sizeof(cl_mem), (void*)&rgbArray_buffer[0]);
status |= clSetKernelArg(kernel_imgProc[0], 3, sizeof(cl_int),  &imgwidth);
status |= clSetKernelArg(kernel_imgProc[0], 4, sizeof(cl_int),  &imgheight);
   status |= clSetKernelArg(kernel_imgProc[0], 6, sizeof(cl_int),  &thre_blue_host);
   status |= clSetKernelArg(kernel_imgProc[0], 7, sizeof(cl_mem),  (void*)&Bjimg_buffer[0]);
   status |= clSetKernelArg(kernel_imgProc[0], 8, sizeof(cl_int),  &thre_dis_host);

   kernel_imgProc[1]=clCreateKernel(program[1],"imgProcess",&status);
status |= clSetKernelArg(kernel_imgProc[1], 0, sizeof(cl_mem), (void*)&rgbArray_buffer[1]);
status |= clSetKernelArg(kernel_imgProc[1], 3, sizeof(cl_int),  &imgwidth);
status |= clSetKernelArg(kernel_imgProc[1], 4, sizeof(cl_int),  &imgheight);
   status |= clSetKernelArg(kernel_imgProc[1], 6, sizeof(cl_int),  &thre_blue_host);
   status |= clSetKernelArg(kernel_imgProc[1], 7, sizeof(cl_mem),  (void*)&Bjimg_buffer[1]);
   status |= clSetKernelArg(kernel_imgProc[1], 8, sizeof(cl_int),  &thre_dis_host);

然后在每个host子线程中进行判断,该子线程是属于device[0]还是device[1]上的任务,运用对应的kernel_imgProc和锁执行启动:usingMultiThreads.cpp中:先对每个线程进行判断,是属于哪个设备:
程序代码: [选择]
cpuorgpu=(m_chuteorder<=(PreparePtr->cpu_thread))?0:1;
sumArray_buffer = clCreateBuffer(PreparePtr->context[cpuorgpu], CL_MEM_WRITE_ONLY| CL_MEM_ALLOC_HOST_PTR, sumsize, NULL,NULL);

然后12个host线程读取图片,再在对应的kernel上启动:
程序代码: [选择]
PreparePtr->mutex_thread.lock();
status |= clSetKernelArg(PreparePtr->kernel_imgProc[cpuorgpu], 1, sizeof(cl_mem), (void*)&(PreparePtr->srcdata_buffer[ind]));
status |= clSetKernelArg(PreparePtr->kernel_imgProc[cpuorgpu], 2, sizeof(cl_mem), (void*)&(PreparePtr->srcdata_back_buffer[ind]));
status |= clSetKernelArg(PreparePtr->kernel_imgProc[cpuorgpu], 5, sizeof(cl_mem),  (void*)&sumArray_buffer);
status |= clSetKernelArg(PreparePtr->kernel_imgProc[cpuorgpu], 9, sizeof(cl_mem), (void*)&(PreparePtr->haveStone_buffer[ind]));

status =clEnqueueNDRangeKernel(PreparePtr->queue[ind], PreparePtr->kernel_imgProc[cpuorgpu], 2, NULL, globalsize, localsize,0,NULL,NULL);
PreparePtr->mutex_thread.unlock();

status=clFinish(PreparePtr->queue[ind]);
if (status != CL_SUCCESS)
{
cout<<"Error:clFinish() failed..."< }

int *sumMap=NULL;
sumMap=(int*)clEnqueueMapBuffer(PreparePtr->queue[ind],sumArray_buffer,CL_TRUE, CL_MAP_READ, 0, sumsize, 0, NULL, NULL, &status);

这样书写,正确吗?我现在是前6个在device[0]上 后6个在device[1]上,后6个都是11秒的样子,前6个是后6个的2倍耗时即22秒的样子,很巧的2倍,我想是不是哪里卡住了设备[1]要等待device[0]所以造成了2倍?还是只是device[0]就是比device[1]慢一倍?

14
OpenCL /
« 于: 四月 17, 2018, 07:34:43 pm »
mutex肯定是只能有一个的,这个必须的。 你有多个kernel对象,也可以上多个, 每个对象对应好1个就可以 ...

加了锁,不再报之前段错误吐核以及数据异常的问题。耗时21秒。
多设备的我明天修改后再反馈。

15
OpenCL /
« 于: 四月 17, 2018, 03:45:24 pm »
Hi, 奈奈,

很高兴看到你回来, 但你的BUG修正代码本身存在BUG。


谢谢[名词2]
(1)我将mutex放在单例里,这样就保证所有的12个host线程共用一个mutxt,然后在每个host子线程里刚刚的地方调用:PreparePtr->mutex_thread.lock();    PreparePtr->mutex_thread.unlock();  这样应该是正确的吧?
(2)双设备--也就是DEVICE_GPU  DEVICE_CPU都要,将12个host线程分给它们,减轻原来的DEVICE_GPU的压力。是这个意思吧?  如果是这个意思,我先去了解一下双设备的例子再修改,我以前没写过双设备的。
(3)不建议host上开12个线程以及大buffer切分,这两点我现在先不修改。我先保证不报错和数据不异常。OK?

页: [1] 2 3 ... 5