2

对于我的论文，我必须使用 CUDA 优化一个特殊的 MPI-Navier Stokes-Solver 程序。原始程序使用 FFTW 求解多个偏微分方程。详细地说，几个上三角矩阵在二维中进行了傅立叶变换，但作为一维数组处理。目前，我正在努力处理部分原始代码：（N 始终设置为 64）

原来的：

//Does the complex to real in place fft an normalizes
void fftC2R(double complex *arr) {
  fftw_execute_dft_c2r(plan_c2r, (fftw_complex*)arr, (double*)arr);
  //Currently ignored: Normalization
  /*  for(int i=0; i<N*(N/2+1); i++)
arr[i] /= (double complex)sqrt((double complex)(N*N));*/
}

void doTimeStepETDRK2_nonlin_original() {
      //calc velocity
      ux[0] = 0;
      uy[0] = 0;

      for(int i=1; i<N*(N/2+1); i++) {
         ux[i] =  I*kvec[1][i]*qvec[i] / kvec[2][i];
         uy[i] = -I*kvec[0][i]*qvec[i] / kvec[2][i];
      }

      fftC2R(ux);
      fftC2R(uy);

      //do some stuff here...
      //...
      return;
}

其中 ux 和 uy 被分配为（双复数数组）：

ux = (double complex*)fftw_malloc(N*(N/2+1) * sizeof(double complex));
uy = (double complex*)fftw_malloc(N*(N/2+1) * sizeof(double complex));

fft 计划创建为：

plan_c2r = fftw_plan_dft_c2r_2d(N, N,(fftw_complex*) qvec, (double*)qvec, FFTW_ESTIMATE);

其中 qvec 的分配方式与 ux 和 uy 相同，并且数据类型为 double complex。


以下是 CUDA 代码的相关部分：

NN2_VecSetZero_and_init<<<block_size,grid_size>>>();
cudaSafeCall(cudaDeviceSynchronize());
cudaSafeCall(cudaGetLastError());

int err = (int)cufftExecZ2D(cu_plan_c2r,(cufftDoubleComplex*)sym_ux,(cufftDoubleReal*)sym_ux);

if (err != CUFFT_SUCCESS ) {
    exit(EXIT_FAILURE);
    return;
}

err = (int)cufftExecZ2D(cu_plan_c2r,(cufftDoubleComplex*)sym_uy,(cufftDoubleReal*)sym_uy);

if (err != CUFFT_SUCCESS ) {
    exit(EXIT_FAILURE);
    return;
}

//do some stuff here...
//...
return;

其中 sim_ux 和 sim_uy 分配为：

cudaMalloc((void**)&sym_ux, N*(N/2+1)*sizeof(cufftDoubleComplex));
cudaMalloc((void**)&sym_uy, N*(N/2+1)*sizeof(cufftDoubleComplex));

相关袖带部分的初始化看起来像

if (cufftPlan2d(&cu_plan_c2r,N,N, CUFFT_Z2D) != CUFFT_SUCCESS){
    exit(EXIT_FAILURE);
    return -1;
}

if (cufftPlan2d(&cu_plan_r2c,N,N, CUFFT_D2Z)  != CUFFT_SUCCESS){
    exit(EXIT_FAILURE);
    return -1;
}

if ( cufftSetCompatibilityMode ( cu_plan_c2r , CUFFT_COMPATIBILITY_FFTW_ALL) != CUFFT_SUCCESS ) {
    exit(EXIT_FAILURE);
    return -1;
}

if ( cufftSetCompatibilityMode ( cu_plan_r2c , CUFFT_COMPATIBILITY_FFTW_ALL) != CUFFT_SUCCESS ) {
    exit(EXIT_FAILURE);
    return -1;
}


所以我使用完整的 FFTW 兼容性，并使用 FFTW 调用模式调用每个函数。当我运行这两个版本时，我收到的 ux 和 uy（sim_ux 和 sim_uy）的结果几乎相同。但是在数组的周期性位置，Cufft 似乎忽略了这些元素，其中 FFTW 将这些元素的实部设置为零并计算复杂部分（数组太大，无法在此处显示）。发生这种情况的步数为 N/2+1。所以我相信，我并不完全理解 Cufft 和 FFTW 之间的 fft-padding 理论。在调用 Cufft-executions 之前，我可以排除这些数组之间的任何以前的差异。所以上面代码的任何其他数组在这里都不相关。
我的问题是：我是否过于乐观地使用几乎 100% 的 FFTW 调用方式？我必须在 ffts 之前处理我的数组吗？Cufft 文档说，我必须调整数据输入和输出数组的大小。但是，当我运行就地转换时，我该怎么做呢？我真的不想离原始代码太远，我不想为每个 fft 调用使用更多的复制指令，因为内存是有限的，并且数组应该尽可能长时间地在 gpu 上保留和处理。

我感谢每一个提示、批判性陈述或想法！

我的配置：

  • 编译器：gcc 4.6（C99 标准）
  • MPI-Package：mvapich2-1.5.1p1（不应该发挥作用，因为减少了单处理调试模式）
  • CUDA 版本：4.2
  • GPU：CUDA-arch-compute_20（NVIDIA GeForce GTX 570）
  • FFTW 3.3
4

1 回答 1

1

一旦我不得不使用CUFFT，我唯一的解决方案就是独家使用“cu_plan_c2c”——实数和复数数组之间有简单的转换：

- 用 0 填充复杂部分以模拟 cu_plan_r2c

- 在复杂结果上使用 atan2（不是 atan）来模拟 cu_plan_c2r

很抱歉没有为您指出任何更好的解决方案，但这就是我最终解决此问题的方式。希望你不会陷入低内存 cpu 方面的任何困难......

于 2013-04-17T14:13:34.607 回答