Вложенные директивы в OpenACC

Question

Я пытаюсь использовать вложенную функцию OpenACC для активного динамического параллелизма моей карты gpu. У меня Tesla 40c, а мой компилятор OpenACC - это версия PGI 15.7.

Мой код настолько прост. Когда я пытаюсь скомпилировать следующий код компилятора возвращает мне эти сообщения

PGCC-S-0155-Illegal context for pragma: acc parallel loop (test.cpp: 158) 
PGCC/x86 Linux 15.7-0: compilation completed with severe errors 

Моя структура кода:

#pragma acc parallel loop 
for(i = 0; i < N; i++) 
{ 
    // <<computation>> 

    int ss = A[tid].start; 
    int ee = A[tid].end; 

    #pragma acc parallel loop 
    for(j = ss; j< (ee + ss); j++) 
    { 
    // <<computation>> 
    } 

Я также пытался изменить свой код, чтобы использовать обычные директивы. Но я не мог собрать снова

#pragma acc routine workers 
foo(...) 
{ 

    #pragma acc parallel loop 
    for(j = ss; j< (ee + ss); j++) 
    { 
    // <<computation>> 
    } 
} 

#pragma acc parallel loop 
for(i = 0; i < N; i++) 
{ 
    // <<computation>> 

    int ss = A[tid].start; 
    int ee = A[tid].end; 

    foo(...); 

} 

Я попытался, конечно, только с рутиной (далее, рабочая, банда) без внутренней параллельной петли директивы. Это был компилятор, но динамический параллелизм не был активирован.

37, Generating acc routine worker 
     Generating Tesla code 
     42, #pragma acc loop vector, worker /* threadIdx.x threadIdx.y */ 
     Loop is parallelizable 

Как я могу использовать динамический параллелизм в OpenACC?

Answer 1

Как я могу использовать динамический параллелизм в OpenACC?

Хотя вложенные области (которые предположительно используют динамический параллелизм) является a new feature in the OpenACC 2.0 specification, я не верю, что это еще не реализовано в PGI 15.7. PGI 15.7 представляет собой частичную реализацию спецификации OpenACC 2.0.

Это ограничение задокументированы в релизе отмечается, PGI 15.7, который должен поставляться с PGI 15.7 компилятором (pgirn157.pdf) в разделе 2.7 (эти примечания к выпуску в настоящее время доступны here):

OpenACC 2.0 Недостающие Особенности

‣ Директива декларации ссылок для глобальных данных не реализована.

‣ Вложенный параллелизм (параллельные и ядерные конструкции в области параллельных или ядер) не выполнен .

На основе комментариев, есть некоторая озабоченность по поводу #pragma acc routine worker, поэтому здесь полностью работает пример с PGI 15.7 этого:

$ cat t1.c 
#include <stdio.h> 
#include <stdlib.h> 
#define D1 4096 
#define D2 4096 
#define OFFS 2 

#pragma acc routine worker 
void my_set(int *d, int len, int val){ 
    int i; 
    for (i = 0; i < len; i++) d[i] += val+OFFS; 
} 

int main(){ 


    int i,*data; 
    data = (int *)malloc(D1*D2*sizeof(int)); 
    for (i = 0; i < D1*D2; i++) data[i] = 1; 

#pragma acc kernels copy(data[0:D1*D2]) 
    for (i = 0; i < D1; i++) 
    my_set(data+(i*D2), D2, 1); 

    printf("%d\n", data[0]); 

    return 0; 
} 
$ pgcc -acc -ta=tesla -Minfo=accel t1.c -o t1 
my_set: 
     8, Generating acc routine worker 
     Generating Tesla code 
     10, #pragma acc loop vector, worker /* threadIdx.x threadIdx.y */ 
     Loop is parallelizable 
main: 
    20, Generating copy(data[:16777216]) 
    21, Loop is parallelizable 
     Accelerator kernel generated 
     Generating Tesla code 
     21, #pragma acc loop gang /* blockIdx.x */ 
$ ./t1 
4 
$ 

Обратите внимание, что банда параллелизм была выполнена на внешнем контуре , а рабочий параллелизм выполнялся во внутреннем (рутинном) цикле.

Этот метод не зависит от динамического параллелизма (вместо этого он relies on a partitioning of parallelism between worker at the routine level and gang at the caller level) и не будет вызывать динамический параллелизм.

Внутреннее использование динамического параллелизма (CDP) не поддерживается в настоящее время в PGI 15.7. Должно быть возможным вызвать другие функции (например, CUDA или библиотеки), которые используют CDP из кода OpenACC, но в настоящее время, изначально не используется (и не поддерживается) в PGI 15.7

Answer 2

попытки заменить «#pragma акк параллельного цикла» с #pragma согла петлей»