不同的线程数下得到的结果不同

加速计算 CUDA
1

用CUDA做的一个高斯滤波的程序，本来定义THREAD_NUM=128的时候计算出来的结果和CPU中完全一致，但是当线程个数达到256的时候就会出现一些细微的不同之处。不知道为什么，当线程个数增加到512的时候，直接只有黑色的图像显示。一个块中最多不是应该可以有512个线程么。测试图像为256256的灰度图。
#define THREAD_NUM 256
int tid = threadIdx.x; int size =nHeight/ THREAD_NUM;
int bid = blockIdx.x;
clock_t start;
if(tid == 0)
{
start=clock();
}
int x;
int y;
for(i=tidsize;i<(tid+1)*size; i++)
{
for(j=0; j<nWidth; j++)
{
double dFilter=0.0;
double dSum = 0.0;
for(x=(-nCenter); x<=nCenter; x++) //行
{
for(y=(-nCenter); y<=nCenter; y++) //列
{

if( (j+x)>=0 && (j+x)<nWidth && (i+y)>=0 && (i+y)<nHeight) //判断边缘
{
dFilter += (double)data[(i+y)*nWidth + (j+x)] * pdKernal_2[(y+nCenter)*nWidowSize+(x+nCenter)];
dSum += pdKernal_2[(y+nCenter)*nWidowSize+(x+nCenter)];
}
}
}
pCanny[i*nWidth+j] = (unsigned char)dFilter/dSum;
}
}

2

CPU中的程序是这样的：
int x;
int y;
for(i=0; i<nHeight; i++)
{
for(j=0; j<nWidth; j++)
{
double dFilter=0.0;
double dSum = 0.0;
for(x=(-nCenter); x<=nCenter; x++) //行
{
for(y=(-nCenter); y<=nCenter; y++) //列
{
if( (j+x)>=0 && (j+x)<nWidth && (i+y)>=0 && (i+y)<nHeight) //判断边缘
{
dFilter += (double)nImageData[(i+y)*nWidth + (j+x)]
* pdKernal_2[(y+nCenter)*nWidowSize+(x+nCenter)];
dSum += pdKernal_2[(y+nCenter)*nWidowSize+(x+nCenter)];
}
}
}
pCanny[i*nWidth+j] = (unsigned char)dFilter/dSum;
}
}

3

楼主您好，

（1）
从代码看没看出问题，不过有个细节：请问您的nHeight能被256和512整除么？
如果不能，这里可能会导致些问题。

（2）在计算能力1.x的卡上，的确一个block最大理论上可以有512个线程，但是也可能达不到（还有其他的因素限制）。在2.0+的卡上，可以最多有1024个。

（以及，话说楼主你将i,j定义成常规概念相反的，即j是横坐标，i是纵坐标。看起来真费劲:)
（以及，如果方便，将kernel的代码发全，现在不全很多数据类型，例如pCanny，pdKernel_2之类的只能靠猜。很费劲）

4

好的，谢谢版主。
我的测试例子是256*256 这一点没有问题。现在我用的线程是128和256进行对比。
全部代码:
#define THREAD_NUM 256

int i,j; //循环变量
int lineByte=(m8/8+3)/44;
int nWidowSize=3;//高斯模板大小
double nSigma = 0.4; //定义高斯函数的标准差
int nCenter = (nWidowSize)/2; //定义滤波窗口中心的索引
int nWidth = m; //获取图像的像素宽度
int nHeight =n; //获取图像的像素高度
//GPU中无法定义数组,数组需要的空间需要全部从CPU中开始定义分配。
double dSum_2 = 0.0; //求和，用于进行归一化
for(i=0; i<nWidowSize; i++)
{
for(int j=0; j<nWidowSize; j++)
{
int nDis_x = i-nCenter;
int nDis_y = j-nCenter;
pdKernal_2[i+j*nWidowSize]=exp(-(1/2)(nDis_xnDis_x+nDis_ynDis_y)
/(nSigmanSigma))/(23.1415926nSigma*nSigma);
dSum_2 += pdKernal_2[i+j*nWidowSize];
}
}
for(i=0; i<nWidowSize; i++)
{
for(int j=0; j<nWidowSize; j++) //进行归一化
{
pdKernal_2[i+j*nWidowSize] /= dSum_2;
}
}
//生成高斯模板
int tid = threadIdx.x;
int size =nHeight/ THREAD_NUM;
int bid = blockIdx.x;

clock_t start;
if(tid == 0)
{
start=clock();
}
int x;
int y;
for(i=tid*size;i<(tid+1)*size; i++)
{
__syncthreads();
for(j=0; j<nWidth; j++)
{
double dFilter=0.0;
double dSum = 0.0;
for(x=(-nCenter); x<=nCenter; x++) //行
{
for(y=(-nCenter); y<=nCenter; y++) //列
{

			   if( (j+x)>=0 && (j+x)<nWidth && (i+y)>=0 && (i+y)<nHeight) //判断边缘  
			   {  
				   dFilter += (double)data[(i+y)*nWidth + (j+x)] * pdKernal_2[(y+nCenter)*nWidowSize+(x+nCenter)];  
				   dSum += pdKernal_2[(y+nCenter)*nWidowSize+(x+nCenter)];  
			   }  
		   }  
	   }  
	   pCanny[i*nWidth+j] = (unsigned char)dFilter/dSum;  
   }

}
if(tid==0)
{
*time=clock()-start;
}

5

不知道为何图像无法上传上去。

6

测试了一下，发现一个很奇怪的问题，在128线程的时候无论是256256的图像还是10241024的图像都不会出现在与CPU计算结果不符合的情况，但是一旦超过128线程就会出现不相符的情况。

7

话说楼主能否在256的时候，给出pCanny中的不对的点的下标？
（无论是线性的或者是(j,i)的形式均可）

有得时候顺着看代码看不出什么，但也许根据计算结果不同的点的坐标逆推，会有启发。

8

以及，楼主你依然发的不全。。唉。。。

Kernel的首行（含有参数列表和类型的那行）那行哪里去了？
调用kernel的那行呢？
等等。

麻烦请发一下。本来直接顺看的就折腾，你还有保留的。。

9

非常感谢版主的耐心，第一次发帖很多地方不懂还望见谅。昨天仔细想了一下修改了下程序，把我原本在kernel中生成的高斯模板放到CPU中去完成，改进后只把完成好的55高斯模板放到GPU中去即可，确是一开始没考虑到这个问题只是简单地把CPU中的程序丢到kernel中。这样处理后，问题基本上解决了，我想问题可能是在线程同步的情况下，每个线程都会去执行高斯模板的生成步骤，会导致高斯模板生成的时候出现一些错误，因为高斯模板的数组存贮在Global Memory中（不知道这样解释是否正确）。每个线程在执行的时候都回去访问这个数组。修改后的代码：
global static void GpuCalculate(double pdKernal_2,unsigned char *pCanny,unsigned char data,clock_t time,int m,int n)
{

int i,j; //循环变量
int lineByte=(m8/8+3)/44;
int nWidowSize=5;
int nCenter = nWidowSize/2; //定义滤波窗口中心的索引
int nWidth = m; //获取图像的像素宽度
int nHeight =n; //获取图像的像素高度
//GPU中无法定义数组,数组需要的空间需要全部从CPU中开始定义分配。
int tid = threadIdx.x;
int size =nHeight/ THREAD_NUM;
int bid = blockIdx.x;

clock_t start;
if(tid == 0)
{
start=clock();
}
int x;
int y;
for(i=tid*size;i<(tid+1)size; i++)
{
for(j=0; j<nWidth; j++)
{
double dFilter=0.0;
double dSum = 0.0;
for(x=(-nCenter); x<=nCenter; x++) //行
{
for(y=(-nCenter); y<=nCenter; y++) //列
{
if( (j+x)>=0 && (j+x)<nWidth && (i+y)>=0 && (i+y)<nHeight) //判断边缘
{
dFilter += (double)data[(i+y)nWidth + (j+x)] * pdKernal_2[(y+nCenter)nWidowSize+(x+nCenter)];
dSum += pdKernal_2[(y+nCenter)nWidowSize+(x+nCenter)];
}
}
}
pCanny[i*nWidth+j] = (unsigned char)dFilter/dSum;
}
}
if(tid==0)
{
time=clock()-start;
}
}
主函数中：
cudaMalloc((void) &time, sizeof(clock_t));
//在显卡内部开辟一个显示时间的内存
cudaMalloc((void) &gpudata, sizeof(unsigned char)xy);
cudaMalloc((void*) &pdKernal, sizeof(double)55);
cudaMalloc((void**) &GpuTempdata,sizeof(unsigned char)xy);
//开辟内存存放数组
cudaMemcpy(gpudata, pBmpbuf, sizeof(unsigned char)xy,cudaMemcpyHostToDevice);
cudaMemcpy(pdKernal, pdKernal_2, sizeof(double)55,cudaMemcpyHostToDevice);
//将CPU中的的数组复制到显卡内存中
GpuCalculate<<<BLOCK_NUM,THREAD_NUM,0>>>(pdKernal,GpuTempdata,gpudata,time,x,y);
//在Gpu中进行计算
clock_t time_used;
//在Cpu中定义Gpu计算的时间变量
cudaMemcpy(GetResult, GpuTempdata, sizeof(unsigned char)xy,cudaMemcpyDeviceToHost);
//取出GPU中计算的数据
cudaMemcpy(&time_used,time,sizeof(clock_t),cudaMemcpyDeviceToHost);
//从Gpu显卡中取出时间
saveBmp(GpuwritePath,GetResult,bmpWidth, bmpheight,8,bmpColorTable);