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

История GPGPU

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

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

© NVIDIA Corporation 2006

Для чего использовать 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

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

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

ALUControl

ALU

Cache

ALUALU

DRAMDRAM

CPU GPU

© NVIDIA Corporation 2006

CPU GPU

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

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

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

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

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

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

© NVIDIA Corporation 2006

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

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

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

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

© NVIDIA Corporation 2006

Недостатки традиционной модели 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

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

Недостатки традиционной модели 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

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

назначения

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

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

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

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

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

© NVIDIA Corporation 2006

Нововведения 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

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

Нововведения 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

Разделяемая память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

Обзор

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

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

CUDA API

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

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

© NVIDIA Corporation 2006

CUFFT

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

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

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

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

© NVIDIA Corporation 2006

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

Модель программирования (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)

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

К бКаждому потоку и блоку потоков присваиваются координаты 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)

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

Локальным регистрам 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

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

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

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

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

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

© NVIDIA Corporation 2006

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

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

М 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

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

М 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

Аппаратная реализация:Набор 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

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

Локальное глобальное 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

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

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

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

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

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

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

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

© NVIDIA Corporation 2006

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

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

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

Устройство = 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

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

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

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

__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

}

Пример: код 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

Пример: параллельный 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

Пример: ядро 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

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

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

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

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

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

© NVIDIA Corporation 2006

ц ф д

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

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

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

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

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

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

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

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

__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

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

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

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

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

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

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

© NVIDIA Corporation 2006

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

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

© NVIDIA Corporation 2006

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

[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

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

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Поддерживаются только скалярные типы

Общая библиотека:текстурный типТекстурный идентификатор (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

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

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

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

© NVIDIA Corporation 2006

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

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

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

© NVIDIA Corporation 2006

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()

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

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

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

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

cudaBindTexture(), cudaUnbindTexture()

© NVIDIA Corporation 2006

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

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

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

© NVIDIA Corporation 2006

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

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

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

void __syncthreads();

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

© NVIDIA Corporation 2006

у

Компиляция

Файлы исходного кода на 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)

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

C/C++ CUDAApplicationpp

CPU C dNVCC CPU Code

PTX Code

PTX to TargetCompiler

G80 … GPU

© NVIDIA Corporation 2006Target code

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

C/C++ CUDAApplicationpp

NVCC

PTX Code Virtual

PTX to TargetCompiler

Physical

G80 … GPU

© NVIDIA Corporation 2006Target code

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;

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

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

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

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

© NVIDIA Corporation 2006

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

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

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

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

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

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

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

© NVIDIA Corporation 2006

Технические спецификации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

Заключение

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

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

© NVIDIA Corporation 2006

Extra Slides

CUDA Libraries

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

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

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