程序要实现的功能是数据滑动加窗,为短时傅里叶变换做准备,比如有2048个数据,每隔32个数据点取512点数据,再和512点的窗相乘,得到512点的数据,然后结果保存在另一块空间中,保存的时候按照1024点保存,其余的512点补零。
代码1:
只利用一段共享内存去存储窗数据,然后每个线程完成一组数据加窗和结果保存,
//只利用一个共享内存去存储窗数据,每一个线程完成一组FFT中的数据的变换
//d_data 原始数据 d_win 窗数据 d_res 结果数据 batch FFT组数
//win_len 窗长 FFT_dot_real 补零0后FFT点数 stride 滑窗点数
__global__ void reshape1(float *d_data, float *d_win, float *d_res, int batch,
int win_len, int FFT_dot_real, int stride,int step)
{
int index = threadIdx.x + blockIdx.x*blockDim.x;
int tidx = threadIdx.x;
//int step = (FFT_dot_real - win_len) >> 1;
extern __shared__ float s_win[];
if (tidx<win_len)
{
s_win[tidx] = d_win[tidx];
}
__syncthreads();
for (int i = 0; i < win_len; i++)
{
d_res[index*FFT_dot_real + step + i] = d_data[index*stride + i] * s_win;
}
}
代码2:
//考虑到数据复用,所以利用两块共享内存分别去存储窗数据和原始数据
//d_data 原始数据 d_win 窗数据 d_res 结果数据 batch FFT组数 win_len 窗长 FFT_dot_real 补零0后FFT点数 stride 滑窗点数
__global__ void reshape2(float *d_data, float *d_win, float *d_res, int batch,
int win_len, int FFT_dot_real, int stride)
{
extern __shared__ float s_win[];
__shared__ float s_data[2 * 1024];
int tidx = threadIdx.x;
int bidx = blockIdx.x;
int step = (FFT_dot_real - win_len) >> 1;
int count = blockDim.x / stride;
s_data[tidx] = d_data[bidx*blockDim.x + tidx];
s_data[1024 + tidx] = d_data[(bidx + 1)*blockDim.x + tidx];
if (tidx<win_len)
{
s_win[tidx] = d_win[tidx];//获取窗数据
}
__syncthreads();//栅栏同步
if (tidx < count)
{
for (int i = 0; i < win_len; i++)
{
d_res[(tidx + count*blockIdx.x)*FFT_dot_real + step + i] = s_data[tidx*stride + i] * s_win;
}
}
}
问题:
1. 这两个核函数为什么核函数1在release模式下没有debug模式快很多,而核函数2在release方式下运行比debug模式快10倍左右。
2. debug方式和release方式大概的区别,我只知道release方式似乎会进行一些优化,那为什么核函数1优化后速度提升不太明显呢?是代码本身没有可优化的空间了吗?
3.cuFFT库有没有可以直接对一段数据做STFT的库呢?我本身的需求是每个32点取512个数据加窗,然后补零到1024点后再进行FFT。