§ NVIDIA CUDA – программно-аппаратная платформа для организации параллельных вычислений на графических процессорах (Graphic processing unit – GPU).


Чтобы посмотреть этот PDF файл с форматированием и разметкой, скачайте его и откройте на своем компьютере.
Лекция
11:
NVIDIA CUDA
Курносов
Михаил Георгиевич
к.т.н. доцент Кафедры вычислительных систем
Сибирский государственный университет
телекоммуникаций и информатики
NVIDIA CUDA
2
2

NVIDIA CUDA

программно
-
аппаратная
платформа для организации параллельных
вычислений на графических процессорах
(
Graphic
processing
unit

GPU
)
http://
developer.nvidia.com/cuda

NVIDIA CUDA
SDK
:
o
архитектура виртуальной машины
CUDA
o
компилятор
C/C++
o
драйвер
GPU

ОС
: GNU/Linux, Apple Mac OS X, Microsoft Windows

201
3

CUDA 5.5
(CUDA 6.0 announced)

2006

CUDA 1.0
Архитектура
NVIDIA CUDA
3
3
CUDA C/C++/Fortran
S
ource
CUDA C/C++/Fortran Compiler
(NVIDIA
nvcc
, PGI
pgfortran
)
PTX (Parallel
Thread
Execution)
(NVIDA GPU assembly language)
GPU Driver
(JIT compiler)
GPU binary code
Host (CPU)
Device (GPU)
CPU
Code
Архитектура
NVIDIA CUDA
4
4

NVIDIA C/C++ Compiler (
nvcc
)

компилятор
c
расширений
языков
C/C++ (
основан на
LLVM)
,
генерирует код для
CPU
и
GPU

NVIDA PTX Assembler
(
ptxas
)

Набор команд
PTX
развивается
:
PTX ISA 3.2 (2013), PTX ISA 3.1 (2012), …

Архитектуры
NVIDIA CUDA
o
NVIDIA Tesla (2007)
o
NVIDIA Fermi (GeForce 400 Series, 2010)
o
NVIDIA
Kepler
(GeForce 600 Series, 2012
)
o
NVIDIA Maxwell (2014)
http://
docs.nvidia.com/cuda/parallel
-
thread
-
execution/index.html
Гетерогенные вычислительные узлы
5
5
CPU1
Core1
Core2
Core3
Core4
GPU1
Cores
GPU Memory
Memory (DDR3)
Memory (DDR3)
QPI
I/O Hub
QPI
QPI
PCI Express Gen. 2
CPU1
Core1
Core2
Core3
Core4
GPU1
Cores
GPU Memory
Гетерогенные вычислительные узлы
6
6
NVIDIA Tesla K10 (2 GPU)

Процессорные ядра:
2 GPU NVIDIA GK104
(2
x 1536
ядер)

Архитектура:
NVIDIA
Kepler

RAM: 8GB GDDR5
PCI Express
3.0
Архитектура
NVIDIA
Kepler
(GK110)
7
7

15
SMX

streaming multiprocessor
(
возможны
конфигурации с 13 и 14
SMX
)

6
контроллеров памяти (
64
-
бит)

Интерфейс подключения к хосту
PCI Express 3.0
Архитектура
NVIDIA
Kepler
(GK110)
8
8
Архитектура
NVIDIA
Kepler
(GK110)
9
9
SMX

streaming
multiprocessor
Архитектура
SMX
(
GK110
)
10
10

192
ядра для выполнения операций с одинарной
точностью (
single precision float, integer
)

64
модуля двойной точности (
double precision, FMA
)

32
модуля специальных функций
(SFU)

32
модул
я
чтения
/
записи (
LD/ST
)

4 планировщика потоков (
warp schedulers
)
Архитектура
SMX
(
GK110
)
11
11
Warp scheduler (GK110)
12
12

Планировщик организует потоки в группы по 32 (
warps
)

Потоки группы выполняются одновременно

Каждый такт потоки группы (
warp
) выполняют две
независимые инструкции (допустимо совмещение
инструкций
double
и
float
)
Warp scheduler (GK110)
13
13
Организация памяти
Kepler
(
GK110
)
14
14

Каждый
SMX
имеет 64
KB
памяти:
o
48
KB shared + 16KB L1 cache
o
16KB shared + 48KB L1 cache

L2
Cache
1536KB

общий для всех
SMX

Global Memory 8GB
Архитектура
NVIDIA
Kepler
(GK110)
15
15
FERMI
GF100
FERMI
GF104
KEPLER
GK104
KEPLER
GK110
Compute Capability
2.0
2.1
3.0
3.5
Threads / Warp
32
32
32
32
Max Warps / Multiprocessor
48
48
64
64
Max Thread Blocks /
Multiprocessor
8
8
16
16
Max Threads / Thread Block
1024
1024
1024
1024
32

bit Registers /
Multiprocessor
32768
32768
65536
65536
Max Registers / Thread
63
63
63
255
Max X Grid Dimension
2^16

1
2^16

1
2^32

1
2^32

1
Hyper

Q
No
No
No
Yes
Dynamic Parallelism
No
No
No
Yes
NVIDIA
Maxwell (2014)
16
16

NVIDIA Maxwell = GPU Cores +
ARM
Core

Интегрированное
ядро
ARM
(
64
бит, проект
Denver
)
o
возможность
загрузки операционной системы на
GPU
o
поддержка унифицированной виртуальной памяти
(
device ←→ host
)
Основные понятия
CUDA
17
17

Хост (
host
)

узел с
CPU
и его память

Устройство
(device)

графический процессор
и его память

Ядро (
kernel
)

это фрагмент программы,
предназначенный для выполнения на
GPU

Пользователь самостоятельно запускает с
CPU
ядра на
GPU

Перед выполнением ядра пользователь копирует данные
из памяти хоста в память
GPU

После выполнения ядра пользователь копирует данные из
памяти
GPU
в память хоста
Основные понятия
CUDA
18
18
CPU
(Host
)
GPU
(Device
)
void
kernelA
()
{
/* Code */
}
void
kernelB
()
{
/* Code */
}
/* Serial code */
/* Serial code */
/* Serial code */
kernelA
()
kernelB
()
Выполнение
CUDA
-
программы
19
19
CPU (Host)
CPU
Memory
PCI Express
GPU (Device)
Memory

Копирование
данных из памяти
хоста в память
GPU
Выполнение
CUDA
-
программы
20
20
CPU (Host)
CPU
Memory
PCI Express
GPU (Device)
Memory

Копирование
данных из памяти
хоста в память
GPU

Загрузка и выполнение
ядра (
kernel
) в
GPU
Выполнение
CUDA
-
программы
21
21
CPU (Host)
CPU
Memory
PCI Express
GPU (Device)
Memory

Копирование
данных из памяти
хоста в память
GPU

Загрузка и выполнение
ядра (
kernel
) в
GPU

Копирование данных из
памяти
GPU
в память хоста
CUDA
HelloWorld
!
22
22
#include
stdio.h

__global__
void
mykernel
()
{
/* Parallel
GPU code (kernel)
*/
}
int
main
()
{
mykernel
1, 1, ;㄀1, ;㄀11, ;㄀();
0;
}
CUDA
HelloWorld
!
23
23
$
nvcc
-
c
-
o
prog.o
./
prog.cu
$ g
++ ./
prog.o

o
prog
-
L/opt/cuda
-
5.5/lib64
\
-
lcudart
Вычислительные ядра (
kernels
)
24
24

Спецификатор
__global__
сообщает
компилятору,
что
функция предназначена для выполнения на
GPU

Компилятор
nvcc
разделяет исходный код

ядра компилируются
nvcc
, а остальной код системным
компилятором (
gcc
, cl,

)

Тройные угловые скобки
“<<< >>>”
сообщают о вызове
ядра на
GPU
и количестве требуемых потоков

Вызов ядра (
kernel
) не блокирует выполнение потока
на
CPU

Функция
cudaThreadSynchronize
()
позволяет реализовать
ожидание завершения
ядра
Вычислительные потоки (
threads
)
25
25

Номер потока (
thread index
)

это трехкомпонентный
вектор (координаты
потока)

Потоки логически сгруппированы в одномерный,
двухмерный или трёхмерный
блок
(
thread block
)

Количество
потоков в блоке ограничено

Kepler
1024
)

Блоки распределяются по потоковым
мультипроцессорам
SMX

Предопределенные переменные
o
threadIdx
.{x, y, z}

номер потока
o
blockDim
.{x, y, z}

размерность блока
Thread Block
Thread
Thread
Thread
Thread
Thread
Thread
Вычислительные потоки (
threads
)
26
26

Блоки группируются одно
-
двух
-
и трехмерную
сетку
(grid)

Блоки распределяются по потоковым мультипроцессорам
SMX

Kepler
15
SMX
)

Предопределенные переменные
o
blockIdx
.{x, y, z}

номер блока потока
o
gridDim
.{x, y, z}

размерность
сетки
Вычислительные потоки (
threads
)
27
27
Grid of thread blocks
Thread Block
Thread
Thread
Thread
Thread
Thread
Thread
Thread Block
Thread
Thread
Thread
Thread
Thread
Thread
Thread Block
Thread
Thread
Thread
Thread
Thread
Thread
Thread Block
Thread
Thread
Thread
Thread
Thread
Thread
Thread Block
Thread
Thread
Thread
Thread
Thread
Thread
Thread Block
Thread
Thread
Thread
Thread
Thread
Thread
gridDim.y
gridDim.x
blockDim.x
blockDim.z
blockDim.y
CUDA Program
Выполнение
CUDA
-
программы
28
28
Block 0
Block 1
Block 2
Block 3
Block 4
Block 5
GPU 1
Block 0
Block 1
Block 2
Block 3
Block 4
Block 5
SMX 1
SMX 2
GPU 2
Block 0
Block 1
Block 2
Block 3
Block 4
Block 5
SMX 1
SMX 2
SMX 3
Пример:
с
ложение векторов
29
29
void
vecadd
(
float
*a,
float
*b,
float
*c,
int
n)
{
int
i;
for
(i = 0; i n; i++) {
c[i] = a[i] + b[i];
}
}
Пример:
с
ложение векторов
30
30
int
main
()
{
int
i, n = 100;
float
*a, *b, *c;
float
*deva, *devb, *devc;
a = (
float
*)malloc(sizeof(float) * n);
b = (
float
*)malloc(sizeof(float) * n);
c = (
float
*)malloc(sizeof(float) * n);
for
(i = 0; i n; i++) {
a[i] = 2.0;
b[i] = 4.0;
}
Пример:
с
ложение векторов
31
//
Выделяем память на
GPU
cudaMalloc
((void **)&deva, sizeof(float) * n);
cudaMalloc
((void **)&devb, sizeof(float) * n);
cudaMalloc
((void **)&devc, sizeof(float) * n);
//
Копируем из памяти узла в память
GPU
cudaMemcpy
(deva
, a, sizeof(float) *
n,
cudaMemcpyHostToDevice);
cudaMemcpy
(devb, b, sizeof(float) * n,
cudaMemcpyHostToDevice);
vecadd_gpu
1, n&#x-213;&#x,3 -;n30;&#x-213;&#x,3 -;n30;&#x-213;&#x,3 -;n30;(deva, devb, devc);
cudaMemcpy
(c, devc, sizeof(float) * n,
cudaMemcpyDeviceToHost
);
cudaFree
(deva);
cudaFree
(devb);
cudaFree
(devc);
free(a);
free(b);
free(c);
return
0;
}
31
Пример:
с
ложение векторов
32
__global__
void
vecadd_gpu
(float
*a,
float
*
b,
float
*c)
{
//
Каждый поток обрабатывает один элемент
int
i =
threadIdx.x
;
c[i] = a[i] + b[i];
}
32

Запускается один блок из
n
потоков (
n
= 1024
)
vecadd_gpu
1, n,-4;&#x n-4;,-4;&#x n-4;,-4;&#x n-4;(deva, devb, devc
);

Каждый поток вычисляет один элемент массива
c
Thread Block
Thread
0
Thread
1
Thread
n
-
1

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

Сложение векторов с количеством элементов

256
int
threadsPerBlock
= 256;
/* Device specific */
int
blocksPerGrid
= (n +
threadsPerBlock
-
1) /
threadsPerBlock
;
vecadd_gpu

blocksPerGrid
,
threadsPerBlock
���(deva,
devb
,
devc
, n);

Будет запущена группа блоков, в каждом блоке по
фиксированному количеству потоков

Потоков может быть больше чем элементов в массиве
Пример:
с
ложение векторов
34
34
__global__ void
vecadd_gpu
(
float
*a,
float
*b,
float
*c,
int
n)
{
int
i =
threadIdx.x
+
blockIdx.x
*
blockDim.x
;
if
(i n)
c[i] = a[i] + b[i];
}
Thread
Block
0
Thread
0
Thread
1
Thread


Thread
Block
1
Thread
0
Thread
1
Thread …

Thread
Block
2
Thread
0
Thread
1
Thread


Thread
Block
3
Thread
0
Thread
1
Thread …

Двухмерный блок потоков
35
35
dim3
threadsPerBlock
(N, N);
matrix
1,
threadsPerBlock
���(A, B, C);

Двухмерный блок потоков
(
threadIdx.x
,
threadIdx.y
)
Информация о
GPU
36
36
(0);
cudaDeviceProp
deviceProp
;
(&
deviceProp
, 0);
/*
*
deviceProp.maxThreadsPerBlock
*
deviceProp.warpSize
*
devProp.totalGlobalMem
* ...
*/
NVIDIA GeForce GTS 250
37
37
CUDA Device Query (Runtime API) version (CUDART static linking)
Device
0: "GeForce GTS 250"
CUDA Driver Version: 3.20
CUDA Runtime Version: 3.20
CUDA Capability Major/Minor version number: 1.1
Total amount of global memory: 1073020928 bytes
Multiprocessors x Cores/MP = Cores: 16 (MP) x 8 (Cores/MP) =
128
(Cores)
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 16384 bytes
Total number of registers available per block: 8192
Warp size: 32
Maximum number of threads per block: 512
Maximum sizes of each dimension of a block: 512 x 512 x 64
Maximum sizes of each dimension of a grid: 65535 x 65535 x 1
Maximum memory pitch: 2147483647 bytes
Texture alignment: 256 bytes
Clock rate: 1.91 GHz
Concurrent copy and execution: Yes
Run time limit on kernels: Yes
Support
host page
-
locked memory mapping: Yes
Compute mode: Default (multiple host
threads
can use this device simultaneously)
Concurrent kernel execution: No
Device has ECC support enabled: No
Иерархия памяти
38
NVidia
GeForce GTS 250

Global memory:
1
GB

Shared
mem
.
per block:
16KB

Constant memory:
64KB

Registers per
block: 8192
38
Data race
39
39
__
global
__
void
race
(
int
* x)
{
int
i
=
threadIdx.x
+
blockDim.x
*
blockIdx.x
;
*
x = *x + 1;
// Data race
}
int
main()
{
int
x;
// ...
race
1, 12A, ;慂䠀A, ;慂䠀8A, ;慂䠀(
d_x
);
// ...
return
0;
}
CUDA Atomics
40
40

CUDA
предоставляет
API
атомарных операций
:

atomicAdd
,
atomicSub
,
atomicInc
,
atomicDec

atomicMax
,
atomicMin

atomicExch

atomicCAS

atomicAnd
,
atomicOr
,
atomicXor
atomicOP
(
a,b
)
{
t1
= *a;
//
read
t2
= t1
OP
b;
//
modify
*a = t2
;
//
write
return
t
;
}
CUDA Atomics
41
41
__
global
__
void
race
(
int
* x)
{
int
i
=
threadIdx.x
+
blockDim.x
*
blockIdx.x
;
int
j =
atomicAdd
(x, 1);
//
j = *x; *x = j + i;
}
int
main()
{
int
x;
// ...
race
1
, 128
���(
d_x
);
// ...
return
0;
}
Умножение матриц
42
42
C = A * B

Результирующая матрица
C
разбивается на подматрицы
размером 16
x16
элементов

Подматрицы параллельно вычисляются блоками потоков

Каждый элемент подматрицы вычисляется отдельным
потоком (в блоке 16
x16 = 256
потоков)

Количество потоков = количеству элементов в матрице
C
int
main
()
{
int
block_size
= 16;
dim3
dimsA
(10 *
block_size
, 10 *
block_size
, 1);
dim3
dimsB
(20 *
block_size
, 10 *
block_size
, 1);
printf
(
"A(%
d,%d
), B(%
d,%d
)
\
n",
dimsA.x
,
dimsA.y
,
dimsB.x
,
dimsB.y
);
unsigned
int
size_A
=
dimsA.x
*
dimsA.y
;
unsigned
int
mem_size_A
=
sizeof
(float) *
size_A
;
float
*
h_A
= (float *)
malloc
(
mem_size_A
);
unsigned
int
size_B
=
dimsB.x
*
dimsB.y
;
unsigned
int
mem_size_B
=
sizeof
(float) *
size_B
;
float
*
h_B
= (float *)
malloc
(
mem_size_B
);
dim3
dimsC
(
dimsB.x
,
dimsA.y
, 1);
unsigned
int
mem_size_C
=
dimsC.x
*
dimsC.y
*
sizeof
(float);
float
*
h_C
= (float *)
malloc
(
mem_size_C
);
Умножение матриц
43
43
Умножение матриц
44
44
const
float
valB
= 0.01f;
constantInit
(
h_A
,
size_A
, 1.0f);
constantInit
(
h_B
,
size_B
, 0.01f);
float
*
d_A
, *
d_B
, *
d_C
;
cudaMalloc
((void **) &
d_A
,
mem_size_A
);
cudaMalloc
((void **) &
d_B
,
mem_size_B
);
cudaMalloc
((void **) &
d_C
,
mem_size_C
);
cudaMemcpy
(
d_A
,
h_A
,
mem_size_A
,
cudaMemcpyHostToDevice
);
cudaMemcpy
(
d_B
,
h_B
,
mem_size_B
,
cudaMemcpyHostToDevice
);
Умножение матриц
45
45
dim3
threads(
block_size
,
block_size
);
dim3
grid(
dimsB.x
/
threads.x
,
dimsA.y
/
threads.y
);
matmul_gpu
兠16grid, threadg5r;&#xid, ;&#xth6r;s00;g5r;&#xid, ;&#xth6r;s00;g5r;&#xid, ;&#xth6r;s00;s(
d_C
,
d_A
,
d_B
,
dimsA.x
,
dimsB.x
);
cudaDeviceSynchronize
();
cudaMemcpy
(
h_C
,
d_C
,
mem_size_C
,
);
Умножение матриц
46
46
free(
h_A
);
free(
h_B
);
free(
h_C
);
cudaFree
(
d_A
);
cudaFree
(
d_B
);
cudaFree
(
d_C
);
0;
}
/* main */
Умножение матриц
47
47
template

int
BLOCK_SI�ZE
__global__ void
matmul_gpu
(
float
*C,
float
*A,
float
*B,
int
wA
,
int
wB
)
{
int
bx
=
blockIdx.x
;
int
by =
blockIdx.y
;
int
tx
=
threadIdx.x
;
int
ty
=
threadIdx.y
;
int
aBegin
=
wA
* BLOCK_SIZE * by;
int
aEnd
=
aBegin
+
wA
-
1;
int
aStep
= BLOCK_SIZE;
int
bBegin
= BLOCK_SIZE *
bx
;
int
bStep
= BLOCK_SIZE *
wB
;
float
Csub
= 0;
Умножение матриц
48
48
for
(
int
a =
aBegin
, b =
bBegin
; a =
aEnd
;
a +=
aStep
, b +=
bStep
)
{
// sub
-
matrix of A
__
shared
__
float
As[BLOCK_SIZE][BLOCK_SIZE];
// sub
-
matrix of B
__
shared
__
float
Bs
[BLOCK_SIZE][BLOCK_SIZE];
// Load from device memory to shared memory
As[
ty
][
tx
] = A[a +
wA
*
ty
+
tx
];
Bs
[
ty
][
tx
] = B[b +
wB
*
ty
+
tx
];
// Synchronize (wait for loading matrices)
__
syncthreads
();
Умножение матриц
49
49
// Multiply the two matrices
#pragma unroll
for
(
int
k = 0; k BLOCK_SIZE; ++k) {
Csub
+= As[
ty
][k] *
Bs
[k][
tx
];
}
__
syncthreads
();
}
/* for
aBegin
... */
// Write the block sub
-
matrix to device memory;
int
c =
wB
* BLOCK_SIZE * by + BLOCK_SIZE *
bx
;
C[c +
wB
*
ty
+
tx
] =
Csub
;
}
Reduction
50
50
O(log
2
n
)
Mark
Harris.
Optimizing
Parallel
Reduction
in
CUDA
//
http
://
www.cuvilib.com/Reduction.pdf
Reduction v1
51
51

Условный
оператор
if
внутри цикла
приводит
к
сильному
ветвлению

Можно перераспределить
данные и
операции по
нитям
Reduction v2
52
52

Количество ветвлений
сокращено

Большое
число конфликтов
банков при
обращении
к
разделяемой
памяти
Dynamic parallelism (CUDA 5.0)
53
53
__
global
__
ChildKernel
(
void
* data)
{
// Operate on data
}
__
global
__
ParentKernel
(void *data)
{
ChildKernel
16, ̖,; 51;̖,; 51;̖,; 51;1(data);
}
// In Host Code
ParentKernel
256, 64㉖, 6;㑐㉖, 6;㑐(data);
Dynamic parallelism (CUDA 5.0)
54
54
__
global
__
RecursiveKernel
(
void
* data)
{
if
(
continueRecursion
== true)
RecursiveKernel
64, 16ͤ,; 16; ͤ,; 16; ͤ,; 16; (data);
}
Dynamic parallelism (CUDA 5.0)
55
55
Литература
56
56

CUDA by Example
//
http://
developer.download.nvidia.com/books/cuda
-
by
-
example/cuda
-
by
-
example
-
sample.pdf

Джейсон
Сандерс
, Эдвард
Кэндрот
.
Технология
CUDA
в
примерах
.
ДМК Пресс, 2011 г
.

А. В. Боресков, А. А. Харламов
.
Основы работы с
технологией CUDA
.
М.:ДМК, 2010 г.
http://www.nvidia.ru/object/cuda
-
parallel
-
computing
-
books
-
ru.html

Приложенные файлы

  • pdf 7734399
    Размер файла: 1 MB Загрузок: 0

Добавить комментарий