Как переменная в памяти устройства используется внешней функцией?

В этом коде:

#include <iostream>

void intfun(int * variable, int value){
    #pragma acc parallel present(variable[:1]) num_gangs(1) num_workers(1)
    {
        *variable = value;
    }
}

int main(){
    int var, value = 29;

    #pragma acc enter data create(var) copyin(value)
        intfun(&var,value);
    #pragma acc exit data copyout(var) delete(value)

    std::cout << var << std::endl;
}

Как int value распознается в памяти устройства в intfun? Если я заменю present(variable[:1]) на present(variable[:1],value) в прагме intfun, я получу следующую ошибку времени выполнения:

FATAL ERROR: data in PRESENT clause was not found on device 1: name=_43144_33_value
 file:/opt/pgi/linux86-64/14.9/include/CC/iostream intfun__FPii line:5
Present table dump for device[1]: NVIDIA Tesla GPU 1, compute capability 3.5
host:0x7fffc11faa28 device:0x2303f20200 size:4 presentcount:1 line:14 name:_43152_14_value
host:0x7fffc11faa34 device:0x2303f20000 size:4 presentcount:2 line:14 name:_43152presentvar

Я не понимаю, почему указание, что value равно present, приводит к ошибке выше. Я проверил с помощью NVVP, что value копируется только один раз в директиве enter data, т.е. не копируется повторно в директиве parallel в intfun. Как OpenACC творит чудеса?


c++ cuda gpgpu openacc pgi
person lodhb    schedule 23.09.2014    source источник

Ответы (1)


arrow_upward
4
arrow_downward

Вы снова запутались в собственном синтаксисе и в том, что вам уже указывали.

value в intfun отличается от value того, что вы сделали copyin(value) в main. Вызов функции передает value по значению, то есть создает его копию. Поэтому добавлять его в предложение present() не имеет смысла, поскольку его нет на устройстве. Компилятор должен скопировать его. (А когда вы вообще не упоминаете об этом, компилятор автоматически распознает, что он нужен, и скопирует его для вас.)

Элемент, который присутствует на устройстве, — это переменная main области видимости value, которая не используется устройством intfun. Присвоение всем вашим переменным одинакового имени, вероятно, не поможет вам понять это.

Чтобы продемонстрировать это, давайте передадим value по ссылке вместо по значению:

$ cat main8.cpp
#include <iostream>

void intfun(int * variable, int &value){
    #pragma acc parallel present(variable[:1],value) num_gangs(1) num_workers(1)
    {
        *variable = value;
    }
}

int main(){
    int var, value = 29;

    #pragma acc enter data create(var) copyin(value)
        intfun(&var,value);
    #pragma acc exit data copyout(var) delete(value)

    std::cout << var << std::endl;
}
[user2@dc12 misc]$ pgcpp -acc -Minfo main8.cpp
intfun(int *, int &):
      5, Generating present(variable[:1])
         Generating present(value[:])
         Accelerator kernel generated
         Generating Tesla code
main:
     14, Generating enter data copyin(value)
         Generating enter data create(var)
     17, Generating exit data delete(value)
         Generating exit data copyout(var)
$ ./a.out
29
$

Теперь переменная основной области видимости value находится на устройстве, компилятор ее знает, она будет использоваться intfun напрямую, и добавление ее в предложение present является законным и функциональным.

person Robert Crovella    schedule 23.09.2014
comment
Тогда я не понимаю, почему в моем исходном примере value не копируется на устройство, если я опускаю copyin(value) и delete(value) в прагмах в main. Откуда устройство получает значение 29? Это хранится в памяти хоста и поэтому должно быть скопировано в память устройства в какой-то момент, верно? - person lodhb; 24.09.2014
comment
Я уже говорил в своем ответе, что если вы пропустите какое-либо упоминание о value, компилятор автоматически распознает, что value требуется ядру, и автоматически сгенерирует для него необходимую копию. Только когда вы указываете present, а затем переопределяете это поведение, все ломается. Если у вас есть новый вопрос, то есть что-то еще, что вы не понимаете, основанное на коде, отличном от того, что показано в вашем вопросе, вы можете задать новый вопрос. copyin(value) в main не имеет отношения к исходному коду, и исходный код прекрасно работает без него. - person Robert Crovella; 24.09.2014
comment
Я проверил с помощью NVVP, что отсутствие явного указания copyin из value в main приводит к отсутствию передачи данных от хоста к устройству. Поэтому я озадачен тем, откуда устройство получает значение value. - person lodhb; 24.09.2014
comment
Его можно передать как параметр ядра. - person Robert Crovella; 24.09.2014
comment
Спасибо за этот ответ, он помог мне кое-что понять. Я хотел отметить, что когда вы говорите то, на что вам уже указали, это относится только к одному из многих людей, которым может быть полезно то, что вы написали. - person Richard; 23.10.2017