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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

Типы данных 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)

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

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

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

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

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

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

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

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

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

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

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

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

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