OpenACC суммы при каждом выполнении

Question

Почему следующий код:

#include <iostream> 

int main(int argc, char const *argv[]) 
{ 
    int sum = 0; 
    int *array; 
    array = new int [100]; 

    #pragma acc enter data create(array[0:100],sum) 

    #pragma acc parallel loop present(array[0:100]) 
    for (int i = 0; i < 100; ++i) 
    { 
     array[i] = 1; 
    } 

    #pragma acc parallel loop present(array[0:100],sum) reduction(+:sum) 
    for (int i = 0; i < 100; ++i) 
    { 
     sum += array[i]; 
    } 

    #pragma acc exit data delete(array[0:100]) copyout(sum) 

    std::cout << sum << std::endl; 

    return 0; 
} 

различных результатах при каждом выполнении?

$ pgcpp -acc -Minfo main.cpp 
main: 
     7, Generating enter data create(sum) 
     Generating enter data create(array[:100]) 
     Generating present(array[:100]) 
     Accelerator kernel generated 
     12, #pragma acc loop gang, vector(256) /* blockIdx.x threadIdx.x */ 
     7, Generating Tesla code 
    15, Generating present(array[:100]) 
     Generating present(sum) 
     Accelerator kernel generated 
     18, #pragma acc loop gang, vector(256) /* blockIdx.x threadIdx.x */ 
     20, Sum reduction generated for sum 
    15, Generating Tesla code 
    25, Generating exit data copyout(sum) 
     Generating exit data delete(array[:100]) 
$ ./a.out 
100 
$ ./a.out 
200 
$ ./a.out 
300 
$ ./a.out 
400 

В соответствии со стандартом OpenACC:

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

Казалось бы, что sum является не перераспределена и вместо используемых повторно (и увеличивается) при каждом запуске программы. Кроме того, оператор + в директиве reduction инициализирует редукционную переменную до 0, поэтому этого не должно произойти, даже если sum не были освобождены между выполнением.

я могу избежать этого либо с помощью copyin вместо create для sum в enter data директивы или установить sum = 0 в одной банде, один рабочий ядро:

#pragma acc parallel present(sum) num_gangs(1) num_workers(1) 
sum = 0; 

Но это не является удовлетворительным, поскольку она требует либо дорогостоящий хост для копирования данных устройства, либо запуск ядра. Почему моя программа ведет себя так?

Answer 1

Вы неверно истолковываете значение инициализации оператора сокращения. Ссылаясь на статью openACC specification, стр. 20-21:

Предложение о сокращении допускается на параллельной конструкции. Он задает оператор сокращения и одну или несколько скалярных переменных. Для каждой переменной создается частная копия для каждой параллельной группы и инициализируется для этого оператора. В конце области значения для каждой группы объединяются с использованием оператора сокращения и результат в сочетании со значением исходной переменной и сохраняется в исходной переменной.

Это означает, что общая проблема с сокращением разбита на кусочки, каждая деталь обрабатывается бандой. Часть проблемы, которую обрабатывает банда, будет использовать указанное значение инициализации для редукционной переменной. Однако, когда окончательный результат создается, индивидуальные результаты от каждой группы будут объединены со значением исходной переменной (sum в вашем случае), и это будет результатом.

Итак, вы должны правильно инициализировать sum, возможно, используя один из методов, которые вы изложите в своем вопросе.

Кроме того, хотя это не является основной проблемой, обратите внимание, что niether освобождение или распределение не влияет на содержимое памяти. Новая переменная, выделенная в этом месте, без правильной инициализации, берет значение в этом месте.