我的程序需要使用每个线程都对一个向量做处理,最后一步是快速傅里叶变换,参照别的资料写的算法是这样的:
__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",意思是核函数内部资源不够用了是吧?我把块内线程数缩小一半,现在可以运行了,还得看看结果对不对,感谢横扫千军斑竹~还有什么问题我再请教
system
10
[
千军斑竹,有个关于程序优化上的问题还想继续请教一下,像我定义的这种结构体,每个线程处理一个向量的情况下,存取全局内存怎么考虑才能达到coalesced?
system
11
(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成员。
system
12
嗯嗯。昨天是举的一个例子,您可以会启动导致其他类型的同步错误的。恭喜你发现了问题所在!
system
13
以及,在上文的例子中为何我用了int4 *, 而不是直接的struct your_struct *, 大家都是16B的存储结构啊?
因为目前nvcc有个怪癖,如果你不使用int4 *之类的内置类型,将自动将所有成员拆分成多次独立访问的,那样就达不到合并的效果了。这种怪癖行为不仅仅降低了性能,也违反了C标准,也许以后的编译器会改进,但目前,使用int4是较为保险的方式。
system
14
感谢版主的热心回复!这个问题我还得好好消化一下您给出的建议,因为我的实际应用情况是这样的,我的那个结构体对象是一个一维数组
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)>>>这样的,每个线程要处理这样一个结构体做一系列数据处理。请问我这样做真的合适吗……
system
15
LZ您好:
直接这样使用会导致访问无法合并,从而影响您的访存效率。
您可能需要考虑更改数据结构,或者修改您的算法实现了。
祝您好运~
system
16
[
艾斯斑竹说的是,我也觉得这样效果不好,但是又不知道该怎么考虑这个问题,看来我得再重新想想~感谢帮助!