Использование OpenACC для распараллеливания вложенных циклов

Question

Использование OpenACC для распараллеливания вложенных циклов

Я очень новичок в openacc и обладаю только знаниями высокого уровня, поэтому любая помощь и объяснение того, что я делаю неправильно, будут оценены.

Я пытаюсь ускорить (распараллелить) не такой простой вложенный цикл, который обновляет уплощенный (от 3D до 1D) массив, используя директивы openacc. Ниже я разместил упрощенный пример кода, который при компиляции с использованием

pgcc -acc -Minfo=accel test.c

выдает следующую ошибку:

call to cuStreamSynchronize returned error 700: Illegal address during kernel execution

Код:

#include <stdio.h>
#include <stdlib.h>

#define min(a,b) (a > b) ? b : a
#define max(a,b) (a < b) ? b : a

#define NX 10
#define NY 10
#define NZ 10

struct phiType {
  double dx, dy, dz;
  double * distance;
};

typedef struct phiType Phi;

#pragma acc routine seq
double solve(Phi *p, int index) {
  // for simplicity just returning a value
  return 2;
}

void fast_sweep(Phi *p) {

  // removing boundaries
  int x = NX - 2; 
  int y = NY - 2;
  int z = NZ - 2;

  int startLevel = 3;
  int endLevel   = x + y + z;

  #pragma acc data copy(p->distance[0:NX*NY*NZ])
  for(int level = startLevel; level <= endLevel; level++){
    int ks = max(1, level-(y + z));
    int ke = min(x, level-2);

    int js = max(1, level-(x + z));
    int je = min(y, level-2);

    #pragma acc region
    {
      #pragma acc loop independent
      for(int k = ks; k <= ke; k++){
        #pragma acc loop independent
        for(int j = js; j <= je; j++){
          int i = level - (k + j);
          if(i > 0 && i <= z){
            int index = i * NX * NY + j * NX + k;
            p->distance[index] = solve(p, index);
          }
        }
      }
    }
  }
}


void create_phi(Phi *p){

  p->dx = 1;
  p->dy = 1;
  p->dz = 1;

  p->distance = (double *) malloc(sizeof(double) * NX * NY * NZ);
  for(int i = 0; i < NZ; i++){
    for(int j = 0; j < NY; j++){
      for(int k = 0; k < NX; k++){
        int index = i * NX * NY + j * NX + k;
        p->distance[index] = (i*j*k == 0) ? 0 : 1;
      }
    }
  }

}


int main()
{
  printf("start \n");

  Phi *p = (Phi *) malloc(sizeof(Phi));
  create_phi(p);

  printf("calling fast sweep \n");
  fast_sweep(p);

  printf(" print the results \n");
  for(int i = 0; i < NZ; i++){
    for(int j = 0; j < NY; j++){
      for(int k = 0; k < NX; k++){
        int index = i * NX * NY + j * NX + k;
        printf("%f ", p->distance[index]);
      }
      printf("\n");
    }
    printf("\n");
  }

  return 0;
}

Вместо использования region а также loop директивы, используя

#pragma acc kernels

выдает следующую ошибку:

solve:
     19, Generating acc routine seq
fast_sweep:
     34, Generating copy(p->distance[:1000])
     42, Generating copy(p[:1])
     45, Loop carried dependence due to exposed use of p[:1] prevents parallelization
         Accelerator scalar kernel generated
     47, Loop carried dependence due to exposed use of p[:i1+1] prevents parallelization

Я запускаю этот код на

GNU/Linux
CentOS release 6.7 (Final)
GeForce GTX Titan
pgcc 15.7-0 64-bit target on x86-64 Linux -tp sandybridge

2

c openacc pgi pgcc

Источник

user4581734 24 авг '15 в 19:56

2 ответа

Решение

Джефф прав, что комитет OpenACC все еще работает над тем, как стандартизировать поддержку агрегированных типов данных с динамическими членами данных. Однако в версии PGI 14.9 или новее мы добавили лучшую поддержку структур, а также классов C++, поэтому в этом случае вы можете упростить код, просто добавив create(p[0:1]), Что произойдет, так это то, что компилятор создаст копию устройства p с памятью, выделенной только для членов данных. Затем, когда вы делаете копию p->distance, память будет выделена для "расстояния", а затем прикрепить его к p, (т.е. время выполнения заполнит указатель устройства в структуре).

Есть предостережения. Во-первых, это поведение не было стандартизировано, поэтому другие компиляторы, такие как Cray, Pathscale, GCC и другие, могут вести себя по-разному. Во-вторых, порядок имеет значение. p должен быть создан до distance можно прикрепить В-третьих, более сложными структурами данных становится очень трудно управлять. Как предполагает Джефф, использование CUDA Unified Memory является хорошей альтернативой для управления сложными структурами данных.

Если вам интересно, большая часть моей презентации GTC2015 обсуждает эту тему ( ссылка). Основное внимание в докладе уделяется управлению данными класса C++, но оно также применимо и к структурам языка Си.

Надеюсь, это поможет, Мат

% cat test1.c
#include <stdio.h>
#include <stdlib.h>

#define min(a,b) (a > b) ? b : a
#define max(a,b) (a < b) ? b : a

#define NX 10
#define NY 10
#define NZ 10

struct phiType {
  double dx, dy, dz;
  double * distance;
};

typedef struct phiType Phi;

#pragma acc routine seq
double solve(Phi *p, int index) {
  // for simplicity just returning a value
  return 2;
}

void fast_sweep(Phi *p) {

  // removing boundaries
  int x = NX - 2;
  int y = NY - 2;
  int z = NZ - 2;

  int startLevel = 3;
  int endLevel   = x + y + z;

  #pragma acc data create(p[0:1]) copy(p->distance[0:NX*NY*NZ])
  for(int level = startLevel; level <= endLevel; level++){
    int ks = max(1, level-(y + z));
    int ke = min(x, level-2);

    int js = max(1, level-(x + z));
    int je = min(y, level-2);

    #pragma acc region
    {
      #pragma acc loop independent
      for(int k = ks; k <= ke; k++){
        #pragma acc loop independent
        for(int j = js; j <= je; j++){
          int i = level - (k + j);
          if(i > 0 && i <= z){
            int index = i * NX * NY + j * NX + k;
            p->distance[index] = solve(p, index);
          }
        }
      }
    }
  }
}


void create_phi(Phi *p){

  p->dx = 1;
  p->dy = 1;
  p->dz = 1;

  p->distance = (double *) malloc(sizeof(double) * NX * NY * NZ);
  for(int i = 0; i < NZ; i++){
    for(int j = 0; j < NY; j++){
      for(int k = 0; k < NX; k++){
        int index = i * NX * NY + j * NX + k;
        p->distance[index] = (i*j*k == 0) ? 0 : 1;
      }
    }
  }

}


int main()
{
  printf("start \n");

  Phi *p = (Phi *) malloc(sizeof(Phi));
  create_phi(p);

  printf("calling fast sweep \n");
  fast_sweep(p);

  printf(" print the results \n");
  for(int i = 0; i < NZ; i++){
    for(int j = 0; j < NY; j++){
      for(int k = 0; k < NX; k++){
        int index = i * NX * NY + j * NX + k;
        printf("%f ", p->distance[index]);
      }
      printf("\n");
    }
    printf("\n");
  }

  return 0;
}

% pgcc -acc -ta=tesla:cc35 -Minfo=accel test1.c -V15.7 ; a.out
solve:
     19, Generating acc routine seq
fast_sweep:
     34, Generating create(p[:1])
         Generating copy(p->distance[:1000])
     45, Loop is parallelizable
     47, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         45, #pragma acc loop gang /* blockIdx.y */
         47, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
start
calling fast sweep
 print the results
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000

0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000
0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000

1

Источник

user3204484 25 авг '15 в 15:31

Другие вопросы по тегам c openacc pgi pgcc

user1658449 24 авг '15 в 21:31 2015-08-24 21:31 · Accepted Answer · 2015-08-24 21:31

Ошибка исходит от вычислительного ядра на графическом процессоре с разыменованием указателя процессора. Это довольно распространенная проблема, над которой работает комитет OpenACC. Подобные динамические структуры данных действительно могут вызвать много проблем, поэтому мы хотим это исправить. Вот два возможных обходных пути для вас.

1) Используйте "управляемую память" через PGI "унифицированный пакет оценки памяти" во время установки компилятора. Это бета- функция, но она помещает все ваши данные в специальный тип памяти, видимый как для процессора, так и для графического процессора. Есть много предостережений, о которых вы должны прочитать в документации, а именно о том, что вы ограничены объемом памяти, доступным на GPU, и что вы не можете получить доступ к памяти из CPU, пока она используется в GPU, но это один из возможных обходных путей. Предполагая, что вы включили эту опцию во время установки, просто добавьте -ta=tesla:managed на флаги вашего компилятора, чтобы включить его. Я попробовал это с вашим кодом, и это сработало.

2) Добавьте указатель на ваш код, чтобы вы не имели доступа distance через p, но вы обращаетесь к нему напрямую, вот так:

double *distance = p->distance;
#pragma acc data copy(p[0:1],distance[0:NX*NY*NZ])
  for(int level = startLevel; level <= endLevel; level++){
    int ks = max(1, level-(y + z));
    int ke = min(x, level-2);

    int js = max(1, level-(x + z));
    int je = min(y, level-2);

    #pragma acc parallel
    {
      #pragma acc loop independent
      for(int k = ks; k <= ke; k++){
        #pragma acc loop independent
        for(int j = js; j <= je; j++){
          int i = level - (k + j);
          if(i > 0 && i <= z){
            int index = i * NX * NY + j * NX + k;
            distance[index] = solve(p, index);
          }
        }
      }
    }

Я знаю, что это может быть проблемой, когда есть много массивов данных для этого, но это обходной путь, который я успешно использовал во многих кодах. К сожалению, это необходимо, поэтому мы хотели бы предложить лучшее решение в будущей версии OpenACC.

Надеюсь, это поможет! Если я смогу найти решение, которое не требует дополнительного указателя, я обновлю этот ответ.