Скачать презентацию Программирование графических ускорителей CUDA Лекция 1 Зачем Скачать презентацию Программирование графических ускорителей CUDA Лекция 1 Зачем

Lection-1.ppt

  • Количество слайдов: 62

Программирование графических ускорителей CUDA Лекция 1 Программирование графических ускорителей CUDA Лекция 1

Зачем нужны многоядерные системы Посмотрим на частоты CPU: ◦ ◦ ◦ ◦ 1978 2004 Зачем нужны многоядерные системы Посмотрим на частоты 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 Динамика роста производительности для CPU и GPU

Зачем нужны многоядерные системы z. Был сделан огромный скачок z. За последние 5 лет Зачем нужны многоядерные системы z. Был сделан огромный скачок z. За последние 5 лет z. Минимальный рост частот z. Развитие многоядерных систем z. Видно, что одноядерные системы зашли в тупик

Зачем нужны многоядерные системы Причины тупика: ◦ Энерговыделение ~ четвертой степени частоты ◦ Ограничения Зачем нужны многоядерные системы Причины тупика: ◦ Энерговыделение ~ четвертой степени частоты ◦ Ограничения техпроцесса

Зачем нужны многоядерные системы Таким образом, повышение быстродействия следует ждать именно от параллельности. На Зачем нужны многоядерные системы Таким образом, повышение быстродействия следует ждать именно от параллельности. На самом деле CPU уже давно используют параллельную обработку для повышения производительности ◦ Конвейер ◦ Multithreading ◦ SSE

Intel Core 2 Duo 32 Кб L 1 кэш данных и команд для каждого Intel Core 2 Duo 32 Кб L 1 кэш данных и команд для каждого ядра 2/4 Мб общий L 2 кэш Единый образ памяти для каждого ядра необходимость синхронизации кэшей

Symmetric Multiprocessor Architecture (SMP) Symmetric Multiprocessor Architecture (SMP)

Symmetric Multiprocessor Architecture (SMP) Каждый процессор zимеет свои L 1 и L 2 кэши Symmetric Multiprocessor Architecture (SMP) Каждый процессор zимеет свои L 1 и L 2 кэши zподсоединен к общей шине zотслеживает доступ других процессоров к памяти для обеспечения единого образа памяти (например, один процессор хочет изменить данные, кэшированные другим процессором)

Blue. Gene/L Blue. Gene/L

Blue. Gene/L z 65536 dual-core nodes znode y 770 Mhz Power. PC y. Double 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 Blue. Gene/L

Архитектура G 80 Массив из потоковых мультипроцессоров Архитектура G 80 Массив из потоковых мультипроцессоров

Архитектура G 80 Архитектура G 80

Примеры многоядерный систем z Необходимо обратить ваше внимание на следующие особенности: 1) Как правило Примеры многоядерный систем z Необходимо обратить ваше внимание на следующие особенности: 1) Как правило вычислительный узел – достаточно маломощный процессор 2) Вычислительные узлы имеют свою оперативную память и свой кэш 3) Вычислительные узлы объединяются в более крупные блоки 4) Крупные блоки могут объединяться с целью наращивания вычислительной мощи

Виды параллелизма z. На уровне задач z. На уровне данных Виды параллелизма z. На уровне задач z. На уровне данных

Классификация Классификация

Классификация CPU – SISD ◦ Multithreading: позволяет запускать множество потоков – параллелизм на уровне Классификация CPU – SISD ◦ Multithreading: позволяет запускать множество потоков – параллелизм на уровне задач (MIMD) или данных (SIMD) ◦ SSE: набор 128 битных регистров ЦПУ можно запаковать 4 32 битных скаляра и проводить над ними операции одновременно (SIMD) GPU – SIMD*

SIMD z. На входе поток однородных элементов, каждый из которых может быть обработан независимо SIMD z. На входе поток однородных элементов, каждый из которых может быть обработан независимо z. На выходе – однородный поток z. Обработкой занимается ядро (kernel)

SIMD • Так каждый элемент может быть обработан независимо от других, то их можно SIMD • Так каждый элемент может быть обработан независимо от других, то их можно обрабатывать параллельно • Можно соединять между собой отдельные ядра для получения более сложной схемы обработки

Эволюция GPU z. Voodoo - растеризация треугольников, наложение текстуры и буфер глубины z. Очень Эволюция GPU z. Voodoo - растеризация треугольников, наложение текстуры и буфер глубины z. Очень легко распараллеливается z. На своих задачах легко обходил CPU

Эволюция GPU: Шейдеры !!ARBvp 1. 0 ATTRIB PARAM TEMP pos = vertex. position; eye Эволюция 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. Компилируется Эволюция GPU: Шейдеры z. Работают с 4 D float-векторами z. Специальный ассемблер z. Компилируется драйвером устройства z. Отсутствие переходов и ветвления

Отличия CPU от GPU z. Очень высокая степень параллелизма z. Основная часть чипа занята Отличия CPU от GPU z. Очень высокая степень параллелизма z. Основная часть чипа занята логикой, а не кэшем z. Ограничения по функциональности

GPGPU z. Использование GPU для решения не графических задач z. Вся работа с GPU GPGPU z. Использование GPU для решения не графических задач z. Вся работа с GPU идет через графический API (Open. GL, D 3 D) z. Программы используют сразу два языка – один традиционный (С++) и один шейдерный z. Ограничения, присущие графическим API

CUDA (Compute Unified Device Architecture) Программирование массивнопараллельных систем требует специальных систем/языков. Программирование ресурсов CPU CUDA (Compute Unified Device Architecture) Программирование массивнопараллельных систем требует специальных систем/языков. Программирование ресурсов CPU ограничено ◦ Multithreading ◦ SSE ◦ Часто bottleneck – в пропускной способности памяти CUDA - система (библиотеки и расширенный C) для программирования GPU

CUDA “Hello World” #define N (1024*1024) __global__ void kernel ( float * data ) 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<<>> ( dev ); cuda. Memcpy ( a, dev, N * sizeof ( float ), cuda. Memcpy. Device. To. Host ); cuda. Free ( dev ); for (int idx = 0; idx < N; idx++) return 0; } printf("a[%d] = %. 5 fn", idx, a[idx]);

CUDA “Hello World” __global__ void kernel ( float * data ) { int idx 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; // выделить память CUDA “Hello World” float a [N]; float * dev = NULL; // выделить память на GPU под N элементов cuda. Malloc ( (void**)&dev, N * sizeof ( float ) ); // запустить N нитей блоками по 512 нитей // выполняемая на нити функция - kernel // массив данных - dev kernel<<>> ( dev ); // скопировать результаты из памяти GPU (DRAM) в // память CPU (N элементов) cuda. Memcpy ( a, dev, N * sizeof ( float ), cuda. Memcpy. Device. To. Host ); cuda. Free ( dev // освободить память GPU );

Подход CUDA Исходная задача разбивается на подзадачи, которые можно решать независимо друг от друга. Подход CUDA Исходная задача разбивается на подзадачи, которые можно решать независимо друг от друга. Каждая из этих подзадач решается набором взаимодействующих между собой нитей

SIMT (Single Instruction, Multiple Threads) z. Параллельно на каждом SM выполняется большое число отдельных SIMT (Single Instruction, Multiple Threads) z. Параллельно на каждом SM выполняется большое число отдельных нитей (threads) z. Нити подряд разбиваются на warp’ы (по 32 нити) и SM управляет выполнением warp’ов z. Нити в пределах одного warp’а выполняются физически параллельно z. Большое число warp’ов покрывает латентность

Программная модель CUDA z GPU (device) это вычислительное устройство, которое: y. Является сопроцессором к Программная модель CUDA z GPU (device) это вычислительное устройство, которое: y. Является сопроцессором к CPU (host) y. Имеет собственную память (DRAM) y. Выполняет одновременно очень много нитей

Программная модель CUDA z Код состоит как из последовательных, так и из параллельных частей Программная модель CUDA z Код состоит как из последовательных, так и из параллельных частей z Последовательные части кода выполняются на CPU z Массивно-параллельные части кода выполняются на GPU как ядра

Программная модель CUDA Отличия нитей между CPU и GPU z Нити на GPU очень Программная модель CUDA Отличия нитей между CPU и GPU z Нити на GPU очень «легкие» z HW планировщик задач z Почти нулевые затраты планировщика z Для полноценной загрузки GPU нужны тысячи нитей x Для покрытия латентностей операций чтения / записи x Для покрытия латентностей sfu инструкций

Программная модель CUDA z. Параллельная часть кода выполняется как очень большое количество нитей z. Программная модель CUDA z. Параллельная часть кода выполняется как очень большое количество нитей z. Нити группируются в блоки (blocks) фиксированного размера z. Блоки объединяются в сеть блоков (grid) z. Ядро выполняется на сетке из блоков z. Каждая нить и блок имеют свой идентификатор z. Нить использует идентификаторы для определения с каким элементом работать

Программная модель CUDA Программная модель CUDA

Программная модель CUDA z. Потоки в CUDA объединяются в блоки: y. Возможна 1 D, Программная модель CUDA z. Потоки в CUDA объединяются в блоки: y. Возможна 1 D, 2 D, 3 D топология блока x. Общее кол-во потоков в блоке ограничено x. В текущем HW это 512 потоков

Программная модель CUDA z. Блоки в CUDA объединяются в сеть: y. Возможна 1 D Программная модель CUDA z. Блоки в CUDA объединяются в сеть: y. Возможна 1 D и 2 D топология сети

Программная модель CUDA z. Каждый блок целиком выполняется на одном SM z. Нити одного Программная модель CUDA z. Каждый блок целиком выполняется на одном SM z. Нити одного блока могут взаимодействовать между собой z. Через shared-память z. Через барьерную синхронизацию z. Нити разным блоков не могут взаимодействовать между собой

Расширения языка C z CUDA – это расширение языка C y [+] спецификаторы для Расширения языка C z CUDA – это расширение языка C y [+] спецификаторы для функций и переменных y [+] новые встроенные типы y [+] встроенные переменные (внутри ядра) y [+] директива для запуска ядра из C кода z Как скомпилировать CUDA код y [+] nvcc компилятор y [+]. cu расширение файла

Расширения языка C z. Спецификатор функций Спецификатор Выполняется на Может вызываться из __device__ device Расширения языка 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. Расширения языка C z. Спецификатор __global__ соответствует ядру z. Может возвращать только void z. Спецификаторы __host__ и __device___ могут использоваться одновременно z. Спецификаторы __global__ и __host__ не могут быть использованы одновременно

Расширения языка C z. Спецификатор __global__ соответствует ядру z. Может возвращать только void z. Расширения языка C z. Спецификатор __global__ соответствует ядру z. Может возвращать только void z. Спецификаторы __host__ и __device___ могут использоваться одновременно z. Спецификаторы __global__ и __host__ не могут быть использованы одновременно

Расширения языка C Ограничения на функции, выполняемые на GPU: z. Нельзя брать адрес (за Расширения языка C Ограничения на функции, выполняемые на GPU: z. Нельзя брать адрес (за исключением __global__) z. Не поддерживается рекурсия z. Не поддерживаются static-переменные внутри функции z. Не поддерживается переменное число входных аргументов

Расширения языка C Ограничения на спецификаторы переменных: z. Нельзя применять к полям структуры или Расширения языка C Ограничения на спецификаторы переменных: z. Нельзя применять к полям структуры или union z. Не могут быть extern z. Запись в __constant__ может выполнять только CPU через специальные функции z__shared__ - переменные не могут инициализироваться при объявлении

Расширения языка C Новые типы данных: z 1/2/3/4 -мерные вектора из базовых типов z(u)char, Расширения языка 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 Расширения языка С 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 Расширения языка С z. В любом CUDA kernel’e доступны: ydim 3 yuint 3 yint grid. Dim; block. Idx; block. Dim; thread. Idx; warp. Size;

Расширения языка С z. Как запустить ядро с общим кол-во тредов равным nx? float Расширения языка С z. Как запустить ядро с общим кол-во тредов равным nx? float * data; dim 3 threads(256, 1, 1); dim 3 blocks(nx / 256, 1); Неявно предпологаем, что nx кратно 256 inc. Kernel<<>> ( data ); <<< , >>> угловые скобки, внутри которых задаются параметры запуска ядра: • Кол-во блоке в сетке • Кол-во потоков в блоке

Расширения языка С Общий вид команды для запуска ядра inc. Kernel<<<bl, th, ns, st>>> Расширения языка С Общий вид команды для запуска ядра inc. Kernel<<>> ( data ); zbl – число блоков в сетке zth – число нитей в сетке zns – количество дополнительной sharedпамяти, выделяемое блоку zst – поток, в котором нужно запустить ядро

Программный стек CUDA Программный стек CUDA

Программный стек CUDA z Программы могут использовать GPU посредством: z Обращения к стандартным функциям Программный стек CUDA z Программы могут использовать GPU посредством: z Обращения к стандартным функциям библиотек (BLAS, FFTW) z Очень просто z Не очень эффективно z Использования CUDA runtime API z Использованием CUDA driver API

Основы CUDA API Многие функции API асинхронны: z. Запуск ядра z. Копирование при помощи Основы CUDA API Многие функции API асинхронны: z. Запуск ядра z. Копирование при помощи функций *Async z. Копирование device <-> device z. Инициализация памяти

Основы CUDA API char * cuda. Get. Error. String cuda. Error_t cuda. Get. Last. Основы 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. CUDA Compute Capability Возможности GPU обозначаются при помощи Compute Capability, например 1. 1 z. Старшая цифра соответствует архитектуре z. Младшая – небольшим архитектурным изменениям z. Можно получить из полей major и minor структуры cuda. Device. Prop

Получение информации о GPU int main ( int argc, char * argv [] ) Получение информации о 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 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 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. Используем утилиту 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 файл Задать правила Создание проекта в 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 на компьютер CUDA driver CUDA Toolkit CUDA SDK Дополнительно: ◦ NVIDIA Nsight ◦ CUDA VS Wisard (http: //sourceforge. net/projects/cudawizard)

Лабораторная работа № 1 Запуск “Hello World” CUDA Получение информации о GPU Написать программу Лабораторная работа № 1 Запуск “Hello World” CUDA Получение информации о GPU Написать программу сложения 2 векторов Написать программу перемножения 2 матриц