CUDA, весна 2009, 01 лекция (от 24 февраля)

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

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

Кроме Б и В будет ещё несколько лиц.

Цель следуюшая: давайте примерно посмотрим, что происх. с тактовыми частотами процессоров. С ростом очень плохо. В районе 2005 года был P4 4ГГц, который умельцы разгоняль до 7—8 мегагерц. С тех пор пошёл резкий скачёк вниз и потихоньку поднимаемся снова. Ждать того, что завтра выйдет новый чип на 10 ггц не надо, не выйдет. Есть несколько весьма серьёзных моментов:

• Как известно, энергия связана с тактовой частотой как четвёртая сетпень: увеличивая частоту в 2 раза, тепловыделение увеличивается в 16. Просто так увеличивать тактовую частоту бессмысленно. Как с ним боролись раньше? (первые процессоры были порядка 5 мгц) борьба шла очень просто --- мы уменьшали размеры элементов процессора, переходили на след. техпроцесс. Сейчас интел переходит на техпроцесс 32 нанометра. Если уменьш. размер элеметнов, то снижается энерговыделение, и при этом возр. скорость, поск. сигналу надо пробегать меньшее расст. Но есть неск. огранич, которые меньшают дальше уменьш. размер.

Рассм., как делают микросхъемы: это набор слоёв, из которых мы вытр. детали. Сущ. жлментов явл. фотолитография: строятся макси, которые наносятся на кристалл, и потом исопльзуем технологию похожую на проявку, и снимаем слои. Постепенно размер элементов прибл. к длине волны, и мы сталкиваемся с волновой оптикой. Со всеми её нелинейными эффектами.

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

Посмотрем примеры на параллельные архитектуры. Рассмотрим, как устроен C2D. У каждого ядра кэш первого уровня, общий кжш второго, контроллер и память. Тут есть интересный момент: представьте, что оба ядра эффективно работают, и модиф. один и тот же участок памяти, пусть даже два близко расп. ядра. Например, оба ядра прочли знач. одной переменной. Это значение легло в кэжши L1, L2. Потом первое ядро решило, что мы в эту переменную что-то запишем. В этот момент нужно сообщить в кэш L1 другого процессора, что старое знач. невалидно. То есть, должен быть единый целостный образ памяти. Т. о., необходимо обесп. синхронизацию кешей.

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

Более интересная архитектура: SMP с CC. То, что мы должны отслеживать кэш, ограничивает количество процессоров.

Другой пример --- архитектура Cell. У него один двухъядерный процессор общего назначения (PowerPC) и 8 спрециализированных процессоров, которые зависимы. Каждый SPE имеет 256 килобайт памяти прямо на чипе. Эта память, как следствие очень быстрая, поскольку обр. он к ней, минуя раздичные контроллеры.

BlueGene/L. Тут строится топология. В качестве элемента исп. Power 370, два таких процессора объединены L3-кэшем. Они уже соединяются в топологию, в частности, в трёхмерный тор.

Рассмотрим сначала двумерный тор. Построим сетку из процессоров и соединим процессоры на противоположных сторонах. Это уменьшает длину пути от одного проуессора до другого. Аналогично строится трёхмерный тор.

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

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

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

О чём мы будем говорить горраздо больше дальше. На самом деле, GPU тоже можно рассм. как MPP. бОЛЕЕ ПОДРОБНО, КАК УСТРОЕНА АРХИТЕКТУРА gpu g80, рассмотрим на слоедующих лекциях, сейчас в общем.

Есть набор потоковых процессоров (16—30), независимых. Каждый устроен след. образом: у него есть 8 выч. юнитов, при этом каждый их наз намного проще, чем то же ядро в core 2. При этом они выполняют одну команду вместе. При этом у них прямо начипе живёт его собственная shared-амять, регистры и кэш.

Вопрос: где находится врешинный кэш? Отвте: в рамках данной лекции они не рассмтариваются, поскольку они не доступны при работе с GPU как с массивным мультипроцессором.

В 60-е годы была предложена классификация архитектур (классфикация по Флинну).

Классическая инструкция — SISD. Есть такая вещь, как MISD, но она редко встречается, а вот SIMD и MISD широко распространены.

Если посмотреть на граф. процессоры пред. поколений, то увидим, что они близки к SIMD. С другой стороны, есть вариант MIMD, когда есть мгного проц. элементов, каждый вып. свои команды со своими данными.

Архитектура SIMD в определённом смысле проще MIMD.

Несколько замечаний: если формально CPU — SISD, но тем не менее, есть несколько вешей: есть Multithreading, есть SSE.

При этом надо понимать, что дополниетльные возможности процессора, такие как SSE и MT, должны использоваться программой, иначе они эффекта не дадут.

Рассмотрим пример многопоточной программы. Процедура вывода "Hello world" запускается в новой нити и в исходном процессе.

CUDA. Что это такое: compute unified device architecture. Это архитектрура, вкл. полный набор всего. Это расширение языка С, которое позв. заниматься массивно-пар. выч. Понятно, что изн. оно затачивалось под G80, хотя может работать и под другими арзитектурами. И CUDA это некое решение, которое позволит вам исп. её. Даже если нет видеолкарты, то можно посмотреть её в режиме эмуляции.

Есть ряд огранчений, например, ограничение по скорости работы памяти. Собственно, кэш это попытка бороться с огр. памяти.

__global__ --- директива, которая гворит nvcc, что функция будет выполняция на каждой нити. Она будет запущена в N (=1048576) экземплярах, каждая из которых выполняет эту функцию.

Поскольку мы запускаем нити на GPU, то должны выделить память на GPU. Для этого есть функция cudaMalloc, в параметре указывается, какое количество видеопамяти необходимо.

Для запуска ядра исп. специальный синтаксис: сначала идёт имя функции (оно должно быть __global__), дальше идут параметры ого, как необх. запустить эти нити, дальше идёт арг., который преедасться каждой нити.

Мы говорим, что хотим создать N/512 блоков нитей по 512 нитей в блоке.

С нашё точки зрения запустился миллион нитей, каждая выполнила свою функцию. Надо как-то отличать, в какой мы нити внутри нити. Для этого есть несколько переменных: threadIdx --- номер нити в блоке, blockIdx --- номер блока, blockDim --- размерность блока.

После того, как ядро посчиталось ...

cuda.cs.msu.su

Будет примерно 11 человек.

Ечсли останется не больше 8 человек, то можно будет разбираться в классе. В противном случае можно сдавать задания, делаете все 5. Есть альтернатива: можно кусок кода переписать на CUDA, подг. презентацию, показать, получить отлично. При этом необх. предупредить заранее.