Ответ 1
Существует несколько способов решения проблем с пропускной способностью CPU-GPU - я надеюсь, что вы подразумеваете под латентностью, а не за латентность самой передачи. Обратите внимание, что я намеренно использовал термин адрес вместо сокращения, так как вам не обязательно уменьшать задержку, если вы можете скрыть его. Также обратите внимание, что я больше знаком с CUDA, поэтому ниже я имею в виду только CUDA, но некоторые функции также доступны в OpenCL.
Как вы уже упоминали блокировка страницы, сама цель увеличения. Кроме того, можно сопоставить память хоста с блокировкой страниц на GPU, механизм, который обеспечивает прямой доступ к данным, выделенным из ядра GPU, без необходимости дополнительной передачи данных. Этот механизм называется переносом с нулевой копией, и это полезно, если данные считываются/записываются только один раз, сопровождаемый значительным количеством вычислений, и для графических процессоров без отдельной памяти (мобильной). Однако, если ядро, обращающееся к нуль-скопированным данным, не сильно связано с вычислением, и поэтому латентность доступа к данным не может быть скрыта, блокировка страницы, но не отображаемая память будет более эффективной. Кроме того, если данные не вписываются в память GPU, нулевая копия все равно будет работать.
Обратите внимание, что чрезмерное количество заблокированных страниц может привести к серьезному замедлению на стороне процессора.
Приближаясь к проблеме под другим углом, как упоминал tkerwin, асинхронная передача (по отношению к потоку процессора, говорящему на GPU) является ключом, чтобы скрыть задержку передачи CPU-GPU путем перекрытия вычислений на CPU с передачей. Это может быть достигнуто с помощью cudaMemcpyAsync(), а также с использованием нулевой копии с асинхронным выполнением ядра.
Это можно сделать еще больше, используя несколько потоков для перекрытия передачи с выполнением ядра. Обратите внимание, что для планирования потока может потребоваться особое внимание для хорошего перекрытия; Карты Tesla и Quadro имеют двухъядерный двигатель DMA, который обеспечивает одновременную передачу данных на GPU и обратно.
Кроме того, с CUDA 4.0 стало проще использовать графический процессор из нескольких потоков ЦП, поэтому в многопоточном коде процессора каждый поток может отправлять свои собственные данные на графический процессор и запускать ядра легче.
Наконец, GMAC реализует асимметричную модель общей памяти для CUDA. Одной из его очень интересных особенностей являются модели когерентности, которые она предоставляет, в частности, ленивое и скользящее обновление, позволяющее передавать только данные, модифицированные на CPU заблокированным способом.
Подробнее см. В следующей статье: Gelado et al. - Асимметричная распределенная общая память
Модель для гетерогенных параллельных систем.