Оптимизация OpenACC против GPU - PullRequest
Новая
       69

Оптимизация OpenACC против GPU

0 голосов
/ 03 июня 2018

Я noob в OpenACC, я пытаюсь оптимизировать код, для процессора я получаю: 

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;

Для моего Intel I7 (я отключил Hyper-Threading) с 6 ядрами, я получаю плохое распараллеливание, разницамежду 6 ядрами против 1 только 30% (это означает, что 70% кода выполняется последовательно, но я не вижу где)

Для графических процессоров: 

...
    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;

И Графический процессор(gtx1070) 3 раза медленнее, чем 6-ядерный процессор! 

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

k = 20000, N = 2 миллиона

ОБНОВЛЕНИЕ:

изменить код 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 в 2 раза медленнее, чем CPU

Вывод: 

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 в 2 раза больше по сравнению с выводом с использованием PGI_ACC_TIME =1?(30 мс против 14 мс)

1 Ответ

0 голосов
/ 11 июня 2018

Я думаю, что большая часть времени на GPU связана с плохим доступом к памяти ваших ядер.В идеале вы хотите, чтобы векторы имели доступ к смежным данным.

Сколько итераций составляет цикл "j"?Если длиннее 32, вы можете попробовать добавить к нему «вектор цикла петли #pragma», чтобы он распараллеливался по векторам и обеспечил вам лучший доступ к данным.

Кроме того, у вас много избыточной памятивыбирает.Рассмотрите возможность установки значений из массивов с индексами «i» во временные переменные, чтобы значения выбирались из памяти только один раз.

...