GPU做的高斯滤波器

加速计算 CUDA
1

我用GPU做的高斯滤波器，为啥计算结果只有一部分对，其他部分是错的呢，是不是我的边缘计算出现错误

global void gaussFilterkernel(const float dev_input,int width,int height,floatdev_output){
int gausstemp[9]={1,2,1,
2,4,2,
1,2,1};
int x=threadIdx.x+blockIdx.xblockDim.x;
int y=threadIdx.y+blockDim.yblockDim.y;

if(x<width&&y<height)
{
	int index=0;
	float sum=0;
	int counterPixel=x+y*width;//中间像素
	for(int j=-1;j<2;j++)
	{
		for(int i=-1;i<2;i++)
		{
			if((x+i)<0||(x+i)>width||(y+j)<0||(y+j)>height)
			{
				continue;
			}
			int currentPixel=x+i+(y+j)*width;//遍历滤波核内的像素

			sum+=dev_input[currentPixel]*gausstemp[index++];
		}
	}
	dev_output[counterPixel]=sum/16.0;
}

}

2

如果不考虑算法和实现的效率之类的问题，一般来说，
对于宽度为W, 高度为H的数组，
宽度的有效范围是0～W-1
而高度的有效范围则是0~H-1
(您想想是吧）

所以您这里的判断条件也许需要修改：
if((x+i)<0||(x+i)>width||(y+j)<0||(y+j)>height)
请确定您的确不需要(x+i) >= width 以及 您不需要 (y+j) >= height的写法。

如果您的确不需要，请保持大于而不是大于等于的写法，并同时请无视此建议。

3

谢谢版主的回答，明白版主的意思

4

版主
我的上边的程序为什么得到的sum值是1967，但是除上16后赋给dev_output[counterPixel]时，这个dev_output[counterPixel]的值始终是0啊

5

首先说，楼主您确定是除以16.0而不是除以9.0么？（我看到您循环了一共9个点）

其次，关于您说的sum是1967, 除以16后变成了0,请问您有证据么？
您怎么知道得到的是0？

6

以及，楼主您修改了昨日我推荐的将>改成>=的建议了吗？
（这样避免越界可能导致的kernel运行失败）

如果您还没有修改，请先修改，否则kernel会中途失败，得不到结果的。

7

谢谢版主的建议。

8

版主你好
我的C函数的程序运行结果完全正确，代码如下
void GaussFilter(floatinput,int width,int height,floatoutput)
{
int index;
float sum;
int gausstemp[9]={1,2,1,
2,4,2,
1,2,1};
for(int i=0;i<height-1;i++)
{
for(int j=0;j<width-1;j++)
{
index=0;
sum=0;
for(int m=i-1;m<i+2;m++)
{
for(int n=j-1;n<j+1;n++)
{
if(n<0||n>=width||m<0||m>height)
{
continue;
}
sum+=input[mwidth+n]gausstemp[index++];
}
}
output[i*width+j]=sum/16.0;
}
}
}[attach]3319[/attach]
但是我的kernel函数得到的结果却是错的，您昨天提到的那地方已经修改了，那地方是应该除以16，那9个元素的和是16，不是除以9，
我的kernel函数如下，版主帮我看看哪地方越界了
global void gaussFilterkernel(const float dev_input,int width,int height,floatdev_output)
{
float gausstemp[9]={1,2,1,
2,4,2,
1,2,1};
int x=threadIdx.x+blockIdx.xblockDim.x;
int y=threadIdx.y+blockDim.yblockDim.y;

if(x<width&&y<height)
{
int index=0;
float sum=0;
int counterPixel=x+y*width;
for(int j=-1;j<=1;j++)
{
for(int i=-1;i<=1;i++)
{
if((x+i)<0||(x+i)>=width||(y+j)<0||(y+j)>=height)
{
continue;
}
int currentPixel=x+i+(y+j)*width;

sum+=dev_input[currentPixel]*gausstemp[index++];
}
}
dev_output[counterPixel]=sum/16.0;
}
}这是用kernel函数的结果，中间那部分是争取的结果
[attach]3320[/attach]

9

看到您补充的图了，你也修改了，这个不是越界。

根据图像看，问题可能出现在您的处理没有覆盖所有点导致的，
能否让我们看下您的host部分代码？
这样让我们看下您的<<<>>>的形状配置？

10

好的谢谢斑竹了
int GaussFilter(floathos_input,int width,int height,floathos_output)
{
floatdev_input=NULL;
floatdev_output=NULL;
checkCudaErrors(cudaMalloc((void**)&dev_input,widthheightsizeof(float)));
checkCudaErrors(cudaMalloc((void**)&dev_output,widthheightsizeof(float)));
checkCudaErrors(cudaMemset(dev_output,0,widthheightsizeof(float)));

checkCudaErrors(cudaMemcpy(dev_input,hos_input,width*height*sizeof(float),cudaMemcpyHostToDevice));

dim3 Dimblock(16,16,1);
dim3 Dimgrid((width+Dimblock.x-1)/16,(height+Dimblock.y-1)/16,1);

gaussFilterkernel<<<Dimgrid,Dimblock>>>(dev_input,width,height,dev_output);
getLastCudaError("don't call gaussFilterkernel");

checkCudaErrors(cudaMemcpy(hos_output,dev_output,width*height*sizeof(float),cudaMemcpyDeviceToHost));

checkCudaErrors(cudaFree(dev_input));
checkCudaErrors(cudaFree(dev_output));

return 0;

}

11

这个真心奇怪了。看您的代码毫无问题的。
我肉眼未能为您找出问题…:frowning:

建议您考虑：
（1）您确定您给GaussFilter()的hos_input缓冲区，具有正确的内容吗？
（2）您确定此函数返回后，负责显示的代码部分正常吗？

也许其他人能否您更多建议以及启发。。。
我同样在关注中。。。

12

额。发现BUG!!

楼主您居然将blockIdx.y写成了blockDim.y
改回去应该立刻可以修复！！

13

之前未能为您指出这点，我深表遗憾，以及，您的
int y=threadIdx.y+blockDim.y*blockDim.y;

我之前看到这行，下意识的认为您不会写错。。。结果您真的写错了。。。
刚才重新回去细读您的代码，发现这里不对。。。

请修复。

14

谢谢版主的提醒，
:3_45:
是我疏忽了，我也同样认为我的那个地方不会出错，我检查了很多遍，
十分感谢版主

15

善哉。

恭喜楼主成功修复BUG…

都怪我惯性思想了。:slight_smile:

16

[attach]3321[/attach]感谢版主，得到了最后的结果

17

嗯嗯。恭喜楼主。

以及，绝对部分的代码都是楼主您辛苦写的，我只是为您拭去了明珠上的一点灰尘。

欢迎再次来访。

18

版主真不容易，这种bug都看出来了。我感觉cuda在中国的发展和横扫版主以及ice版主的努力和细心有一定关系。

19

qhyjerry您好：

在节假日以赞扬斑竹的内容灌水在本版也是不可以的，敬请注意。

祝您中秋快乐~