我的程序需要使用每个线程都对一个向量做处理，最后一步是快速傅里叶变换，参照别的资料写的算法是这样的：
__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
[
艾斯斑竹说的是，我也觉得这样效果不好，但是又不知道该怎么考虑这个问题，看来我得再重新想想~感谢帮助!