openacc - async - получение ускорения при передаче данных при использовании множества очередей асинхронизации

Question

Мой вопрос касается влияния передачи нескольких массивов в разных очередях асинхронизации между хостом и устройством.

Предположим, у нас есть четыре массива:

double *a, *b, *c, *d; 

И каждый из них был выделен с размером N.

a = (double*) malloc(N * sizeof(double)); 
b = (double*) malloc(N * sizeof(double)); 
c = (double*) malloc(N * sizeof(double)); 
d = (double*) malloc(N * sizeof(double)); 

Теперь мы можем передавать их между устройством и хостом следующим образом в один пункте:

#pragma acc enter data copyin(a[0:N], b[0:N], c[0:N], d[0:N]) async 
#pragma acc wait 

Или мы можем использовать много асинхронных положений и передавать их по разным очередям:

#pragma acc enter data copyin(a[0:N]) async(1) 
#pragma acc enter data copyin(b[0:N]) async(2) 
#pragma acc enter data copyin(c[0:N]) async(3) 
#pragma acc enter data copyin(d[0:N]) async(4) 
#pragma acc wait 

Результаты обоих вышеуказанных подходов являются одинаковыми. Однако с точки зрения производительности, второй, по-видимому, лучше в некоторых случаях.

Я сделал некоторые измерения и обнаружил, что для copyout ING и обновления хоста кажется, что использование more чем одной очереди лучше one с точкой зрения производительности.

Давай называть первый подходом one, а второй подходом more, и следующий подходом more_nonumber (не заметить никакого номера для пункта асинхронного):

#pragma acc enter data copyin(a[0:N]) async 
#pragma acc enter data copyin(b[0:N]) async 
#pragma acc enter data copyin(c[0:N]) async 
#pragma acc enter data copyin(d[0:N]) async 
#pragma acc wait 

Тогда здесь измерения для 10000 итераций (за исключением 100 первых и 100 последних из них, что приводит к средним 9,800 итераций между ними):

one

CopyIn: 64.273us

Обновление устройства: 60.928us

Обновление самостоятельно: 69.502us

CopyOut: 70.929us

---

more

CopyIn: 65.944us

Обновление устройства: 62.271us

Обновление самостоятельно: 60.592us

CopyOut: 59.565us

---

more_nonumber

CopyIn: 66.018us

Обновление устройства: 62.735us

Update самостоятельно: 70.862us

CopyOut: 72.317us

---

Средняя 9800 работает!

При использовании метода more для копирования (70.929/59.565) наблюдается ускорение в сравнении с one или 14% для самообслуживания (69.502/60.592).

Мой вопрос: Являются ли эти цифры законными? Можем ли мы полагаться на эти цифры?

Для вашего удобства я поставил свой код на github. Вы можете взглянуть на него.

Answer 1

Async наиболее полезен при перемещении данных и вычислении данных на устройстве. Следовательно, это упражнение немного постороннее, но я сделаю все возможное, чтобы объяснить, что происходит. Я должен отметить, что именно так PGI в настоящее время (v16.10) реализует «асинхронный» и не обязательно, как другие реализации OpenACC будут реализовывать «асинхронный».

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

Если вы используете свой код без каких-либо асинхронных предложений, вы увидите, что добавление 4 переменных в одну прагму значительно быстрее, чем у каждого из них. Причина в том, что с четырьмя прагмами хост ждет, пока последний буфер не будет отправлен, прежде чем перейти к следующей переменной. Когда все они находятся в одной и той же прагме, как только один буфер последнего массива запускает его передачу, среда выполнения может начать заполнять другой буфер данными следующего массива.

Когда вы добавляете «асинхронный» к единственной прагме, а затем ожидаете, вы не должны видеть разницу в производительности, не используя «асинхронный» или «ждать» вообще. Другими словами, это то же самое:

#pragma acc update device(a,b,c,d) 

#pragma acc update device(a,b,c,d) async 
#pragma acc wait 

При добавлении «асинхр» в 4-х отдельных прагмах, вы получите примерно такую ​​же производительность, как вы бы положить их все в том же прагме с CPU Безразлично 't ждать, чтобы начать буферизацию следующего массива.

Я не могу объяснить, почему копия назад (обновление self, copyout) с устройства происходит быстрее, когда каждый массив находится в собственной очереди async. Мне кажется, что это не имеет большого значения.

Для этого теста оптимальным было бы не использовать двойные буферы и вместо этого просто привязывать весь массив в физической памяти. Таким образом, вы сэкономите расходы на копирование из виртуальной в закрепленную память. Компилятор не может сделать это по умолчанию, так как физическая память ограничена и не может гарантировать, что память программы будет соответствовать. Для большинства кодов, где большая часть времени тратится на устройство, используя закрепленную память, не помогает, если вообще. Наконец, есть некоторые накладные расходы на производительность для освобождения закрепленной памяти (устройство необходимо синхронизировать, чтобы гарантировать отсутствие передачи в закрепленную память). Следовательно, наиболее полезно, если массив находится в одной области данных, но имеет множество обновлений.

Вот времена, используя возлагали память:

% pgcc -fast -acc -ta=tesla:cc35,pinned transfer.c -o tpinned.out 
% numactl -C 2 ./tpinned.out 
one 
CopyIn: 50.161us 
Update device: 49.679us 
Update self: 47.595us 
CopyOut: 47.631us 
--------- 
more 
CopyIn: 52.448us 
Update device: 52.135us 
Update self: 49.904us 
CopyOut: 47.926us 
--------- 
more_nonumber 
CopyIn: 52.172us 
Update device: 51.712us 
Update self: 49.363us 
CopyOut: 49.430us 
---------