Подождите...Каталог
Новые концепции для CUDA. Часть 1Кэш параллельных данных (parallel data cache - PDC), также известный как "совместно используемая память", является одним из основных элементов архитектуры G80, который позволяет работать с более эффективными примитивами синхронизации локальных данных. Похоже, что это основной элемент архитектуры NVIDIA для ускорения работы геометрических шейдеров DirectX 10. Перед тем как идти дальше, посмотрите на архитектуру NVIDIA G8x и CUDA. Эта блок-схема очень похожа на одну из диаграмм из патентов NVIDIA, зарегистрированных в 2003 году.
В G80 варианте этой архитектуры на каждый мультипроцессор приходится 8 процессоров (ALU-арифметико-логических устройств) и 16 мультипроцессоров на чип. Кроме того, как видно на блок-схеме, у каждого процессора есть минимум один регистр. Опять же, на блок-схеме можно заметить, что на каждый мультипроцессор приходится 1пул общей памяти (shared memory). А общаться между собой разные мультипроцессоры могут только через память устройства (device memory). К тому же, нет никаких естественных примитивов, чтобы упростить эту ситуацию. Вот почему мы называем возможности синхронизации G80 как "локальные" – они не распространяются на весь чип. Однако с другой стороны, это невероятно эффективно для имеющихся возможностей. Так что же из себя представляет PDC? Известно, что каждый блок общей памяти состоит из 16 банков однопортовой SRAM. Каждый банк имеет объем 1 KiB (1024 байт) и пропускную способность 32 бита на такт. Так как у G80 всего 16 мультипроцессоров – всего это 256 KiB памяти с пропускной способностью более 675 GiB/s. Для всех задач и целей это можно рассматривать как логическое и очень гибкое расширение регистрового файла. В дополнение к помощи в синхронизации, кэш параллельных данных всегда может сохранять имеющуюся пропускную способность. Принцип этого сценария, похоже, такой же, как у ячейки, куда Вы вручную загружаете данные и затем их используете много раз, минимизируя доступ к DRAM. Ячейка в несколько раз меньше, но фундаментальная идея та же самая.
Нужно заметить, что кэш параллельных данных поддерживает связь внутри группы из 16 тредов, без какой-либо явной синхронизации. Это можно назвать "суперлокальная неявная синхронизация", и это очень похоже на архитектуру пиксельных шейдеров, которая позволяет выполнять ddx и ddy команды с высокой пропускной способностью и большой скоростью. В большей группе тредов тоже не трудно организовать синхронную работу и передачу данных, но для этого уже нужна точная синхронизация. Поэтому давайте посмотрим, как это работает, и какие имеет ограничения. [N4-Новые концепции для CUDA. Часть 2] Давайте посмотрим, как NVIDIA описывает среду программирования CUDA. Довольно наглядно это показано на следующем рисунке:
Термин 'kernel' (ядро ОС) взят от потоковых процессоров. CPU делегирует "ядро" на GPU, и оно там разбивается на сети, блоки и треды. Кроме того, на этом рисунке не показано, что треды (Thread) в блоках группируются в "четверки". На первый взгляд это может показаться нагромождением, но в действительности это не так. Четверки соответствуют группе тредов, которые работают синхронно, то есть в идеале в пределах одной четверки ветвление полностью когерентное, что обеспечивает максимальную производительность. Помните, что когерентная ветвь у G80 состоит из 32 тредов. Четверки также группируются в блоки. Идея здесь в том, что одиночный блок будет гарантированно исполняться на том же самом мультипроцессоре, что позволяет тредам в пределах одного блока поддерживать связь и синхронизироваться друг с другом. Программист может задавать количество тредов на блок – это влияет на допустимые задержки и максимальное число доступных регистров, а также на объем доступной общей памяти. Вернемся к этому чуть позже. После этого все блоки группируются в одну сеть, которая объединяет все треды, выделенные GPU под одно программное ядро. Программа GPU в любое время имеет доступ к уникальному идентификатору задействованного блока и треда. Если Вы хотите, чтобы одновременно работали несколько программных ядер (что действительно нужно в мульти-GPU конфигурации), нужно, чтобы CPU работал в многопоточном режиме. СинхронизацияСейчас вы, может быть, думаете, что для организации работы примитивов синхронизации потребуется уйма переменных и функций. Это не так – все это делается вызовом единственной функции, у которой даже нет параметров! __syncthreads() просто позволяет Вам устанавливать точку синхронизации – барьер, и до того, как все треды блока не закончат выполнение всех предыдущих инструкций, ни один последующий код не будет выполняться (блокирование кода, относящегося как к общей памяти, так и к памяти устройства). Это позволяет эффективно предотвратить всевозможные отказы общей памяти "чтение ещё до записи", а также памяти устройства, если другие блоки не работают с теми же самыми ячейками памяти. Теоретически, эта команда почти не заметна; она занимает не больше времени, чем любая другая одиночная скалярная операция (2 цикла для 32 тредов), но здесь есть небольшая загвоздка. Чтобы понять проблему, представьте себе, что только арифметико-логические устройства имеют около 10 ступеней (с точки зрения шедулера это 5 ступеней, так как он работает на половинной частоте). Когда Вы запускаете синхронизацию, Вы хотите сделать так, чтобы все треды были в одной и той же точке процедуры. Это означает, что вы сдвигаете конвейер! При сдвиге конвейера Вы теряете минимум столько рабочих тактов, сколько стадий в этом конвейере. И это даже без учета задержек, что ещё больше усложняет ситуацию. Если, например, считывание данных из памяти для одной из четверок не закончилось, то все другие четверки для поддержки синхронной работы будут ждать окончания этой операции! Понятно, что это плохо. Можно ли эту проблему полностью устранить? Нет. Можно ли устранить её хотя бы частично? Да. Основная идея в том, что если у вас для этой операции задействованы 2 блока на один мультипроцессор, можно надеяться, что у Вас ещё достаточно много других тредов, чтобы обеспечить приличное быстродействие, в то время как другие блоки засинхронизированы. Если, например, у вас работает 8 блоков на одном мультипроцессоре и синхронизации достаточно редки, то есть практически никогда не совпадают по времени, для того чтобы спрятать задержки памяти, у Вас все еще имеется 87,5% тредов. Так это работает для локальной синхронизации. Если Вам "без вариантов" нужна глобальная синхронизация, помните, что Вы можете просто распределить нагрузку на несколько программных ядер и заставить CPU осуществлять синхронизацию. Если в каком-то конкретном случае это сделать нельзя, или это будет неэффективно, то треды могут коммутировать друг с другом через память устройства, что очень похоже на то, как если бы задача решалась с помощью CPU. Только убедитесь в том, что Вы знаете, какие блоки в данный момент в вашей архитектуре GPU работают, так как число мультипроцессоров может быть разным! Методики предотвращения сбоев "чтение потом запись" в этом случае такие же, как и для CPU, хотя у последнего может быть задействован общий кэш L2. У CUDA вообще нет кэша с чтением/записью, поэтому для организации общей синхронизации нужно учитывать многочисленные задержки памяти. Очевидно, что в идеале нужно иметь быстродействующую локальную синхронизацию и глобальную синхронизацию, организованную на уровне кэша второго уровня, но, похоже, что с G80 мы этого ещё не имеем. [N5-Что в сумме] Теперь, когда мы знаем, что из себя представляет CUDA и что она может, давайте подведем краткие итоги:
|
|||||||
Источник: www.beyond3d.com
