Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек...

59
CUDA: Новая архитектура для вычислений на GPU

Transcript of Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек...

Page 1: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CUDA:Новая архитектура для вычислений на GPU

Page 2: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

История GPGPU

Графические процессоры (GPUs) использовались для неграфических вычисленийиспользовались для неграфических вычислений в течение нескольких летGeneral-Purpose computation on GPUs: GPGPUGeneral-Purpose computation on GPUs: GPGPUПриложения GPGPU:

Симуляция физикиСимуляция физикиОбработка сигналовВычислительная математика/геометрияОперации с базами данныхВычислительная биологияВВычислительная экономикаКомпьютерное зрение

© NVIDIA Corporation 2006

Page 3: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Для чего использовать GPU?

GPU является программируемым процессором:С поддержкой языков высокого уровняС поддержкой 32-bit floating point IEEE-754Большой вычислительной мощностью:Большой вычислительной мощностью:

LOP

SG

FL G80GL = Quadro FX 5600

G80 = GeForce 8800 GTX

G71 = GeForce 7900 GTX

G70 = GeForce 7800 GTX

NV40 = GeForce 6800 Ultra

NV35 = GeForce FX 5950 Ultra

NV30 = GeForce FX 5800

© NVIDIA Corporation 2006

Page 4: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Что стоит за эволюцией GPU?GPU предназначен для вычислений с большим параллелизмом и интенсивной арифметикойпараллелизмом и интенсивной арифметикой

Гораздо большее число транзисторов отведено на обработку данных, а не на управление исполнением (f )(flow control)То, чем является графика

ALUControl

ALU

Cache

ALUALU

DRAMDRAM

CPU GPU

© NVIDIA Corporation 2006

CPU GPU

Page 5: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Каково применение GPU?

GPU демонстрируют хорошие результаты вПараллельной обработке данныхПараллельной обработке данных

С одной и той же последовательностью действий,применяемых к большому объёму данных

С высокой плотностью арифметикиС высокой плотностью арифметикиДостаточно большим отношением числа арифметических инструкций к числу обращений к памяти

Одни и те же вычисления означают меньшие требования к управлению исполнением (flow control)

Высокая плотность арифметики и большой объём данныхВысокая плотность арифметики и большой объём данных означают возможность покрытия латентности памяти вычислениями (вместо больших кэшей на CPU)

© NVIDIA Corporation 2006

Page 6: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Недостатки традиционной модели GPGPU

До недавнего времени GPU могли программироваться только посредством графических APIграфических API

Длительное время требуемое дляДлительное время, требуемое для изучения

Избыточность и накладные расходы графических APIграфических API

© NVIDIA Corporation 2006

Page 7: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Недостатки традиционной модели GPGPU: аппаратные ограничения

Нет линейной адресации памяти DRAMНет линейной адресации памяти DRAMТолько операция ‘gather’

Control Control

DRAM

ALUControl

CacheALU ALU ...

d d d d

ALUControl

CacheALU ALU ...

d d d d

Нет возможности записи в произвольные адреса

DRAM d0 d1 d2 d3 d4 d5 d6 d7 …

ALUControl

CacheALU ALU ... ALU

Control

CacheALU ALU ... …

Ограниченная гибкость

DRAM d0 d1 d2 d3 d4 d5 d6 d7 …

© NVIDIA Corporation 2006

Ограниченная гибкость

Page 8: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Недостатки традиционной модели GPGPU: аппаратные ограничения

Приложения часто ограничены пропускной способностью памяти

ALUControl

ALU ALU ... ALUControl

ALU ALU ...

DRAM

Cache

d0 d1 d2 d3

Cache

d4 d5 d6 d7 …d0 d1 d2 d3 d4 d5 d6 d7

Простой вычислительных мощностей

© NVIDIA Corporation 2006

Page 9: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CUDA

Compute Unified Device Architecture: новая программно аппаратная

CPUApplication

программно-аппаратная архитектура для вычислений на GPU

CUDA Libraries(FFT, BLAS)

Присутствует в GeForce 8800, Quadro FX 5600/4600 и вышеН ф

CUDA Runtime

Независима от графических APIРяд особенностей

CUDA Driver

дпредназначенных для вычислений общего назначения

GPU

© NVIDIA Corporation 2006

назначения

Page 10: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Нововведения CUDA:Простота и легковесность

Среда разработки является расширением языка программирования C

ММеньшее время на изучение

О й йОтдельный программно-аппаратный стек для вычислений

Высокая производительностьВысокая производительность

© NVIDIA Corporation 2006

Page 11: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Нововведения CUDA:байт-адресация и scatter

Имеется возможность байт-адресацииИмеется возможность байт-адресации Gather:

Control Control

DRAM

ALUControl

CacheALU ALU ...

d d d d

ALUControl

CacheALU ALU ...

d d d d

+ Scatter:

DRAM d0 d1 d2 d3 d4 d5 d6 d7 …

ALUControl

CacheALU ALU ... ALU

Control

CacheALU ALU ... …

Большая гибкость

DRAM d0 d1 d2 d3 d4 d5 d6 d7 …

© NVIDIA Corporation 2006

Большая гибкость

Page 12: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Нововведения CUDA:Триггерная разделяемая память

CUDA йCUDA предоставляет доступ к регистровой разделяемой памяти для обмена данными между потокамиду

ALUControl

CacheALU ALU ... ALU

Control

CacheALU ALU ...

Sharedmemory

Cache

d0 d1 d2 d3

Sharedmemory

Cache

d4 d5 d6 d7

DRAM d0 d1 d2 d3 d4 d5 d6 d7 …

Big memory bandwidth saving

© NVIDIA Corporation 2006

Page 13: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Разделяемая памятьSharedMemory

ThreadExecutionManager

Уменьшает расстояние между АЛУ иControl

ALUShared

Data

р ду

Данными, позволяя

•Сократить число обращений к

Control

DRAM

P1P2P

DataPn’=P1+P2+P3+P4

•Сократить число обращений к внешней памяти

•Минимизировать избыточностьALU

P3P4P5Pn’=P1+P2+P3+P4

•Минимизировать избыточность загрузки данных и вычислений

•Увеличить интенсивность

Control

Увеличить интенсивность арифметики держа данные ближе к АЛУ keeping

ALUALU

Pn’=P1+P2+P3+P4

p g

© NVIDIA Corporation 2006

Parallel execution through Shared Memory

Page 14: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Обзор

Модель программирования CUDA

Аппаратная реализация

CUDA API

Технические спецификации Tesla / GeForce 8 Series / Quadro 5600/4600

CUDA-библиотеки прикладного уровня: CUBLAS, CUFFT

© NVIDIA Corporation 2006

CUFFT

Page 15: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Модель программирования:Массивно-параллельный процессор

GPU - вычислительное устройство которое:GPU - вычислительное устройство которое:Является сопроцессором (device) для центрального процессора (CPU / host).Об б й DRAM ( йОбладает собственной памятью DRAM (память устройства, device memory)Параллельно обрабатывает множество потоков исполнения

Ядро (kernel) – функция (последовательность команд GPU)с параллелизмом данных исполняемая над большимс параллелизмом данных, исполняемая над большим количеством потоков.

Отличия потоков GPU от CPUДля полной эффективности GPU нужны тысячи потоковМногоядерным CPU нужно гораздо меньше

© NVIDIA Corporation 2006

Многоядерным CPU нужно гораздо меньше

Page 16: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Модель программирования (SPMD + SIMD): группирование потоков

Ядро (kernel) : Host Device

исполняется над сеткой (grid) блоков потоков (thread blocks)

Kernel 1

Grid 1

Block(0, 0)

Block(1, 0)

Block(2, 0)

Блок потоков (thread block):Набор потоков способных общаться между собой

Block(0, 1)

Block(1, 1)

Block(2, 1)

общаться между собой посредством:

разделяемой памяти (shared memory)

Kernel 2

Grid 2

(shared memory)точек синхронизации (барьеров) Block (1, 1)

Thread Thread Thread Thread ThreadОдновременно исполняется только одна сетка (grid)

Нет координаты сеткиThread(0, 1)

Thread(1, 1)

Thread(2, 1)

Thread(3, 1)

Thread(4, 1)

Thread(0, 0)

Thread(1, 0)

Thread(2, 0)

Thread(3, 0)

Thread(4, 0)

© NVIDIA Corporation 2006

Thread(0, 2)

Thread(1, 2)

Thread(2, 2)

Thread(3, 2)

Thread(4, 2)

Page 17: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Модель программирования:Координаты потоков

К бКаждому потоку и блоку потоков присваиваются координаты Device

Grid 1

Координата блока: 2D (1D) Block(0, 0)

Block(1, 0)

Block(2, 0)

Block Block BlockКоордината потока: 3D (2D, 1D)

В зависимости от задачи может

Block(0, 1)

Block(1, 1)

Block(2, 1)

В зависимости от задачи может упростить адресацию данных

Обработка изображений

Block (1, 1)

Thread(0, 0)

Thread(1, 0)

Thread(2, 0)

Thread(3, 0)

Thread(4, 0)

Решение уравнений в ч.п. (PDE) Thread(0, 1)

Thread(1, 1)

Thread(2, 1)

Thread(3, 1)

Thread(4, 1)

Thread(0 2)

Thread(1 2)

Thread(2 2)

Thread(3 2)

Thread(4 2)

© NVIDIA Corporation 2006

(0, 2) (1, 2) (2, 2) (3, 2) (4, 2)

Page 18: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Модель программирования:пространства памятиКаждый поток имеет доступ к :

Локальным регистрам GridЛокальным регистрам (R/W, on-chip, per-thread)Локальной памяти (R/W, DRAM, uncached, per-th d)

Grid

Block (0, 0) Block (1, 0)

thread)Разделяемой памяти (R/W, on-chip, per-block)Глобальной памяти

Shared Memory

Registers Registers

Shared Memory

Registers Registers

Глобальной памяти (R/W, DRAM, uncahed, per-context)Константной памяти (R only DRAM cached per

Thread (0, 0) Thread (1, 0) Thread (0, 0) Thread (1, 0)

(R-only, DRAM, cached, per-context)Текстурной памяти (R-only, cached, per-context)

Global

LocalMemory

LocalMemory

LocalMemory

LocalMemory

Host

Центральный процессор может обновлять/запрашивать

Глобальную память (DRAM)ConstantMemory

GlobalMemory

Host

© NVIDIA Corporation 2006

у ( )Константную память (DRAM)Текстурную память (DRAM)

TextureMemory

Page 19: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Модель программирования:подход к разработке

Л б DRAMЛокальная и глобальная память находится в DRAM –латентность гораздо выше, чем у регистровой памяти

Некешируемару

Задача разбивается на подзадачи В б бВходные данные разбиваются на подблоки вмещающиеся в разделяемую памятьКаждый подблок обрабатывается блоком потоков:

Подблок загружается в разделяему память из глобальнойПроводятся вычисления над данными в разделяемой памятиКаждый поток может делать много “проходов” над любыми элементами данныхэлементами данныхРезультаты копируются обратно из разделяемой памяти в глобальнуюТипичный но не строго обязательный шаблон

© NVIDIA Corporation 2006

Типичный, но не строго обязательный шаблон

Page 20: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Модель программирования:пример умножения матриц

М P M * N { NМатрица P = M * N {WIDTH x WIDTH}Без выделения блоков:

Один поток рассчитывает один

N

Один поток рассчитывает один элемент результата PM и N эффективно загружаютсяWIDTH раз из глобальной памяти

WID

TH

WIDTH раз из глобальной памяти

M P

HW

IDT

H

© NVIDIA Corporation 2006

WIDTH WIDTH

Page 21: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Модель программированияпример умножения матриц

М P M * N { N ZEМатрица P = M * N {WIDTH x WIDTH}С выделением блоков:

Один блок потоков вычисляет одну

N

IZE

BL

OC

K_S

IZ

Один блок потоков вычисляет одну Psub - подматрицу P размера BLOCK_SIZE x BLOCK_SIZEM и N эффективно загружаются IZ

EB

LO

CK

_SI

WID

TH

M и N эффективно загружаются только WIDTH / BLOCK_SIZE раз из глобальной памяти

M P

BL

OC

K_S

I

Улучшение в BLOCK SIZE

Psub SIZE

HBLOCK_SIZE раз

BLOCK_SIZEBLOCK SIZE BLOCK SIZE BLOCK SIZE

BL

OC

K_S

WID

TH

© NVIDIA Corporation 2006

__ _ _

WIDTH WIDTH

Page 22: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Аппаратная реализация:Набор SIMD мультипроцессоров

GPU представляет из себяDevice

GPU представляет из себя набор мультипроцессоров

Каждый мультипроцессор

Multiprocessor N

Multiprocessor 2Каждый мультипроцессор является SlMD-набором скалярных 32-х битных процессоров

Multiprocessor 2

Multiprocessor 1

р ц р

На каждом такте мультипроцессор исполняет

InstructionUnit

Processor 1 …Processor 2 Processor Mу р ц родну и ту же инструкцию над группой потоков,называемой warp

Число потоков в warp - warp size

© NVIDIA Corporation 2006

Page 23: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Аппаратная реализация:Архитектура памяти

Локальное глобальное DeviceЛокальное, глобальное, константное пространства имён находятся в DRAM-памяти устройства (GPU)

Multiprocessor N

Multiprocessor 2Каждый мультипроцессор обладает:

Файлом локальных 32-х битных регистров

Multiprocessor 2

Multiprocessor 1

Shared Memoryбитных регистров

Разделяемой памятью

y

InstructionUnit

Processor 1

Registers

…Processor 2

Registers

Processor M

Registers

Кэшем констант

Текстурным кэшемConstant

Cache

Device memory

TextureCache

© NVIDIA Corporation 2006

Device memory

Page 24: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Аппаратная реализация :Модель исполнения

Каждый блок потоков состоит из warp’овW SIMD фWarp – SIMD-группа потоков фиксированного размера, состоящая из скалярных потоков с последовательными координатами.

Блок потоков всегда исполняется только на одном мультипроцессоре

Разделяемые переменные блока потоков хранятся в on-chip модуле разделяемой памяти мультипроцессораФайл локальных регистров делится между всеми потоками,Файл локальных регистров делится между всеми потоками,обрабатываемыми мультипроцессором

Слишком большого размер блока потоков для ядра, использующего слишком много локальных регистров, вызывает сбой этапа исполнения

Мультипроцессор может обрабатывать несколько блоков потоков одновременно

Локальный регистровый файл и разделяемая память делятся между всеми потоками / блоками потоков, работающимимежду всеми потоками / блоками потоков, работающими одновременноТем самым, c уменьшением числа используемых локальных регистров (на поток) и размера используемой разделяемой памяти (на блок) увеличивается число потоков / блоков потоков, способных

© NVIDIA Corporation 2006

( ) у ,одновременно находиться в обработкеНе влияет на логику работы, только на производительность

На логическом уровне блоки потоков всегда изолированы

Page 25: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Быстрый обзор

Устройство = GPU = набор мультипроцессоровМультипроцессор = набор скалярных процессоров с файлом разделяемой памятиЯдро = GPU программаЯдро = GPU-программаСетка = массив блоков потоков над которыми исполняется ядроБлок потоков = группа SIMD-потоков исполняющих ядро и имеющих возможность быстрого обмена данными

Memory Location Cached Access WhoLocal Off-chip No Read/write One threadShared On-chip N/A Read/write All threads in a blockGlobal Off-chip No Read/write All threads + hostConstant Off-chip Yes Read All threads + host

© NVIDIA Corporation 2006

pTexture Off-chip Yes Read All threads + host

Page 26: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CUDA: C для GPUПростота и универсальность

Дополнения только там, где нужно

global void KernelFunc( );__global__ void KernelFunc(...);

__shared__ int SharedVar;

KernelFunc<<<500, 128>>>(...);

Явное выделение памяти GPUcudaMalloc(), cudaFree()

Копирование CPU->GPU, GPU->CPU, GPU->GPUcudaMemcpy(), cudaMemcpy2D(), ...

© NVIDIA Corporation 2006

Page 27: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Пример: ядро сложения векторов

////Попарное сложение элементов вектора//Один поток на сложение

__global__ void vectorAdd(void vectorAdd(

float *oC, float *iA, ,float *iB

){int idx = blockDim.x * blockId.x + threadIdx.x;oC[idx] = iA[idx] + iB[idx];

}

© NVIDIA Corporation 2006

}

Page 28: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Пример: код CPU для сложения векторов//Размер вектора в элементахconst int N = 1048576;// б й//размер вектора в байтахconst int dataSize = N * sizeof(float);

//Выделение памяти CPUfloat *h_A = (float *)malloc(dataSize);float *h_B = (float *)malloc(dataSize);_float *h_C = (float *)malloc(dataSize);

//Выделение памяти GPUfloat *d_A, *d_B, *d_C;cudaMalloc((void **)&d_A, dataSize));cudaMalloc((void **)&d B dataSize));cudaMalloc((void **)&d_B, dataSize));cudaMalloc((void **)&d_C, dataSize));

//Инициализировать h_A[], h_B[]…

//Скопировать входные данные в GPU для обработкиcudaMemcpy(d_A, h_A, dataSize, cudaMemcpyHostToDevice) );cudaMemcpy(d_B, h_B, dataSize, cudaMemcpyHostToDevice) );

//Запустить ядро из N / 256 блоков по 256 потоков//Предполагая, что N кратно 256vectorAdd<<<N / 256, 256>>>(d C, d A, d B);/ , ( _ , _ , _ );

//Считать результаты GPUcudaMemcpy(h_C, d_C, dataSize, cudaMemcpyDeviceToHost) );

© NVIDIA Corporation 2006

Page 29: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Пример: параллельный reduce

0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15Индексы

10 1 8 -1 0 -2 3 5 -2 -3 2 7 0 11 0 2Значения

ThreadШ 10 1 2 3 4 5 6 7

8 -2 10 6 0 9 3 7 -2 -3 2 7 0 11 0 2Значения

Thread IDs

Шаг 1 stride = 8

8 -2 10 6 0 9 3 7 -2 -3 2 7 0 11 0 2Значения

0 1 2 3Шаг 2

stride = 4Thread

IDs

8 7 13 13 0 9 3 7 -2 -3 2 7 0 11 0 2Значения

0 1Шаг 3 Thread 0 1

21 20 13 13 0 9 3 7 -2 -3 2 7 0 11 0 2Значенияstride = 2

Шаг 4

IDs

Thread

© NVIDIA Corporation 2006

0

41 20 13 13 0 9 3 7 -2 -3 2 7 0 11 0 2Значения

Шаг 4 stride = 1

Thread IDs

Page 30: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Пример: ядро reduce

global reduce(int *g odata int *g idata){__global__ reduce(int *g_odata, int *g_idata){//Размер данных равен размеру блокаextern __shared__ int sdata[];

//Загрузить данные в разделяемую память//Загрузить данные в разделяемую памятьint tid = threadIdx.x;int i = blockDim.x * blockIdx.x + threadIdx.x;sdata[tid] = g_idata[i];

//Reduce в разделяемой памятиfor(int stride = blockDim.x / 2; stride > 0; stride >>= 1){

__syncthreads();if(tid < stride)if(tid < stride)

sdata[tid] = sdata[tid] + sdata[tid + stride];}

//Записать результат блока в глобальную память//Записать результат блока в глобальную память__syncthreads();if(tid == 0) g_odata[blockIdx.x] = sdata[0];

}

© NVIDIA Corporation 2006

Page 31: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Среда программирования

Расширение среды программирования языка C

Состоит изРасширений синтаксиса

Для написания кода GPUДля написания кода GPURun-time библиотеки:

Общая часть, предоставляющая встроенные векторные типа и подмножеств вызовов RTL поддерживаемых как на CPU, так и на GPUCPU-компонента, для доступа и управления одним или у у рнесколькими GPUGPU-компонента, предоставляющая функции специфические только для GPU

© NVIDIA Corporation 2006

ц ф д

Page 32: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Расширения языка:Виды функций

Executed Only callableExecuted on the:

Only callable from the:

__device__ float DeviceFunc() GPU GPU__ __

__global__ void KernelFunc() GPU CPU__host__ float HostFunc() CPU CPU

__global__ является функцией-ядромВсегда void

__ __

д__device__ и __host__ может быть использовано вместеНельзя взять адрес __device__ функции

В функциях, исполняемых на GPU:Нет рекурсииНет статических переменных внутри функций

© NVIDIA Corporation 2006

Нет статических переменных внутри функцийНет переменного числа аргументов

Page 33: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Расширения языка:Виды переменных

Memory Scope LifetimeMemory Scope Lifetime__device__ __shared__ int SharedVar; shared block block

__device__ int GlobalVar; global grid context

К d i б

__device__ __constant__ int ConstantVar; constant grid context

Ключевое слово __device__ необязательно совместно с__shared__ или __constant__

Локальные переменные без идентификатора вида хранятся в регистрах

За исключением больших структур или массивов, ру уррасполагающихся в локальной памяти

Вид памяти указателя адаптируется к типу присвоенного

© NVIDIA Corporation 2006

д у д ру у рвыражения

Page 34: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Расширения языка:запуск ядра

При запуске ядру передаются обязательные р у ру рпараметры конфигурации сетки

__global__ void KernelFunc(...);__ __dim3 DimGrid(100, 50); //5000 блоков потоковdim3 DimBlock(4, 8, 8); //256 потоков на блокi t Sh dM B t 64 //64 б йsize_t SharedMemBytes = 64; //64 байта разд. памятиKernelFunc<<<DimGrid, DimBlock, SharedMemBytes>>>(...);

О Sh dM B t байт:Опциональные SharedMemBytes байт:Выделяются в дополнение к статически объявленным разделяемым переменнымОтображаются на любую переменную вида:

extern __shared__ float DynamicSharedMem[];

© NVIDIA Corporation 2006

Вызов ядра асинхроненУправление немедленно возвращается к CPU

Page 35: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Расширения языка:Встроенные переменные

dim3 gridDim;Р б ( )Размер сетки в блоках (gridDim.z неиспользовано)

dim3 blockDim;Раз ер о о о б о а о о аРазмеры одного блока в потоках

dim3 blockIdx;Индекс блока внутри сеткеИндекс блока внутри сетке

dim3 threadIdx;Индекс потока внутри блокаИндекс потока внутри блока

Доступно только для кода GPU

© NVIDIA Corporation 2006

Page 36: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Общая библиотека

Встроенные векторные типыПодмножество библиотеки C доступное и для CPU иПодмножество библиотеки C, доступное и для CPU, идля GPU

© NVIDIA Corporation 2006

Page 37: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Общая библиотека:Встроенные векторные типы

[u]char[1..4], [u]short[1..4], [u]int[1..4], [ ]l [1 4] fl t[1 4][u]long[1..4], float[1..4]

Доступ по x, y, z, w:uint4 param;uint4 param;int y = param.y;

dim3Синоним uint3Используется для задания параметров сетки исполнения

© NVIDIA Corporation 2006

Page 38: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Общая библиотека:математические функции

powf sqrtf cbrtf hypotfpowf, sqrtf, cbrtf, hypotfexpf, exp2f, expm1flogf, log2f, log10f, log1pfsinf, cosf, tanfasinf, acosf, atanf, atan2fsinhf, coshf, tanhfasinhf, acoshf, atanhfceil floor trunc roundceil, floor, trunc, roundИ т.д.

Если есть возможность, для CPU-функций используются существующие реализации используемого C-компилятора

© NVIDIA Corporation 2006Поддерживаются только скалярные типы

Page 39: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Общая библиотека:текстурный типТекстурный идентификатор (texture reference)Текстурный идентификатор (texture reference)

texture<float, 2> myTexRef; //2D-текстура типа float

Настройка режимов текстуры на CPUmyTexRef.addressMode[0] = cudaAddressModeWrap;myTexRef.addressMode[1] = cudaAddressModeWrap;myTexRef.filterMode = cudaFilterModeLinear;

Обращение к текстуре из кода GPUfloat4 al e te 2D(m Te Ref )float4 value = tex2D(myTexRef, u, v);

© NVIDIA Corporation 2006

Page 40: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CPU-компонента библиотеки

Функции для:Управления GPU (одним или несколькими)Управления памятьюУправления памятьюУправления текстурамиВзаимодействия с OpenGL и Direct3D9д pОбработки ошибок

Один поток CPU управляет только одним GPUКаждый поток CPU управляет своим GPU

© NVIDIA Corporation 2006

Page 41: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CPU-компонента библиотеки: управление GPU

Запрос параметров GPU в системеcudaGetDeviceCount(), cudaGetDeviceProperties()

В б GPUВыбор GPUcudaChooseDevice(), cudaSetDevice()

© NVIDIA Corporation 2006

Page 42: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CPU-компонента библиотеки: управление памятью

Два вида памяти:Д дЛинейная память: доступ по 32-х битным указателямCUDA-массивы: непрозрачные контейнеры данных

ф бдля текстур; доступ только через функции выборки текстуры

Выделение памятиВыделение памятиcudaMalloc(), cudaMallocPitch(), cudaFree(), cudaMallocArray(), cudaFreeArray()

Копирование CPU->GPU, GPU->GPU, GPU->CPUcudaMemcpy(), cudaMemcpy2D(), cudaMemcpyToArray() cudaMemcpyFromArray() etccudaMemcpyToArray(), cudaMemcpyFromArray(), etc. cudaMemcpyToSymbol(), cudaMemcpyFromSymbol()

Получение адреса device переменных

© NVIDIA Corporation 2006

Получение адреса __device__ переменныхcudaGetSymbolAddress()

Page 43: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CPU-компонента библиотеки: управление текстурами

Текстурный идентификатор (texture reference) может быть привязан к:

CUDA массивамCUDA-массивамЛинейной памяти

Только 1D-текстуры, без фильтрации, по целым ур ф ртекстурным координатам

cudaBindTexture(), cudaUnbindTexture()

© NVIDIA Corporation 2006

Page 44: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CPU-компонента библиотеки: взаимодействие с графическими API

OpenGL buffer objects и Direct3D9 vertex buffersмогут быть отображены в адресное у р дрпространство CUDA-контекста:

Для передачи (отображения) данных в OpenGLДля считывания результатов рендеринга OpenGLcudaGLMapBufferObject(), cudaGLUnmapBufferObject()cudaGLUnmapBufferObject() cudaD3D9MapVertexBuffer(), cudaD3D9UnmapVertexBuffer()

© NVIDIA Corporation 2006

Page 45: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

GPU-компонента библиотеки:математические функции

Некоторые математические функции (e.g. sin(x)) обладают менее точной, но более быстрой реализацией (e g sin(x))быстрой реализацией (e.g. __sin(x))

__powlog, log2, log10__ g, __ g , __ g

__exp__sin, __cos, __tan

© NVIDIA Corporation 2006

Page 46: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

GPU-компонента библиотеки:текстурные выборки

Для идентификаторов текстур, привязанных к 2D CUDA-массивам:

float u v;float u, v;float4 value = tex2D(myTexRef, u, v);

Для текстур, привязанных к линейной памяти:int i;int i;float4 value = tex1Dfetch(myTexRef, i);

© NVIDIA Corporation 2006

Page 47: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

GPU-компонента библиотеки:функции синхронизации

void __syncthreads();

Синхронизует все потоки внутри блока потоковКак только все потоки достигли точки, исполнение продолжаетсяОбычно используется для предотвращения ошибок RAW / WAR / WAW при доступе кошибок RAW / WAR / WAW при доступе к разделяемой или глобальной памятиДопускается внутри ветвлений только если всеДопускается внутри ветвлений только если все потоки блока гарантированно вычисляют одно и то же значение условия

© NVIDIA Corporation 2006

у

Page 48: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Компиляция

Файлы исходного кода на CUDA C компилируются при помощи nvccпомощи nvcc

NVCC является оболочкой над другими инструментамиВызывает соответствующие компиляторы и инструменты,cudacc, g++, cl, ...

NVCC генерирует:C-код (код CPU)

Компилируется вместе с остальными частями приложения, ру р ,написанными на ‘чистом’ C

Объектный код PTX (код GPU)

Любой исполняемый файл с кодом на CUDA требует:CUDA runtime library (cudart)CUDA core library (cuda)

© NVIDIA Corporation 2006

CUDA core library (cuda)

Page 49: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Стадии компиляции CUDA-приложения

C/C++ CUDAApplicationpp

CPU C dNVCC CPU Code

PTX Code

PTX to TargetCompiler

G80 … GPU

© NVIDIA Corporation 2006Target code

Page 50: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Стадии компиляции CUDA-приложения

C/C++ CUDAApplicationpp

NVCC

PTX Code Virtual

PTX to TargetCompiler

Physical

G80 … GPU

© NVIDIA Corporation 2006Target code

Page 51: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

NVCC и виртуальная машина PTX

C/C++ CUDAfloat4 me = gx[gtid];me.x += me.y * me.z;

EDG

C/C++ CUDAApplication

Разделяет код GPU и CPU code Open64

Генерирует PTX-ассемблер GPU EDG CPU Code

Parallel Thread eXecution (PTX)Виртуальная машина и набор инструкцийOpen64Модель программированияРесурсы исполнения и состояниеPTX Code

ld global v4 f32 {$f1 $f3 $f5 $f7} [$r9 0];

© NVIDIA Corporation 2006

ld.global.v4.f32 {$f1,$f3,$f5,$f7}, [$r9+0];mad.f32 $f1, $f5, $f3, $f1;

Page 52: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Отладка в режиме эмуляции

Исполняемый файл скомпилированный вИсполняемый файл, скомпилированный в режим эмуляции (nvcc -deviceemu) работает целиком на CPU

Не требуется драйвер CUDA и GPUКаждый поток GPU эмулируется потоком CPU

При работе в режиме эмуляции можно:Использовать средства отладки CPU (точки останова, просмотр, и т.д.)Обращаться к любым данным “GPU” с CPU и наоборотДелать любые CPU-вызовы из код GPU и наоборотДелать любые CPU вызовы из код GPU и наоборот (например printf())Выявлять ситуации зависания, возникающие из-за неправильного применения syncthreads()

© NVIDIA Corporation 2006

неправильного применения __syncthreads()

Page 53: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Недостатки отладки в режиме эмуляции

Часто работает очень медленно

Разыменование указателей GPU на стороне CPU, илиуказателей CPU на стороне GPU может давать ожидаемыеуказателей CPU на стороне GPU может давать ожидаемые результаты в режиме эмуляции, но генерирует ошибку на настоящем GPU

Результаты операций с плавающей точкой CPU и “настоящего” GPU почти всегда различаются из-за:настоящего GPU почти всегда различаются из за:

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

Опции компилятора для включения строго одинарной точности

© NVIDIA Corporation 2006

Page 54: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Технические спецификацииG F 8800 Q d FX 5600/4600GeForce 8800 и Quadro FX 5600/4600

Мультипроцессоров

Тактовая частота АЛУ

(GHz)

Объём памяти (MB)

GeForce 8800 GTX 16 1.35 768

GeForce 8800 GTS 12 1.2 640

Quadro FX 5600 16 1.35 1500

Quadro FX 4600 12 1.2 768

Максимальный размер блока потоков : 512Максимальный размер сетки блоков в одном изм. : 65535Размер warp’а : 32Размер warp’а : 32Локального регистровый файл мультипроцессора: 32KBРазделяемой память на мультипроцессор : 16 KB

© NVIDIA Corporation 2006

д у р ц рОбщий объём константной памяти : 64 KB

Page 55: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Заключение

Массивно-параллельная архитектура GPU, р р ур ,обладающая высокой и постоянно возрастающей производительностью, хорошо подходит для решения задач с параллелизмом данных.

CUDA на GeForce 8800 и Quadro 5600/4600 улучшает существующую модельулучшает существующую модель программирования GPU (GPGPU), упрощая и расширяя её введением быстрой разделяемой р р д р р дпамяти и возможности синхронизации потоков.

© NVIDIA Corporation 2006

Page 56: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

Extra Slides

CUDA Libraries

Page 57: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CUDA Libraries

CUBLASCUDA “Basic Linear Algebra Subprograms”Implementation of BLAS standard on CUDAImplementation of BLAS standard on CUDAFor details see cublas_library.pdf and cublas.h

CUFFTCUDA Fast Fourier Transform (FFT)FFT one of the most important and widely used numerical algorithmsFor details see cufft library pdf and cufft hFor details see cufft_library.pdf and cufft.h

© NVIDIA Corporation 2006

Page 58: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CUBLAS Library

Self-contained at API levelApplication needs no direct interaction with CUDA driver

Currently only a subset of CUBLAS core functions are implemented

Simple to use:Create matrix and vector objects in GPU memoryFill th ith d tFill them with dataCall sequence of CUBLAS functionsUpload results back from GPU to hostUpload results back from GPU to host

Column-major storage and 1-based indexingFor maximum compatibility with existing Fortran apps

© NVIDIA Corporation 2006

p y g pp

Page 59: Новая архитектура для ... - NVIDIA · Grid 2 shared memory точек синхронизации (барьеров) Block (1, 1) Одновременно исполняется

CUFFT Library

Efficient implementation of FFT on CUDA

F tFeatures1D, 2D, and 3D FFTs of complex-valued signal dataBatch execution for multiple 1D transforms in parallelBatch execution for multiple 1D transforms in parallelTransform sizes (in any dimension) in the range [2, 16384]

© NVIDIA Corporation 2006