Часть I: Введение в CUDA Александр Межов Ассистент кафедры Информатики as.mezhov@gmail.com 30 сентября 2011 Кафедра Информатики. - презентация

1 Часть I: Введение в CUDA Александр Межов Ассистент кафедры Информатики 30 сентября 2011 Кафедра Информатики

2 АППАРАТНЫЕ ОСОБЕННОСТИ GPU Краткий обзор архитектурных особенностей GPU

3 Основные тенденции Переход к многопроцессорным системам Развития технологий параллельного программирования – OpenMP, MPI, TPL etc. – Простота в использовании

4 Классификация архитектур Виды параллелизма – На уровне данных (Data) – На уровне задач (Instruction) Single Instruction (SI)Multiple Instruction (MI) Single Data (SD)SISDMISD Multiple Data (MD)SIMDMIMD *GPU: SIMT – Single Instruction Multiple Thread

5 Архитектура многоядерных CPU Кэш первого уровня – для инструкций (L1-I) – для данных (L1-D) Кэш второго уровня – на одном кристалле – используется раздельно Проблема синхронизации кэш-памяти Processor 1 Processor 2 L1-I L2 L1-D L1-I L1-D

6 Архитектура GPU: Device L2... Texture Processing Cluster Texture SM... Texture Processing Cluster Texture SM... Device RAM

7 Архитектура GPU: TPC Кластер текстурных блоков (TPC) – Память для текстур – Потоковый мультипроцессор Texture Processing Cluster (TPC) Texture memory Streaming Multiprocessor... Streaming Multiprocessor

8 Архитектура GPU: SM Память констант Память инструкций Регистровая память Разделяемая память 8 скалярных процессоров 2 суперфункциональных блока Streaming Multiprocessor Instructions Constants Shared Memory SP SFU Registers

9 Основные отличия GPU от CPU Высокая степень параллелизма (SIMT) Минимальные затраты на кэш-память Ограничения функциональности

10 РАЗВИТИЕ ВЫЧИСЛЕНИЙ НА GPU Развитие технологии неграфических вычислений

11 Эволюция GPU Voodoo T&L Shader Floating- point GPGPU

12 General-Purpose Computation on GPU – Вычисления на GPU общего (неграфического) назначения – AMD FireStream – NVIDIA CUDA – DirectCompute (DirectX 10) – OpenCL

13 ПРОГРАММНАЯ МОДЕЛЬ CUDA Основные понятия и определения CUDA

14 CUDA – Compute Unified Device Architecture Host – CPU (Central Processing Unit) Device – GPU (Graphics Processing Unit)

15 Организация работы CUDA GPU Исходные данные Блок Потоки Блок Потоки Блок Потоки Ядро (Kernel)

16 Warp и латентность Warp – Порция потоков для выполнения на потоковом мультипроцессоре (SM) Латентность – Общая задержка всех потоков warpа при выполнении инструкции

17 Топология блоков (block) Возможна 1, 2 и 3-мерная топология Количество потоков в блоке ограничено (512) 1D 2D 3D

18 Топология сетки блоков (grid) Возможна 1 и 2-мерная топология Количество блоков в каждом измерении ограничено 65536=2 16 1D 2D

19 Адресация элементов данных CUDA предоставляет встроенные переменные, которые идентифицируют блоки и потоки – blockIdx – blockDim – threadIdx 1D Grid & 2D Block: int dataIdnex = blockIdx.x * blockDim.x + threadIdx.x

20 Барьерная синхронизация Синхронизация потоков блока осуществляется встроенным оператором __synchronize Блок Потоки

21 CUDA: РАСШИРЕНИЕ C++ Особенности написания программ для GPU CUDA

22 Расширение языка С++ Новые типы данных Спецификаторы для функций Спецификаторы для переменных Встроенные переменные (для ядра) Директива для запуска ядра

23 Процесс компиляции Файлы CUDA (GPU) *.cu Файлы CUDA (GPU) *.cu Файлы CPU *.cpp, *.h Файлы CPU *.cpp, *.h Исполняемый модуль *.dll, *.exe Исполняемый модуль *.dll, *.exe nvccVC90

24 Типы данных CUDA 1, 2, 3 и 4-мерные вектора базовых типов – Целые: (u)char, (u)int, (u)short, (u)long, longlong – Дробные: float, double – Пример: float(1), float2, float3, float4 dim3 ~ uint3 – Пример: dim3(n) = uint(n,1,1)

25 Спецификаторы функций СпецификаторВыполняетсяВызывается __device__device __global__devicehost __host__host

26 Спецификаторы функций Ядро помечается __global__ Ядро не может возвращать значение Возможно совместное использование __host__ и __device__ Спецификаторы __global__ и __host__ не могут использоваться совместно

27 Ограничения функций GPU Не поддерживается рекурсия Не поддерживаются static-переменные Нельзя брать адрес функции __device__ Не поддерживается переменное число аргументов

28 Спецификаторы переменных СпецификаторНаходитсяДоступнаВид доступа __device__device R __constant__devicedevice / hostR / R/W __shared__deviceblockR/W

29 Ограничения переменных GPU Переменные __shared__ не могут инициализироваться при объявлении Запись в __constant__ может производить только host через CUDA API Спецификаторы нельзя применять к полям структур и union

30 Переменные ядра dim3 gridDim unit3 blockIdx dim3 blockDim uint3 threadIdx int warpSize

31 Директива запуска ядра Kernel >>(data) – blocks – число блоков в сетке – threads – число потоков в блоке

32 Общая структура программы CUDA __global__ void Kernel(float* data) {... } void main() {... Kernel >>(data);... }

33 Предустановки Видеокарта NVIDIA с поддержкой CUDA Драйвера устройства с поддержкой CUDA NVIDIA CUDA Toolkit NVIDIA CUDA SDK NVIDIA Nsight Visual Studio 2008 / 2010 Компилятор Visual C++ 9.0

34 Пример программы CUDA

35 Литература NVIDIA Developer Zone – NVIDAI CUDA – Неграфические вычисления на графических процессорах – Создание простого приложения CUDA в Visual Studio 2010 – studio-2010.html studio-2010.html

36 Спасибо за внимание Александр Межов 30 сентября 2011 Кафедра Информатики

37 Часть II: Обзор CUDA API Александр Межов Ассистент кафедры Информатики Кафедра Информатики