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 回答