
Lection-1.ppt
- Количество слайдов: 62
Программирование графических ускорителей CUDA Лекция 1
Зачем нужны многоядерные системы Посмотрим на частоты CPU: ◦ ◦ ◦ ◦ 1978 2004 2005 2006 2007 2008 2009 г. г. - 4. 77 МГц, Intel 8086 Pentium 4, 3. 46 GHz Pentium 4, 3. 8 GHz Core Duo T 2700, 2333 MHz Core 2 Duo E 6700, 2. 66 GHz Core 2 Duo E 6800, 3 GHz Core 2 Duo E 8600, 3. 33 Ghz Core i 7 950, 3. 06 GHz
Динамика роста производительности для CPU и GPU
Зачем нужны многоядерные системы z. Был сделан огромный скачок z. За последние 5 лет z. Минимальный рост частот z. Развитие многоядерных систем z. Видно, что одноядерные системы зашли в тупик
Зачем нужны многоядерные системы Причины тупика: ◦ Энерговыделение ~ четвертой степени частоты ◦ Ограничения техпроцесса
Зачем нужны многоядерные системы Таким образом, повышение быстродействия следует ждать именно от параллельности. На самом деле CPU уже давно используют параллельную обработку для повышения производительности ◦ Конвейер ◦ Multithreading ◦ SSE
Intel Core 2 Duo 32 Кб L 1 кэш данных и команд для каждого ядра 2/4 Мб общий L 2 кэш Единый образ памяти для каждого ядра необходимость синхронизации кэшей
Symmetric Multiprocessor Architecture (SMP)
Symmetric Multiprocessor Architecture (SMP) Каждый процессор zимеет свои L 1 и L 2 кэши zподсоединен к общей шине zотслеживает доступ других процессоров к памяти для обеспечения единого образа памяти (например, один процессор хочет изменить данные, кэшированные другим процессором)
Blue. Gene/L
Blue. Gene/L z 65536 dual-core nodes znode y 770 Mhz Power. PC y. Double Hammer FPU (4 Flop/cycle) y 4 Mb on-chip L 3 кэш y 512 Mb off-chip RAM y 6 двухсторонних портов для 3 D-тора y 3 двухсторонних порта для collective network y 4 двухсторонних порта для barrier/interrupt
Blue. Gene/L
Архитектура G 80 Массив из потоковых мультипроцессоров
Архитектура G 80
Примеры многоядерный систем z Необходимо обратить ваше внимание на следующие особенности: 1) Как правило вычислительный узел – достаточно маломощный процессор 2) Вычислительные узлы имеют свою оперативную память и свой кэш 3) Вычислительные узлы объединяются в более крупные блоки 4) Крупные блоки могут объединяться с целью наращивания вычислительной мощи
Виды параллелизма z. На уровне задач z. На уровне данных
Классификация
Классификация CPU – SISD ◦ Multithreading: позволяет запускать множество потоков – параллелизм на уровне задач (MIMD) или данных (SIMD) ◦ SSE: набор 128 битных регистров ЦПУ можно запаковать 4 32 битных скаляра и проводить над ними операции одновременно (SIMD) GPU – SIMD*
SIMD z. На входе поток однородных элементов, каждый из которых может быть обработан независимо z. На выходе – однородный поток z. Обработкой занимается ядро (kernel)
SIMD • Так каждый элемент может быть обработан независимо от других, то их можно обрабатывать параллельно • Можно соединять между собой отдельные ядра для получения более сложной схемы обработки
Эволюция GPU z. Voodoo - растеризация треугольников, наложение текстуры и буфер глубины z. Очень легко распараллеливается z. На своих задачах легко обходил CPU
Эволюция GPU: Шейдеры !!ARBvp 1. 0 ATTRIB PARAM TEMP pos = vertex. position; eye = program. local [0]; light = program. local [1]; mvp [4] = { state. matrix. mvp }; mv 0 [4] = { state. matrix. program [0]. inverse }; l, lt, et, lt 2, et 2; DP 4 DP 4 ADD lt 2. x, light, mv 0 [0]; lt 2. y, light, mv 0 [1]; lt 2. z, light, mv 0 [2]; lt 2. w, light, mv 0 [3]; et. x, eye, mv 0 [0]; et. y, eye, mv 0 [1]; et. z, eye, mv 0 [2]; et. w, eye, mv 0 [3]; l, -pos, lt 2; DP 3 MOV lt. x, lt. y, lt. z, lt. w, l, vertex. texcoord [1]; l, vertex. texcoord [2]; l, vertex. normal; l. w;
Эволюция GPU: Шейдеры z. Работают с 4 D float-векторами z. Специальный ассемблер z. Компилируется драйвером устройства z. Отсутствие переходов и ветвления
Отличия CPU от GPU z. Очень высокая степень параллелизма z. Основная часть чипа занята логикой, а не кэшем z. Ограничения по функциональности
GPGPU z. Использование GPU для решения не графических задач z. Вся работа с GPU идет через графический API (Open. GL, D 3 D) z. Программы используют сразу два языка – один традиционный (С++) и один шейдерный z. Ограничения, присущие графическим API
CUDA (Compute Unified Device Architecture) Программирование массивнопараллельных систем требует специальных систем/языков. Программирование ресурсов CPU ограничено ◦ Multithreading ◦ SSE ◦ Часто bottleneck – в пропускной способности памяти CUDA - система (библиотеки и расширенный C) для программирования GPU
CUDA “Hello World” #define N (1024*1024) __global__ void kernel ( float * data ) { int idx = block. Idx. x * block. Dim. x + thread. Idx. x; float x = 2. 0 f * 3. 1415926 f * (float) idx / (float) N; data [idx] = sinf ( sqrtf ( x ) ); } int main ( int argc, char * { float a [N]; float * dev = NULL; argv [] ) cuda. Malloc ( (void**)&dev, N * sizeof ( float ) ); kernel<<
CUDA “Hello World” __global__ void kernel ( float * data ) { int idx = block. Idx. x * block. Dim. x + thread. Idx. x; // номер текущей нити float x = 2. 0 f * 3. 1415926 f * ( float) idx / (float) N; // значение аргумента data [idx] = sinf ( sqrtf ( x ) ); } // найти значение и // записать его в массив z Для каждого элемента массива (всего N) запускается отдельная нить, вычисляющая требуемое значение. z Каждая нить обладает уникальным id
CUDA “Hello World” float a [N]; float * dev = NULL; // выделить память на GPU под N элементов cuda. Malloc ( (void**)&dev, N * sizeof ( float ) ); // запустить N нитей блоками по 512 нитей // выполняемая на нити функция - kernel // массив данных - dev kernel<<
Подход CUDA Исходная задача разбивается на подзадачи, которые можно решать независимо друг от друга. Каждая из этих подзадач решается набором взаимодействующих между собой нитей
SIMT (Single Instruction, Multiple Threads) z. Параллельно на каждом SM выполняется большое число отдельных нитей (threads) z. Нити подряд разбиваются на warp’ы (по 32 нити) и SM управляет выполнением warp’ов z. Нити в пределах одного warp’а выполняются физически параллельно z. Большое число warp’ов покрывает латентность
Программная модель CUDA z GPU (device) это вычислительное устройство, которое: y. Является сопроцессором к CPU (host) y. Имеет собственную память (DRAM) y. Выполняет одновременно очень много нитей
Программная модель CUDA z Код состоит как из последовательных, так и из параллельных частей z Последовательные части кода выполняются на CPU z Массивно-параллельные части кода выполняются на GPU как ядра
Программная модель CUDA Отличия нитей между CPU и GPU z Нити на GPU очень «легкие» z HW планировщик задач z Почти нулевые затраты планировщика z Для полноценной загрузки GPU нужны тысячи нитей x Для покрытия латентностей операций чтения / записи x Для покрытия латентностей sfu инструкций
Программная модель CUDA z. Параллельная часть кода выполняется как очень большое количество нитей z. Нити группируются в блоки (blocks) фиксированного размера z. Блоки объединяются в сеть блоков (grid) z. Ядро выполняется на сетке из блоков z. Каждая нить и блок имеют свой идентификатор z. Нить использует идентификаторы для определения с каким элементом работать
Программная модель CUDA
Программная модель CUDA z. Потоки в CUDA объединяются в блоки: y. Возможна 1 D, 2 D, 3 D топология блока x. Общее кол-во потоков в блоке ограничено x. В текущем HW это 512 потоков
Программная модель CUDA z. Блоки в CUDA объединяются в сеть: y. Возможна 1 D и 2 D топология сети
Программная модель CUDA z. Каждый блок целиком выполняется на одном SM z. Нити одного блока могут взаимодействовать между собой z. Через shared-память z. Через барьерную синхронизацию z. Нити разным блоков не могут взаимодействовать между собой
Расширения языка C z CUDA – это расширение языка C y [+] спецификаторы для функций и переменных y [+] новые встроенные типы y [+] встроенные переменные (внутри ядра) y [+] директива для запуска ядра из C кода z Как скомпилировать CUDA код y [+] nvcc компилятор y [+]. cu расширение файла
Расширения языка C z. Спецификатор функций Спецификатор Выполняется на Может вызываться из __device__ device __global__ device host __host__ host z. Спецификатор переменных Спецификатор Находится Доступна Вид доступа __device__ device R __constant__ device / host R/W __shared__ device block R/W
Расширения языка C z. Спецификатор __global__ соответствует ядру z. Может возвращать только void z. Спецификаторы __host__ и __device___ могут использоваться одновременно z. Спецификаторы __global__ и __host__ не могут быть использованы одновременно
Расширения языка C z. Спецификатор __global__ соответствует ядру z. Может возвращать только void z. Спецификаторы __host__ и __device___ могут использоваться одновременно z. Спецификаторы __global__ и __host__ не могут быть использованы одновременно
Расширения языка C Ограничения на функции, выполняемые на GPU: z. Нельзя брать адрес (за исключением __global__) z. Не поддерживается рекурсия z. Не поддерживаются static-переменные внутри функции z. Не поддерживается переменное число входных аргументов
Расширения языка C Ограничения на спецификаторы переменных: z. Нельзя применять к полям структуры или union z. Не могут быть extern z. Запись в __constant__ может выполнять только CPU через специальные функции z__shared__ - переменные не могут инициализироваться при объявлении
Расширения языка C Новые типы данных: z 1/2/3/4 -мерные вектора из базовых типов z(u)char, (u)int, (u)short, (u)long, long zfloat, double zdim 3 – uint 3 с конструктором, позволяющим задавать не все компоненты z. Не заданные инициализируются единицей
Расширения языка С int 2 float 4 float 2 dim 3 a = make_int 2 ( b = make_float 4 x = make_float 2 grid = dim 3 ( blocks = dim 3 ( 1, 7 ); ( a. x, a. y, 1. 0 f, 7 ); ( b. z, b. w ); 10 ); 16, 16 ); Для векторов не определены покомпонентные операции Для double и long возможны только вектора размера 1 и 2.
Расширения языка С z. В любом CUDA kernel’e доступны: ydim 3 yuint 3 yint grid. Dim; block. Idx; block. Dim; thread. Idx; warp. Size;
Расширения языка С z. Как запустить ядро с общим кол-во тредов равным nx? float * data; dim 3 threads(256, 1, 1); dim 3 blocks(nx / 256, 1); Неявно предпологаем, что nx кратно 256 inc. Kernel<<
Расширения языка С Общий вид команды для запуска ядра inc. Kernel<<
Программный стек CUDA
Программный стек CUDA z Программы могут использовать GPU посредством: z Обращения к стандартным функциям библиотек (BLAS, FFTW) z Очень просто z Не очень эффективно z Использования CUDA runtime API z Использованием CUDA driver API
Основы CUDA API Многие функции API асинхронны: z. Запуск ядра z. Копирование при помощи функций *Async z. Копирование device <-> device z. Инициализация памяти
Основы CUDA API char * cuda. Get. Error. String cuda. Error_t cuda. Get. Last. Error cuda. Error_t cuda. Thread. Synchronize cuda. Error_t cuda. Event. Create cuda. Error_t cuda. Event. Record cuda. Error_t ( cuda. Error_t ); (); ( cuda. Event_t * ); cuda. Stream_t ); cuda. Event. Query ( cuda. Event_t ); cuda. Event. Synchronize ( cuda. Event_t ); cude. Event. Elapsed. Time ( float * time, cuda. Event_t start, cuda. Event_t stop ); cuda. Event. Destroy ( cuda. Event_t ); cuda. Get. Device. Count ( int * ); cuda. Get. Device. Propertis ( cuda. Device. Prop * props, int device. No );
CUDA Compute Capability Возможности GPU обозначаются при помощи Compute Capability, например 1. 1 z. Старшая цифра соответствует архитектуре z. Младшая – небольшим архитектурным изменениям z. Можно получить из полей major и minor структуры cuda. Device. Prop
Получение информации о GPU int main ( int argc, char * argv [] ) { int device. Count; cuda. Device. Prop dev. Prop; cuda. Get. Device. Count ( &device. Count ); printf ( "Found %d devicesn", device. Count ); for ( int device = 0; device < device. Count; device++ ) { cuda. Get. Device. Properties ( &dev. Prop, device ); printf ( "Device %dn", device ); printf ( "Compute capability : %d. %dn", dev. Prop. major, dev. Prop. minor ); printf ( "Name : %sn", dev. Prop. name ); printf ( "Total Global Memory : %dn", dev. Prop. total. Global. Mem ); printf ( "Shared memory per block: %dn", dev. Prop. shared. Mem. Per. Block ); printf ( "Registers per block : %dn", dev. Prop. regs. Per. Block ); printf ( "Warp size : %dn", dev. Prop. warp. Size ); printf ( "Max threads per block : %dn", dev. Prop. max. Threads. Per. Block ); printf ( "Total constant memory : %dn", dev. Prop. total. Const. Mem ); } return 0; }
Compute Capability GPU Tesla S 1070 Compute Capability 1. 3 Ge. Force GTX 260 1. 3 Ge. Force 9800 GX 2 1. 1 Ge. Force 9800 GTX 1. 1 Ge. Force 8800 GTX 1. 0 RTM Appendix A. 1 CUDA Programming Guide
Compute Capability z Compute Caps. – доступная версия CUDA y Разные возможности HW y Пример: x В 1. 1 добавлены атомарные операции в global memory x В 1. 2 добавлены атомарные операции в shared memory x В 1. 3 добавлены вычисления в double z Узнать доступный Compute Caps. можно через cuda. Get. Device. Properties() y См. CUDAHello. World z Сегодня Compute Caps: y Влияет на правила работы с глобальной памятью
Компиляция программ z. Используем утилиту nvcc (c: Program FilesNVIDIA GPU Computing ToolkitCUDAv 4. 2bin) z. Используем MS Visual Studio и создаем вручную проект z. Используем NVIDIA Nsight z. New->Project->NVIDIA->CUDA 4. 2 Runtime API
Создание проекта в VS Создать новый проект Добавить к нему. cu файл Задать правила компиляции. cu файлов (CUDA SDK – Cuda. rules) Project Properties -> Linker -> General -> Additional Library = $(CUDA_LIB_PATH) Linker->Input->Additional Dependencies>cudart. lib
Установка CUDA на компьютер CUDA driver CUDA Toolkit CUDA SDK Дополнительно: ◦ NVIDIA Nsight ◦ CUDA VS Wisard (http: //sourceforge. net/projects/cudawizard)
Лабораторная работа № 1 Запуск “Hello World” CUDA Получение информации о GPU Написать программу сложения 2 векторов Написать программу перемножения 2 матриц