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

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

В прошлой лекции была допущена неточность: разные варпы выполняются независимо, иначе бы всё теряло смысл. Поэтому это SIMT.

О чём сегодня поговорим:

Лектор рад сказать, что народу по-прежнему много, но вакантные места ещё есть, по сравнению с прошлым разом.

Рдин из самых важных моментов в cuda, как при написании, так и при оптимизации — работа с памятью, что тормозит работу везде, и на cpu, и на gpu. В цп это решается большим количеством кэшей. В gpu разница во времени доступа к своей памяти и к dram может отлю на два порядка. При этом, как упоминалось, с кэшем довольно плохо. Какие типы памяти есть в cuda (память, есть её название, далее вид доступа, как она выделяется, на каком уровне осущ. владение, и какова пропуск. способность):

• Регистры. Расп. непоср. на мультипроцессоре, к ним имеется полный доступ, выделяются они отдельно на каждый тред, причём выделение делается на этапе компиляции, длинамически деление менять нельзя, регистры одной нити невидимы другой. Можно запустить nvcc, и онвыдаст ассемблер, во что компилир. часть, выоплн. на gpu. Это такой ассемблер, там с первого-второго раза понятно, там можно посмотреть, сколько регистров требует ваше ярдо. Поск. регистры на чипе, скорость доступа практич. мгновеннная.
• Локальная память. Обычно необх. в ней возн., когда не зватает регистров. Она выделяется компилятором, но в dram, и скорость к ней очень низкая, обр. в ней не кэшируется. Поэтому иногд бывает полезно откомпилировать код в ассемблер, посмотреть, сколько реально регистров исп., и из этого можно понять, сколько регистров можно исп. в блоке.
• Shared память. К ней доступ на чтение и запись, выдялется на блок, и доступна на чтение и запись всем нитям блока. Поск. в блоке может бежать неск. сот нитей, и что произойдёт, если все они захотят писать в одно место памяти, целиком на вашей совести. Надо понимать при работе с shmem, что другие ниити тоже с ней работают, и когда они к ней обр., в общем случчае неизв., и нужно избегать конфликтов. Классич. пример: чтение из локальноцй памяти в шаред, и потом работа с прочит. данными. Чтобы обесп. доступ к шмем тдолько после чтения, используется синхронизация. Shared память расп. прямо на мультипроцессоре, есть тенденция, что объём её растёт. Скорость очень высокая, в идеале как к регистрам. Распред. shmem: заранее говорите, сколько нужно, то есть на этапе компиляции.
• Глобальная память. dram. Память, расп. вне чипа gpu, но соед. с ним. Можно в неё читать и писать, кроме того, в cuda, начиная с более поздних моделей gpu пояавл. атомарные способы работы с глобальной памятью. Пример: счётчик, который необх. инкрементировать. Скорость по ср. с ост. памятью низкая, опять два порядка по ср. с on-chip памятью. Но, при этом, её зато очень много, сколько на gpu стоит, столько можно использовать, она не кжшируется
• Есть ещё два типа памяти, которые расп. в dram. Это два типа памяти, которые можно только читать, но доступ к ним идёт через кэш. Модуль tex, который рассм. ранее, отвечате за доступ к этим типам памяти. За счёт того, что доступ кэшируется, скорость довольно высокая.

Сегодня в основном поговорим про глобальную память. След. лекция посвязена shared-памяти. Есть способы повыш. произв. за счёт применения спец. паттернов доступа.

Итак:

• самая быстрая шаред, самая медленная глобальная
• Можно исп. кэширование (например, в случае image processing)
• Варп разбивается на две половинки, и доступ к памяти осущ. независимо для каждой половины варпа

Чтобы получ. высокую скорость доступА, необх. исп. опр. шаблоны доступа, которые накл. огр. на то, как вы лезете в память. Рассм. этих огр. делается тлько по варпам. Если в пределах полуварпа всё нормально, то варп нормально обр. к памяти.

Краткое резюме: просто за счёт правильной работы с памятью можно повысить в разы, десятки раз скорость работы. Важно исп. shared-память. Сущ. спец. паттерны доступа, которые гарантируют то, что доступ произойдёт опр. доступом. Паатерны меняются с поколением gpu, самые жёсткие огр. для первых поколений, для последних самые слабые. Память работает независимо в рамках полуварпа.

Рассмотрим примеры. Особой оптимизации не будет, рассм. чуть более сложные примеры на cuda.

Пусть есть две матрица N×N, для удобства считаем, что N кратно 16. Эти матр. необх. расп. в глоб памяти. Для этоо необх. аллоц. глоб. память и скопировать из памяти хоста. Если мы считаем произв. матриц, то на каждый элемент произв. запустим одну нить, одна нить будет считать ровно один элемент произведения. Пусть блок нитей будет квадратным и состоять из 16×16 нитей. На след. матрице станет ясно, почему такая организация будет наиболее подходящая, сейчас лектор предлагает поверить на слово. Grid тоже организуем двумерно. То есть, если есть исх. матрица, разб. её на клеточки 16×16, и так считается. Каждая клеточка считается независимо.

Рассмотрим самый простой способ. Ядро: берём на вход указатели, где наход. матр. a, b, их размер и ук., куда писать результат.

Поск. матр. расп. в памяти линейно, то нам необх. эффективно считать адреса элементов. Для матр. a, где нужна строка, считаем нач. смещ. и бежим по ней, для матр. b, где нужен столбец, чуть хуже.

Понятно, что на умн. при итер. по столюбцу можно сэкономить, но проблема не в этом. На что необх. обр. внимание: что нам нужно для выч. элемента: 2*n чтений из глоб. памяти и 2*n вычислений. Из такого баланса понятно, что огр. будет доступ в память.

как нам это ядро запустить на выполнение: сперва нужно выделить память на девайсе под массивы.

С чем связана передача укащателя на результат в параметре: все функции cuda возвр. всегда знач. одного и того же типа: код ошибки. В примерах из sdk любят заворачивать все вызовы в макрос, который в случае ошибки печатает диагностику.

Вопрос: если поставить неск. видеокарт, как из идентифицировать, они же не явл. логич. единицей?

Ответ: есть спец. api, которые возвр. разл. информацию, в том числе количество девайсов и их параметры. Далее можно задавать, какой девайс является активным. Можно сначала поставить констр., что сейчас работаем с dev0

Вопрос: то есть, работа с девайсами последовательная?

Ответ: нет, большинство вызовов асинхронные. На саомм деле, ближе к концу курса запланирована лекцйия, как работать с неск. gpu

Обр. внимание, что у memcpy, в отл. от обычного вызова добавился ещё один параметр: откуда и куда копируем (host to dev, dev to host, dev to dev и host to host).

После этих двух вызовов скопировали две матрицы в память gpu, и всё, что осталось — запустить ядро, что мы и делаем.

После этого копируем обр. с девайса на хост и осв. память.

Тут есть вызов thread sync., который обесп. синхр. потока с gpu. Это один из таких способов. В начале в неких подозр. местах его можно ставить часто.

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

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

Вопрос: можно ли по ходу выполнения скидывать результат в память

Ответ: нафиг? Вы получите массу проблем с тем, что готово, что нет. Можно теоретически запустить много ядер, но проще от этого не будет, выигрыш не получите. Лезть во внутр. детали, каким обр. идёт scheduling, не нужно. В cuda так не принято, но формально так сделать можно. В cudaa есть понятие потока, где можно объед. асинхр. операции в цепочку, можно созд. несколько stream, и опрашивать их, где готово. Не надо плодить сущности без надобности. Есть высказывание, что самая большая оптимизация в программе — когда она впервые заработает, перейдёт из сост. нерабочего кода в сост. рабочего.

Итак, давайте посм... обр. внимание, что операции с глоб. памяти довольно простые. Есть три осн. функции: malloc, memcpy, free, которые позв. цп выделить память на gpu, скопировать что-то, и освободить глоб. память.

Вопрос: зачем нужен последний параметр в memcpy, если понятно, что копир. с хоста на девайс.

Ответ: как, это же просто указатели, числа? Тенденция сэкономить на буковках приводит к тому, что получается нечтиаемый и неподдерж. код. Есть язык Smalltalk, где очень длинные идентификаторы, но зато всё понятно.

Это общий некий скелет простейш. программы на cuda, что-то считающей: заводите массивы, копируете их, что-то считаете, копирукете обратно результат.

Краткий итог рассм. программы:

• 2×N арифм. операций, что для устр. с большик кол-вом АЛУ очень мало, и 2×N обр. к памяти, что очень много
• Тормозит обр. к памяти, memory bound

Что самое непр., обращаемся мы к одним и тем же участкам памяти. Если возьмём два соседних элемента, то окажется, что половина булет общая. Оптимизация закл. в том, чтобы попытаться, раз уж мысчитаем блоком вот такого (16×16) размера, понятно, что чтобы посчитать этот блок, то для него нужна полоса в 16 строк и плоса в 16 столбцов. Понятно, что если мы их щагрузим в shred, то быстр. повысится. То есть, сначала хорошо бы их считать в память, потом считать-считать, и потом один раз записать рез-т. К сожалению, не всё так просто, поск. если n большое, то оно может и не влезть.

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

Ображения идут через 32-64-128-битовые слова. отсюда следует, что при обр. к t[i] его размер должен быть развен 4/8/16 байтам и оно должно быть выровнено по данной границе.

Что произойдёт, если элемент не выровнен: будет две команды на чтение.

Что хорошо: вся выделяемая память выравн. по 256 байтам.

Как это используется: Допустим, у нас есть структура из трёх 4-байтовых элементов. Тогда размер у неё будет 12 байт, и через одну оно читаться будет плохо. У структуры можно указывать __align__(), при этом увеличился размер на треть, но скорость доступа выросла.

Есть такое понятие, как device compute capability, чтобы не надо было каждый раз думать: вот gf8800 и gf9600, что между ними разного: это два числа: основная версия точка дополнительная. Первое число — основная архитектура ядра (пока не меняется и равен 1), второе число — revision number, которое характеризует незначительные изменения. Соответственно, есть целый ряд фич, которые указан по compute capability, напимре, что есть в 1.1 и выше. CC показывает глобальные изменения, которые появились.

Вопрос: есть ли средства программной эмуляций compute capability старших версий на девайсах с младшей версией

Ответ: насколько известно лектору, нет.

Ранее приводилась ссылка на табличку (в CUDA Programming Guide), в которой указывалось, какая версия СС у какого девайса.

Паттерны доступа к памяти тоже зависят от СС.

Что очень хрошего есть в работе с глоб. памятью: в gpu есть такая полезная фича, как coalescing. Обычно потоки в варпах ломаятся запросами к памяти толпой. При опр. условиях gpu имеет их склеивать, когда, например, ломится полварпа, тогда они склеиваются в один. Если вы правильно сумели его использовать, вы получили выигрыш в доступе к памяти. Как оно работает: поск. запросы к памяти об. независимо для каждого полуварпа, это происх. независимо для каждого полуварпа (то есть, у одного могут объед., у другого — нет). Для того, чтобы это произошло, должны быть вып. минимум два условия:

• Длина блока должна обр. 32/64/128 байт
• Этот блок должен быть выровнен по своему размеру

Плюс накл. дополн. огр., которые привзяаны к СС

• Для 1.0 и 1.1: каждая нить либо обр. к 32-битовым словам, давая 64-байт. блок, илибо е 64-битовым, давая 128Б. Каждая нить обр. к своему слову: 0-я нить к слову в начале блока, вторая к след. и так далее. При этом, если некоторые нити чтение не делают, можно сделать фиктивное стение и свести ситуацию к передыдущей.

Вопрос: а coalecing для записи?

Ответ: тут проблем нет, поск. послали запрос. на запись и считаете дальше, тормозит именно чтение, поск. пока не прочитаете зн., дальше работать не можете.

Вопрос: а если записать, а потом прочитать?

Ответ: для разных варпов неизвестно что, для одного варпа вероятнее всего корректно.

Если у нас эти правила для устройств с CC 1.0/1.1 не вып. эти усл., то coalescing не будет.

Условия для устр. 1.2/1.3: cbnefwbz pfvtnyj ekexibkfcm^

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

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

В чём отлич., если хотяодно из усл. не вып.:

• 1.0/1.1: 16 отдельных транзакций
• 1.2/1.3: Пытается отдельные запросы объед. в меньшие блоки (2,3,...) и для каждого идёт отдельная транзакция

В этом смысле устр. с большим СС оказываются более гуманными. Кроме того, порядок обр. в блоке не важен. (в отл. от 1.0/1.1)

Вопрос: кто делает это объед? Контроллер?

Ответ: нет, контр. только вып. запросы, объед. делается жедлезно раньше.

Простые, известные грабли, показ. отл. арх. gpu и cpu: предст., что у вас есть данные, которые орг. как структура, и есть массив структур, и задача осущ. чтение данных. На ЦП всё просто: читается линейка кэша, и получается ряд элементов массива, для ЦП подобный подход очень удобен, удобно орг. массив структур. Посмотрим то же на gpu: есть структура, есть массив структур, теперь предст., что внутри ядра нужно по очереди прочесть три компонента. Понятно, что coalescing происходить не будет. При том, что на cpu это работало бы очень эффективно: мы бы это загрузили в кэш при первом обращении и дальше всё б работало.

"Почему мы всё время наступаем на одни и те же грабли? Давайте скинемся и купим новые"

Давайте теперь разобъём на три массива (которые можно объед. в структуру. А можно и не объед.). Теперь coalecing происх. будет.

В первом случае для чтение 3 float нам понадобилось 48 обращений, во втором случае 3.

Для cpu подобный подход может быть неудобен: первое обр. поместит кусок массива ax, поэтому второе обр. попадёт мимо кэша и третье тоже.

Вопрос: а если выровнять и читать целую структуру:

Ответ: мы можемч читать максимум 8 байт за раз. 2 float объединить можно, больше уже нет.

Основной паттерн работы с глоб. памятью:

• Постараться польз. как можно меньше
• Размеры и выравнивание. Следить, чтобы это всегда было
• Обязательно смотреть, стараться, чтобы получился coalescing

Вопрос: по ходу вып. программы есть ли возм. посмотреть, есть ли coalescing? Есть ли профайлеры для gpu?

Ответ: Насколько лектору известно, нет, но, говорят, что-то готовится. Предп., что ядро простое, и в случае чего можно посчитать на бумаге. Набор специльных тулзов это хорошо, но ни одна тулза не заменяет мозг.

Рассмотрим итерат. методы решения СЛАУ. Зачем они вообще нужны, если есть ЛУ-разложение? В каких случаях итер. подход выгоден? Можно поставить по одной нити на каждый эл-т массива x, и тогда цена по времени будет O(N). Второй вариант, когда матр. сильно разрежена. Классич. пример, когда берём дифур и пишем разн. схему. При этом полученная матрица будет сильно разреж. и иметь рег. структуру, чаще всего диагональную. Если матр. разр. и структ. её известна, то цена итерации — O(N). Если не так, то какая она, нафиг, сильно разреженая?

Обычно важно, сколько итер. вып., и сколько времени занимает итерация. Обыно итер. занимает O(N), и за 10 итераций (если метод сходится) получается решение. В случае ЛУ-разлож. оно выполн. за O(N^2), что есть две больгие разницы.

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

Небольшой матем. шажок: что нужно, чтобы это фигня сходилась: пусть x* — точно реш. Тогда возьмём и всюду добавим и вычтем x*, тогда наша погрешность описывается след. уравнением: ... . Что нам даст сходимость, причём сходимость быструю: |\alpha| \times ||A|| < 1. Тогда получим очень хорошую сходимость (например, если на кажд. итер. точность удваивается). Для cuda гораздо лучше подходит подобное. Есть методы более сложны, есть просто идущий с cuda пакет, который много чего умеет делать с матр. и линейными уравнениями. Если что-то нужно считать, то лучше наверное сначала попробовать его. Тем не менее, лоргика довольно проста, можно для примера посм. ядро, чтобы сделать одну итерацию цикла. Как видите, ничего особо сложного нет. Обратите внимание, что coalescing здесь происходит. ЗА один шаг выполняется одна итерация. Дальше есть варианты: можно звать неск. раз ядро для повыш. точности, можно сделать пинг-понг: поменяв местами x1 и x0, добавить ещё один цикл, что позволит вып. две итерации за раз, единст., что понадобится, синхронизация нитей перед началом второго шага. Для более серьёзных расчётов лучше таки использовать cublas(??).