1

我是 OpenACC 的菜鸟，我尝试优化代码，对于 CPU，我得到：

Time = Time + omp_get_wtime();
    {
      #pragma acc parallel loop
      for (int i = 1;i < k-1; i++)
      {
        jcount[i]=((int)(MLT[i]/dt))+1;
      }
      jcount[0]=0;
      jcount[k-1]=N;

          #pragma acc parallel loop collapse(2)
            for (int i = 0;i < k - 1; i++)
            {
                for(int j=jcount[i];j < jcount[i+1];j++)
                {
                    w[j] = (j*dt - MLT[i])/(MLT[i+1]-MLT[i]);
                    X[j] = MLX[i]*(1-w[j])+MLX[i+1]*w[j];
                    Y[j] = MLY[i]*(1-w[j])+MLY[i+1]*w[j];
                }
            }
    }
Time = omp_get_wtime() - Time;

对于我的 6 核 Intel I7（我关闭了超线程），我的并行化很差，6 核与 1 核之间的差异只有 30%（这意味着 70% 的代码按顺序运行，但我看不出在哪里）

对于 GPU：

...
    acc_init( acc_device_nvidia );
...
TimeGPU = TimeGPU + omp_get_wtime();
    {
      #pragma acc kernels loop independent  copyout(jcount[0:k]) copyin(MLT[0:k],dt)
      for (int i = 1;i < k-1; i++)
      {
        jcount[i]=((int)(MLT[i]/dt))+1;
      }
      jcount[0]=0;
      jcount[k-1]=N;

          #pragma acc kernels loop independent copyout(X[0:N+1],Y[0:N+1]) copyin(MLT[0:k],MLX[0:k],MLY[0:k],dt) copy(w[0:N])
            for (int i = 0;i < k - 1; i++)
            {
                for(int j=jcount[i];j < jcount[i+1];j++)
                {
                    w[j] = (j*dt - MLT[i])/(MLT[i+1]-MLT[i]);
                    X[j] = MLX[i]*(1-w[j])+MLX[i+1]*w[j];
                    Y[j] = MLY[i]*(1-w[j])+MLY[i+1]*w[j];
                }
            }
    }
TimeGPU = omp_get_wtime() - TimeGPU;

并且GPU(gtx1070)比 6 核处理器慢 3 倍！

Launch parameters:
GPU: pgc++ -ta=tesla:cuda9.0 -Minfo=accel -O4
CPU: pgc++ -ta=multicore -Minfo=accel -O4

k = 20000,N = 200万

更新：

更改 GPU 代码：

TimeGPU = TimeGPU + omp_get_wtime();
#pragma acc data create(jcount[0:k],w[0:N]) copyout(X[0:N+1],Y[0:N+1]) copyin(MLT[0:k],MLX[0:k],MLY[0:k],dt)
    {
      #pragma acc parallel loop
      for (int i = 1;i < k-1; i++)
      {
        jcount[i]=((int)(MLT[i]/dt))+1;
      }
      jcount[0]=0;
      jcount[k-1]=N;

          #pragma acc parallel loop
            for (int i = 0;i < k - 1; i++)
            {
                for(int j=jcount[i];j < jcount[i+1];j++)
                {
                    w[j] = (j*dt - MLT[i])/(MLT[i+1]-MLT[i]);
                    X[j] = MLX[i]*(1-w[j])+MLX[i+1]*w[j];
                    Y[j] = MLY[i]*(1-w[j])+MLY[i+1]*w[j];
                }
            }
    }
TimeGPU = omp_get_wtime() - TimeGPU;
    Launch parameters:
    pgc++ -ta=tesla:managed:cuda9.0 -Minfo=accel -O4

现在 GPU 比 CPU 慢 2 倍

输出：

139: compute region reached 1 time
        139: kernel launched 1 time
            grid: [157]  block: [128]
             device time(us): total=425 max=425 min=425 avg=425
            elapsed time(us): total=509 max=509 min=509 avg=509
    139: data region reached 2 times
        139: data copyin transfers: 1
             device time(us): total=13 max=13 min=13 avg=13
    146: compute region reached 1 time
        146: kernel launched 1 time
            grid: [157]  block: [128]
             device time(us): total=13,173 max=13,173 min=13,173 avg=13,173
            elapsed time(us): total=13,212 max=13,212 min=13,212 avg=13,212

为什么我的 TimeGPU 比使用 PGI_ACC_TIME=1 的输出大 2 倍？（30 毫秒对 14 毫秒）

4

1 回答 1

1

我认为很多 GPU 时间是由于内核的内存访问不佳造成的。理想情况下，您希望向量访问连续数据。

“j”循环有多少次迭代？如果超过 32，那么您可以尝试在其上添加一个“#pragma acc 循环向量”，以便它将在向量之间并行化并为您提供更好的数据访问。

此外，您还有很多冗余的内存提取。考虑将具有“i”索引的数组中的值设置为临时变量，以便仅从内存中获取一次值。

于 2018-06-11T20:04:00.847 回答