CUDA, весна 2009, 04 лекция (от 17 марта)

Материал из eSyr's wiki.

• Диктофонная запись: http://esyr.org/lections/audio/cuda_2009_summer/cuda_09_03_17.ogg

Иерархия памяти CUDA. shared-память и эффективное использование.

Один из самых важных типов памяти — shared-память. Этот память, расп. на чипе, нах. в области потокового мультипроцессора. Доступ к ней, за искл. некоторых случаев, идёт так же быстро, как к регистрам. В нормальных условиях доступ практически мгновенен.

Память выделяется на уровне блока, и это определяет то, сколько блоков может одновременно выполняться на мультипроцессоре. Если у вас есть много блоков и мало мультипроцессоров, то есть желание неск. блоков запихнуть на каждый. Основное ограничение при этом — чтобы этим блокам было достаточно shared-памяти. Shared-память распределяется на этапе коспиляции. Тот кусок, которые вы выделили блоку, доступен всем нитям блока. Каждая нить может читать и писать, и ответственность за отстутсвие при этом проблем лежит на плечах программиста.

Резюмировать осн. свойства shred-памяти:

• Самая быстрая среди тех, которой можно управлять
• Сейчас всего 16К на мультипроцессор
• Память совместно исп. нитями блока
• Каждый warp рабю на две половинки, кажджая половинка обр. к памяти независимо. У двух полуварпов не может быть конфликтов по памяти.
• В силу ого, что доступна каждой нити, нкобходима явная синхронизация

Какой стандартный паттерн работы с shmem:

посмотрим, когда она нам нужна, когда даёт выигрыш — когда нужно одни и те же знач. перемалывать. Соответственно, в начале каждый процесс загружает свой кусочек в shmem. поск. нити выполн. не одновременно, то нужно быть уверенным, что каждый процесс загрузил свой кусок. Для этого используется явная синхронизация, __syncthreads(), и выполнение продолжается толдько после того, как все нити дошли за эту команду. Соответственно, после этого можно быть уверенным, что каждая нить свой кусок загрузила. После этого делаются вычисления над данными, при этом исп. то, что память быстрая, вып. некие вычисления, след. шаг — записать данные в глоб. память. Для того каждая нить сохзр. свой куок shmem в глобальную, после этого нужно вставить __syncthreads(), чтобы убедиться, что все данные к этому моменту сохранены.

Вот такой вот шаблон. Рассмотрим его на старом примере: умножении матриц. Мы считаем, что у нас есть две кв. матрицы N*N, при этом для простоты считаем, что N кратно 16. Мы устр. по нити на элемент произведения, блок имеет размер 16*16 нитей, сами блоки орг. также в двумерный grid.

Посмотрим на старую реализацию. Что мы делаем: мы опр. необх. адр. в памяти, после чего циклом считаем произведение. Чтобы посчитать произв., нам необх. две арифм. операции и два бр. к памяти. Понятно, что обр. к памяти по времени перевешивает все вычисления, более того, на 1 элемент нужно 2N обращения. Можно добиться coalescing, но это дорого. Понятно, что это можно улучшить. Посмотрим, в чём основная проблема нашей реализации. Если мы посмотрим на то, как мы вычисляли, то каждому блоку соотв. квадратная подматр. C', лдя её вычисления нам нужны 2 полосы — 16*N из А и N*16 из B. При этом обр. внимание на след. момент: когда мы будем считать C', мы будем многократно обр. к одним и тем же элементам из этих полос. Если два эл-та расп. один над другим, то у них будут один. жл-ты из А, елси рядо по горизонтали — из В. При том идёт обр. к глобальной памяти — удовольствие очень дорогое. Казалось бы, давайте загрузим в shmem и будем с ними работать. Понятно, что этот подход не подзодит, поск. для больших N мы даже одну полосму не поместим. Но, на самом еле, это нас не останавливает: обратим внимание на след. фактор — обр. внимание, что полосы можно побить на квадратики 16*16 и что мы получим: можно всю матр. С' предст. след. образом: берём квадр. из А, из В, и умножаем, потом прибавл. произв. следующих, и так далее. Таким обр., цикл от 0..N можно разбить на два цикла: цикл от 0..N/16 по квадратам и цикл 0..16 по каждому квадратику.

Понятно, что мы можем легко засунуть две таких матр. 16*16 в память. Следующий момент — при таком разбиении нам элементы из этого квадратика понадобятся только один раз (в пределах блока).

Тогда, что у нас получается: дост. внутри ядра в цикле вып. след. операции — загружаем по одному эл-ту в каждой нити (блок — 16*16 нитей, поэтому так мы загрузим сразу блок), после чего считаем произв., после чего загр. слде. квадратики и так далее. Единст., что необх. — выполнить синхронизацию.

Вот он, код ядра на cuda. Сначала мы запоминаем номер блока, далее идёт цикл по блокам. Каждая нить загр. один элемент, при этом польз. тем, что ост. нити тоже загружает свой жэлемент. В какой очке мы выделяем shmem (вне или внутри цикла) не важно, так как память распр. статически. После загрузки делаем __syncthreads(). Что дальше? Два блока обр. можно загружать следующие. Но перед этим необх. убедиться, что данные, которыем мы исп., более не нужны.

Это позволило нам сократить количество обр. к памяти в 16 раз.

Вопрос: __syncthreads() синхронизыет нити в каких рамках?

Ответ: __syncthreads() — синхронизация всех нитей блока. Синхронизации нитей грида нет. Елси необх. синхр. все нити грида, то проще разбить ядро на две части.

__syncthreads() — операция крайне дешёвая, факт., мы просто помечаем, что данный варп готов. В случае, если ждать ничего не надо, она занимает один такт.

Код стал чуть больше, но зато мы здорово сэкономили на бр. к глобальной памяти. То есть, это себя оправдывает.

Только за счёт эого получили 10-кратное приращ. производительности. Это типичный паттерн — мы строим блоки так, чтобы ему нужен был огр. объём данных, чтобы его можно было втянуть в shmem, что-то с ними делать, потом выплюнуть. Если все данные всосать не можем, то точно также делим на чатси и обр. по частям.

Мы эффективно исп. shmem для уменьш. кол. обр. к памяти.

Однако, shmem тоже несёт опр. ньюансы. Если в случае глоб. памяти была фича по объед. запросов, что позв. получить выигрыш, то здесь при непр. обращении налагаются штрафные санкции.

Shmem созд. с целью увел. пропуск. способности. Для этого было принято след. арх. решение: память разб. на 16 частей, наз. банками, при этом, за 1 такт можно осущ. 16 обр., по одному на банк. Что случится, если две нити лезут в одну банку? Происз. bank conflict, при котором запросы будут отработаны один за другим. BC это то, что способно понизить быстродействие кода. Самый худший вариант — когда все нити лезут в один банк. Для того, чтобы уменьш. вер. этого конфликта, было принято след. решение: как правило, соседние нити обр. к соседним ячейкам памяти. Поэтомы распр. по банкам не последовательные куски памяти (в этом случае были бы сплошные конфликты), было принято след. решение: всю память делим на ур. 32-битовых слов, и разд. 32-битовые слова по банкам. При этом что нам то дало: осн. типы данных, с к-рыми работает cuda — float/int (оба по 4 байта), и если 16 нитей полуварпа обр. к посл. куску данных, то каждая нить обр. к своему банку и конфликтов нет.

Есть понятие "порядок конфл." — сколько нитей лезут в один банк. Понятно, что у одного полуварпа может быть неск. конфликтов, и понятно, что важен конфликт с наиб. порядком, потому что он опр. замедл. в быстр.

Единст. ньюанс: если все обр. в одно слово, то происходит броадкаст (если меньше, то вроде нет).

Типичный пример — обр. в каждый второй банк. Типичный пример: a[2*i] + a[2*i+1].

Рассмотрим перемножение матриц 16*16 в плане конфликта банков. Доступ к одной матр. идёт по стр., к другой — по столбцам. По строкам всё очень сдорово, потому что каждая нить обр. своему банку. Что произ., если все 16 нитей обр. к одному столбцу. А вот тут ждёт непр. сюрприз — конфликт 16-го порядка, пос.к всё в одном банке. Для такого случая есть довольно простой хак, позв. получить и для обр. по строкам, и по столбца получить бесконфл. обращение. Давайте добавим к матр. ещё один столбец. Тогда соседние эл-ты в строке аккуратно лягут в 16 банков, то же в столбце. В итоге, мы изб. от конфл. как при доступе по строкам, так и по столбцам. Это на самом деле довольно расп. способ в CUDA. Это может иметь смысл и при работе с глоб. памятью.

Это довольно расп. приём, когда мы вводим фиктивные данные, но которые позв. обесп. удачный способ доступа к памяти.

Вопрос: почему нельзя матрицу повернуть?

Ответ: да, можно. Но иногда бывает, что нужно обр. и так, и так. Это тоже хороший вариант.

Последняя часть лекции будет посвязена довольно простой проблеме: есть массив, нужно найти его сумму. Или есть миллион эл-тов, нужно их просуммировать. В общем виде эта задача наз. задачей редукции. Есть массив и есть бинарная операция. ДАльше лектор будет писать для суммы, но можно заменить её на любую бинарную операцию.

Как на бычном С написать сумму массива: берём и суммируем. Как написать параллельно? Возникает интересный вопрос.

Есть уровен, который решается блоками и отдельными нитями. Поск. блоки взаимод. не могут, то что мы делаем: берём отд. массив, разб. его на куски и каждому куску сопост. один блок.

Что мы можем сделать внутри блока? Тут уже есть shmem, есть взаимод. нитей... Поэтому, первый шаг понятен: копируем данные в shmem. Далее нужно уметь эти данные просуммировать. Есть простой способ суммирования — иерарзическое суммирование, которое сумму преобр. в дерево, и сумма считается за log n. Поск. операция будет происх. внутри блока, то ln <размер блока>, иого операций — n/<разм. блока> ln <разм. блока>. Следующий шаг — собрать в массив поменьше и ещё раз применить эту операцию.

Как посчитать сумму, если есть пар. арх.: суммируем пары элементов, потом получ. суммы парами и так далее, пока не останется одно число.

Алгоритм ясен, остался вопрос, ка это реализовать на cuda. Часть вещей, которые сейчасб удут упоминаться, будут рассм. на лекции по оптимизации.

Первый этап: думать лектор не хотел, расписал просто по нитям. Самый тупой вариант. Лектор оговорится, что блоков хватает, у нас одномерный грид (по каждому разм. не более 16 бит; если не хватает, делаем грид двумерный, но тогда изм. способ адр.). Чем это плохо:

• Heavy branching. У же на первой итерации цикла мы не исп. половину нитей

Второй вариант: немного схитрим — каждая нить будет теперь получать два элемента. Код стал гораздо аккуратнее: теперь всместо номера нити исп. пороговое значение, так что бранчинг может идти по варпам. Но всё равно не хорошо:

• Посмотрим, к каким эл-там мы обращаемся. Вспомним про конфликт банков кратности 2. Он ровно здесь и есть. То есть, на первой итерации конфликт 2 порядка, на второй памяти — 4 порядка, и так до 16 на последних неск. итерациях.

Посмотрим, в чём засада: мы начинали сумм. сосед. эл-ты и расст. удваивали. Понятно, что при этом шли пост. конфликты. Попробуем сделать наоборот: начнём суммировать с наиб. удалённых, с тех, кто расп. на половине размера блока. При то кол. работающих нитей всё время уменьш. Но проблемы с бранчингом не будет, поск. с худшем случае деление ляжет внутрь одного ворпа. Понятно, что если кол-во элементов 2^n, то бранчинга не будет, кроме последних шагов, когда работает один варп.

Вроде бы при этом стало всё хорошо. За искл. одной проблемы: на первой итерации половина нтией простаивает.

Для испр. этого сделаем посл. шаг — назначим каждому блоку в два раза больше элементов, и дальше всё также.

Остался последний штрих: что будет, если s <= 32, оно стало меньше размера варпа. При этом условие (tid < s) понятно, что чётко выполнено, синхронизация уже не нужна (остался один варп), и можно ручками развернуть цикл. В cuda есть средства сделать полный анролл, когда всё известно, но этот вариант тоже работает.

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

Что являлось здесь ключевым: использование shmem — основную часть суммирование делаем в shmem. Далее — борьба с бранчингом и банк конфликтами. Последние шаги — более полная загрузка и разворачивание циклов.