求助:内核的fft算法通不过

我的程序需要使用每个线程都对一个向量做处理,最后一步是快速傅里叶变换,参照别的资料写的算法是这样的:

__global__ void kernel_fft(PixelType* d_UnCube, int width, int height)
{
	int idx = threadIdx.x + blockDim.x*blockIdx.x;
	int idy = threadIdx.y + blockDim.y*blockIdx.y;
	int i,j,k,r;
	int la;
	float2 t , temp1, temp2;
	if ((idx < width) && (idy < height))
	{
		/*----FFT算法----*/
		for( i = 1; i <= fft_layer; i++)
		{
			la = (1<<(fft_layer+1-i));//la=2^i代表第i级每个分组所含节点数																			
			//蝶形运算
			for( j = 1; j <= la/2; j++)//la/2代表第i级每个分组所含蝶形运算单元数,同时它也表示每个蝶形运算单元上下节点之间的距离
			{
				r = (j-1)*(1<<(i - 1));
				for( k = j-1; k < fft_dim - 1; k = k + la)//遍历每个分组,分组总数为N/la
				{																			
					t.x = cos(2*Pi*r/fft_dim);
					t.y = -sin(2*Pi*r/fft_dim);
					temp1 = d_UnCube[idy*width + idx].spe[k];
					temp2 = d_UnCube[idy*width + idx].spe[k+la/2];
					d_UnCube[idy*width + idx].spe[k+la/2] = MulFloat2(SubFloat2(temp1 , temp2), t);
					d_UnCube[idy*width + idx].spe[k] = AddFloat2(temp1 , temp2);
				}
			}
		}
		/*----按照倒位序重新排列变换后信号----*/
		for(i=1,j=fft_dim/2;i<=fft_dim-2;i++)
		{
			if(i<j)
			{
				t = d_UnCube[idy*width + idx].spe[j];
				d_UnCube[idy*width + idx].spe[j] = d_UnCube[idy*width + idx].spe[i];
				d_UnCube->spectrum[i] = t;
			}
			k = fft_dim/2;
			while( k <= j)
			{
				j = j-k;
				k = k/2;
			}
			j = j+k;
		}	
		d_UnCube[idy*width + idx].color = 2;
	}
}

PixelType是这样的:

struct PixelType{
   float2   spe[fft_dim];
   unsigned char color;
};

fft_dim是256,fft_layer是8,

 kernel_fft<<< blocksPerGrid, threadsPerBlock>>>( d_UnCube, scubeWidth, scubeHeight);

我屏蔽这段程序中的不同部分,发现红色字体部分是报错的地方,(这段内容换成对color的赋值是可以运行的)。

cudaHyperpro.exe 中的 0x7548b9bc 处最可能的异常: Microsoft C++ 异常: 内存位置 0x0033f888 处的 cudaError_enum。
线程 'Win32 线程' (0xe5c) 已退出,返回值为 0 (0x0)。
线程 'Win32 线程' (0x3b0) 已退出,返回值为 0 (0x0)。

单步启动这个核函数,vs报的错误是这个。
是不是我的访存方式有问题?求大神帮助!

(1)“红色的地方”在哪里?

(2)请无视“最可能的异常”,那个是runtime library内部处理的时候用到的,将被转换成返回值,您应该使用cudaError_t r = cudaDeviceResult();在您的<<<>>>后面来获取返回值的,这样您可以直接知道是什么错误,而不是在那里抓虾。

[

(1)不好意思,,红色的代码段是第三层for循环里面的语句,执行碟形运算的主体。
(2)我这里怎么没有cudaDeviceResult()这个函数,,??

LZ您好:

第二点那个可能是笔误了,您可以在kernel后面加上cudaDeviceSynchronize(),并检查这个函数的返回值。

祝您好运~

不好意思,是cudaDeviceSynchronize(); 这个真心笔误了。我绝对不是故意的。楼主您也不是,对吗?

[

这个返回值是cudaSuccess,但是从结果来看这个kernel并没有执行,不知道是为什么?

那可能是这样,您的<<<>>>导致一个无法被cudaDeviceSynchronize()的错误。

换句话说,您的<<<>>>少见的导致了一次同步错误,而<<<>>>无返回值,同时又不是异步错误,导致用cudaDeviceSynchronzie()无法返回,您可以这样:

cudaError_t result = cudaGetLastError();放置在您的<<<>>>行后,然后这个用来检测您的<<<>>>立刻导致的错误,一般情况下,您的问题会得到cudaErrorInvalidConfiguration之类的。

这表明您的<<<>>>中执行的形状配置硬件无法满足。
(例如,5000个线程的block)

请立刻尝试此建议。

感谢您的建议,明天上班马上试一试。刚开始摸索cuda,还是有很多东西需要学习啊~

[

千军斑竹,早上试了结果发现返回的是"cudaErrorLaunchOutofResources",意思是核函数内部资源不够用了是吧?我把块内线程数缩小一半,现在可以运行了,还得看看结果对不对,感谢横扫千军斑竹~还有什么问题我再请教

[

千军斑竹,有个关于程序优化上的问题还想继续请教一下,像我定义的这种结构体,每个线程处理一个向量的情况下,存取全局内存怎么考虑才能达到coalesced?

(1)那最好将struct拆开,将里面的每个成员单独构成数组。这样方可。例如:
struct your_struct {int a; int b; int c; int d;};

直接一个线程如果访问一个struct的实例的话,例如先访问a, 再访问b,
那么都是不合并的(因为大家都间隔了16B访问4B),
但如果a,b,c,d单独抽出,构成数组, 那么访问a,b,c,d是合并的。

(2)或者,一次性合并的将struct读取到寄存器也可以,还是上文这个例子,
int4 value = *(int4 *)&your_struct的第i个实例;
这样一次性合并读取,然后使用value.x/value.y/value.z/value.w来访问原a,b,c,d成员。

嗯嗯。昨天是举的一个例子,您可以会启动导致其他类型的同步错误的。恭喜你发现了问题所在!

以及,在上文的例子中为何我用了int4 *, 而不是直接的struct your_struct *, 大家都是16B的存储结构啊?

因为目前nvcc有个怪癖,如果你不使用int4 *之类的内置类型,将自动将所有成员拆分成多次独立访问的,那样就达不到合并的效果了。这种怪癖行为不仅仅降低了性能,也违反了C标准,也许以后的编译器会改进,但目前,使用int4是较为保险的方式。

感谢版主的热心回复!这个问题我还得好好消化一下您给出的建议,因为我的实际应用情况是这样的,我的那个结构体对象是一个一维数组

struct PixelType{
   float2   spe[fft_dim];
   unsigned char color;
}pt[640*480];
fft_dim有可能是256,512,甚至是1024

然后线程规划是<<<((640+16-1)/16,(480+16-1)/16),1),(16,16,1)>>>这样的,每个线程要处理这样一个结构体做一系列数据处理。请问我这样做真的合适吗……

LZ您好:

直接这样使用会导致访问无法合并,从而影响您的访存效率。

您可能需要考虑更改数据结构,或者修改您的算法实现了。

祝您好运~

[

艾斯斑竹说的是,我也觉得这样效果不好,但是又不知道该怎么考虑这个问题,看来我得再重新想想~感谢帮助!