两个CUDA程序,一个在release方式运行比debug方式可以快很多,另一个则基本不变,大概原因是什么呢

  • 2 replies
  • 844 views
程序要实现的功能是数据滑动加窗,为短时傅里叶变换做准备,比如有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。

程序要实现的功能是数据滑动加窗,为短时傅里叶变换做准备,比如有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。

虽然我们一般不讨论Debug情况下的性能(而只讨论它对调试的辅助和正确性问题). 但是第一个Kernel看起来颇为明显. 应当是明显的卡在访存上的.

虽然编译器可以生成劣化很多的指令在Debug下, 或者说一些没有生成一些很好的优化指令. 但是显存自身并不存在明显的所谓"Debug版本", 编译器也不存在明显的会"疯狂读取N倍显存之类的"故意捣鬼. 所以kernel如果是受限于访存自身, 而不在于计算指令之类的上有明显的短板, 则可以Debug版本降速不明显. 楼主可以直接用profiler看下debug版本的计算和访存情况(虽然这样做往往无意义)

嗯嗯,好的知道了,十分感谢 :)