高斯金字塔

///这是做高斯金字塔,怎样让这两个kernel函数合成一个呢,我就简单的一合并,结果图像是黑色的
global void GaussianSF_kernel(unsigned char* SrcBuf,unsigned char* DstBuf,float* Weight,int w,int h,float S,int nbands)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
int j = blockDim.y * blockIdx.y + threadIdx.y;

if (i > 1  && i < w - 2   && j > 1 && j < h - 2 )
{
	for (int n=0;n < nbands;n++)
	{
		/************************************************************************/
		/* 计算权重矩阵                                                         */
		/************************************************************************/
		Weight[0] = 1/(C*S*S)*(exp(-(8/2*S*S)));
		Weight[1] = 1/(C*S*S)*(exp(-(5/2*S*S)));
		Weight[2] = 1/(C*S*S)*(exp(-(4/2*S*S)));
		Weight[3] = 1/(C*S*S)*(exp(-(5/2*S*S)));
		Weight[4] = 1/(C*S*S)*(exp(-(8/2*S*S)));

		Weight[5] = 1/(C*S*S)*(exp(-(5/2*S*S)));
		Weight[6] = 1/(C*S*S)*(exp(-(2/2*S*S)));
		Weight[7] = 1/(C*S*S)*(exp(-(1/2*S*S)));
		Weight[8] = 1/(C*S*S)*(exp(-(2/2*S*S)));
		Weight[9] = 1/(C*S*S)*(exp(-(5/2*S*S)));

		Weight[10] = 1/(C*S*S)*(exp(-(4/2*S*S)));
		Weight[11] = 1/(C*S*S)*(exp(-(1/2*S*S)));
		Weight[12] = 1/(C*S*S);
		Weight[13] = 1/(C*S*S)*(exp(-(1/2*S*S)));
		Weight[14] = 1/(C*S*S)*(exp(-(4/2*S*S)));

		Weight[15] = 1/(C*S*S)*(exp(-(5/2*S*S)));
		Weight[16] = 1/(C*S*S)*(exp(-(2/2*S*S)));
		Weight[17] = 1/(C*S*S)*(exp(-(1/2*S*S)));
		Weight[18] = 1/(C*S*S)*(exp(-(2/2*S*S)));
		Weight[19] = 1/(C*S*S)*(exp(-(5/2*S*S)));

		Weight[20] = 1/(C*S*S)*(exp(-(8/2*S*S)));
		Weight[21] = 1/(C*S*S)*(exp(-(5/2*S*S)));
		Weight[22] = 1/(C*S*S)*(exp(-(4/2*S*S)));
		Weight[23] = 1/(C*S*S)*(exp(-(5/2*S*S)));
		Weight[24] = 1/(C*S*S)*(exp(-(8/2*S*S)));


		float sum = Weight[0] + Weight[1] + Weight[2] + Weight[3] + Weight[4] + Weight[5] + Weight[6] + Weight[7] + Weight[8]
		+Weight[9] + Weight[10] + Weight[11] + Weight[12] + Weight[13] + Weight[14] + Weight[15] + Weight[16] + Weight[17]
		+Weight[18] + Weight[19] + Weight[20] + Weight[21] + Weight[22] + Weight[23] + Weight[24];
		/************************************************************************/
		/* 九个点的像素值                                                       */
		/************************************************************************/
		float b[25];
		float G = 0.0;	
		b[0] = (float)SrcBuf[(j+2)*w + i-2+n*w*h];
		b[1] = (float)SrcBuf[(j+2)*w + i-1+n*w*h];
		b[2] = (float)SrcBuf[(j+2)*w + i+n*w*h];
		b[3] = (float)SrcBuf[(j+2)*w + i+1+n*w*h];
		b[4] = (float)SrcBuf[(j+2)*w + i+2+n*w*h];

		b[5] = (float)SrcBuf[(j+1)*w + i-2+n*w*h];
		b[6] = (float)SrcBuf[(j+1)*w + i-1+n*w*h];
		b[7] = (float)SrcBuf[(j+1)*w + i+n*w*h];
		b[8] = (float)SrcBuf[(j+1)*w + i+1+n*w*h];
		b[9] = (float)SrcBuf[(j+1)*w + i+2+n*w*h];

		b[10] = (float)SrcBuf[j*w + i-2+n*w*h];
		b[11] = (float)SrcBuf[j*w + i-1+n*w*h];
		b[12] = (float)SrcBuf[j*w + i+n*w*h];
		b[13] = (float)SrcBuf[j*w + (i+1)+n*w*h];
		b[14] = (float)SrcBuf[j*w + (i+2)+n*w*h];

		b[15] = (float)SrcBuf[(j-1)*w + i-2+n*w*h];
		b[16] = (float)SrcBuf[(j-1)*w + i-1+n*w*h];
		b[17] = (float)SrcBuf[(j-1)*w + i+n*w*h];
		b[18] = (float)SrcBuf[(j-1)*w + i+1+n*w*h];
		b[19] = (float)SrcBuf[(j+2)*w + i-1+n*w*h];

		b[20] = (float)SrcBuf[(j-2)*w + i-2+n*w*h];
		b[21] = (float)SrcBuf[(j-2)*w + i-1+n*w*h];
		b[22] = (float)SrcBuf[(j-2)*w + i+n*w*h];
		b[23] = (float)SrcBuf[(j-2)*w + i+1+n*w*h];
		b[24] = (float)SrcBuf[(j+2)*w + i-2+n*w*h];
		/************************************************************************/
		/* 计算高斯模糊的值                                                     */
		/************************************************************************/

		for (int k=0;k<25;k++)
		{
			G += b[k] * Weight[k]/sum;
		}
		if (G >= 255.0) G = 255.0;
		DstBuf[j*w + i+n*w*h] = G;
	}
}

}
global void Zoom_kernel(unsigned char* d_SrcBuf, unsigned char * d_DstBuf, int w, int h,int bands,int z)
{
//降采样
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
int i0,j0,nw,nh;
nw = (int)(w/pow(2.0,z)+0.5);
nh = (int)(h/pow(2.0,z)+0.5);
i0 = (int)(ipow(2.0,z));
j0 = (int)(j
pow(2.0,z));
if(i < nw && j < nh)
{
for(int k = 0;k < bands;k++)
{
d_DstBuf[kwh+wj+i] = d_SrcBuf[kwh+wj0+i0];
}
}
}

LZ您好:

请您在充分理解算法实际情况的前提下,改写合适的实现。

以及,如果后面一个kernel需要依赖于前面kernel的全局结果,那么可能是无法合并的,必须通过前面一个kernel的结束达成全局的同步。

以及,新invoke kernel并不十分耗费时间的,您也可以选择不合并。

以及,本版仅讨论CUDA自身的问题,而不能就您的具体问题直接告诉您正确的CUDA代码实现,请谅解。

祝您好运~

恩,后面的确实依赖前面的全局结果,经过版主的指点,我知道该继续往什么方向做了,非常感谢

LZ您好:

不客气的,稍微再补充一下,如果是后面kernel中每个线程只使用前面kernel中对应的一个线程的结果,那么是可以的(当然这个也不叫“依赖于全局结果”),如果不是这样,那么就不行。

因为全局同步的唯一手段是kernel结束,如果直接合并,某些运行进度比较靠前的线程需要读取原来第一个kernel的结果的时候,生成这些结果的线程并不一定已经将这些结果生成完毕(同时也没有方法全局同步来保证这一点),那么前面那些线程读取到的都是错误的数据,结果自然不能保证。

祝您编码顺利~

您好,如果我的第二个kernel中线程数量是第一个中的一半,这是不是就不可以合并啊,就是依赖于全局结果?

LZ您好:

这个取决于您的算法逻辑的,不能简单从线程数目角度直接予以判定。

希望您能理解这一点。

祝您好运~