【IT168 文档】今天尝试用cuda把FFT实现，遇到了难题。直接调用cufft库的话，内存拷贝与数据处理的时间比大约是1：2。但是据说cufft并不是最高效的，所以想自己锻炼一下。
我的思路是将二维的每一行映射到一个block，每个点都是一个thread。
先将数据拷贝到显卡全局内存，然后拷贝到每个block的共享内存，这是因为读取global memory会占用更多的指令周期。
然后就是处理这段shared memory，可是每当写的时候就会很慢，后来仔细阅读了一下bank conflict的部分，将相邻的线程所对应的memory错开存储，果然好多了，可是这个虽然只包括一维fft(还没做矩阵转置和另一个一维变换)，还是要耗费掉和cufft库差不多的时间。
仔细检查发现是在写内存的时候占用的时间比例很大，不知道怎么才能缩短时间。
view plaincopy to clipboardprint?
__global__
void FFT_2D_Radix2(DATA_TYPE* dg_buffer, int N )
{
int tid, rev, pos, pre, stride = 33;
tid = threadIdx.x;
rev = bit_reverse3(tid, tail_zero_nums(N));
__shared__ DATA_TYPE s_DataR[MATRIX_SIZE]; // 512*4 = 2kB
__shared__ DATA_TYPE s_DataI[MATRIX_SIZE]; // 512*4 = 2kB
__shared__ DATA_TYPE s_CosTable[MATRIX_SIZE]; // 512*4 = 2kB
__shared__ DATA_TYPE s_SinTable[MATRIX_SIZE]; // 512*4 = 2kB
pos = tid * stride % MATRIX_SIZE;
s_DataR[pos] = dg_buffer[blockIdx.x * BLOCK_SIZE + rev];
s_DataI[pos] = dg_buffer[N*N + blockIdx.x * BLOCK_SIZE + rev];
float theta = GV_2PI / N;
s_SinTable[pos] = __sinf( theta * tid );
s_CosTable[pos] = __cosf( theta * tid );
__syncthreads();
int step, w;
for(step = 1;step<N;step=step*2)
{
if(tid & step)
{
w = ( tid & ( step - 1 ) ) * stride % MATRIX_SIZE;
DATA_TYPE tempR = s_DataR[pos] * s_CosTable[w] + s_DataI[pos] * s_SinTable[w];
DATA_TYPE tempI = s_DataI[pos] * s_CosTable[w] - s_DataR[pos] * s_SinTable[w];
pre = ( tid - step ) * stride % MATRIX_SIZE;
s_DataR[pos] = s_DataR[pre] - tempR;
s_DataI[pos] = s_DataI[pre] - tempI;
s_DataR[pre] += tempR;
s_DataI[pre] += tempI;
}
__syncthreads();
}
}
void FFT_2D_Radix2(DATA_TYPE* dg_buffer, int N )
{
int tid, rev, pos, pre, stride = 33;
tid = threadIdx.x;
rev = bit_reverse3(tid, tail_zero_nums(N));
__shared__ DATA_TYPE s_DataR[MATRIX_SIZE]; // 512*4 = 2kB
__shared__ DATA_TYPE s_DataI[MATRIX_SIZE]; // 512*4 = 2kB
__shared__ DATA_TYPE s_CosTable[MATRIX_SIZE]; // 512*4 = 2kB
__shared__ DATA_TYPE s_SinTable[MATRIX_SIZE]; // 512*4 = 2kB
pos = tid * stride % MATRIX_SIZE;
s_DataR[pos] = dg_buffer[blockIdx.x * BLOCK_SIZE + rev];
s_DataI[pos] = dg_buffer[N*N + blockIdx.x * BLOCK_SIZE + rev];
float theta = GV_2PI / N;
s_SinTable[pos] = __sinf( theta * tid );
s_CosTable[pos] = __cosf( theta * tid );
__syncthreads();
int step, w;
for(step = 1;step<N;step=step*2)
{
if(tid & step)
{
w = ( tid & ( step - 1 ) ) * stride % MATRIX_SIZE;
DATA_TYPE tempR = s_DataR[pos] * s_CosTable[w] + s_DataI[pos] * s_SinTable[w];
DATA_TYPE tempI = s_DataI[pos] * s_CosTable[w] - s_DataR[pos] * s_SinTable[w];
pre = ( tid - step ) * stride % MATRIX_SIZE;
s_DataR[pos] = s_DataR[pre] - tempR;
s_DataI[pos] = s_DataI[pre] - tempI;
s_DataR[pre] += tempR;
s_DataI[pre] += tempI;
}
__syncthreads();
}
}
