Порівняння кількості cuda процесорів. CUDA ми котимося: технологія NVIDIA CUDA. Історія розвитку CUDA

І призначений для трансляції host-коду (головного, керуючого коду) і device-коду (апаратного коду) (файлів з розширенням.cu) в об'єктні файли, придатні в процесі складання кінцевої програми або бібліотеки в будь-якому середовищі програмування, наприклад NetBeans.

В архітектурі CUDA використовується модель пам'яті грид, кластерне моделювання потоків та SIMD-інструкції. Застосовується не тільки високопродуктивних графічних обчислень, але й різних наукових обчислень з використанням відеокарт nVidia. Вчені та дослідники широко використовують CUDA у різних галузях, включаючи астрофізику, обчислювальну біологію та хімію, моделювання динаміки рідин, електромагнітних взаємодій, комп'ютерну томографію, сейсмічний аналіз та багато іншого. У CUDA є можливість підключення до програм, що використовують OpenGL і Direct3D. CUDA - кросплатформне програмне забезпечення для таких операційних систем як Linux, Mac OS X та Windows.

22 березня 2010 nVidia випустила CUDA Toolkit 3.0, який містив підтримку OpenCL.

Обладнання

Платформа CUDA Вперше з'явилися на ринку з виходом чіпа NVIDIA восьмого покоління G80 і стала присутня у всіх наступних серіях графічних чіпів, які використовуються в сімействах прискорювачів GeForce, Quadro та NVidia Tesla.

Перша серія обладнання, що підтримує CUDA SDK, G8x, мала 32-бітний векторний процесор одинарної точності, що використовує CUDA SDK як API (CUDA підтримує тип double мови Сі, проте зараз його точність знижена до 32-бітного з плаваючою комою). Пізніші процесори GT200 мають підтримку 64-бітної точності (тільки для SFU), але продуктивність значно гірша, ніж для 32-бітної точності (через те, що SFU всього два на кожен потоковий мультипроцесор, а скалярних процесорів – вісім). Графічний процесор організує апаратну багатопоточність, що дозволяє використовувати всі ресурси графічного процесора. Таким чином, відкривається перспектива перекласти функції фізичного прискорювача на графічний прискорювач (приклад реалізації – nVidia PhysX). Також відкриваються широкі можливості використання графічного устаткування комп'ютера до виконання складних неграфічних обчислень: наприклад, в обчислювальної біології та інших галузях науки.

Переваги

Порівняно з традиційним підходом до організації обчислень загального призначення за допомогою можливостей графічних API, архітектура CUDA має такі переваги в цій галузі:

Обмеження

• Усі функції, що виконуються на пристрої, не підтримують рекурсії (у версії CUDA Toolkit 3.1 підтримує покажчики та рекурсію) та мають деякі інші обмеження

Підтримувані GPU та графічні прискорювачі

Перелік пристроїв від виробника обладнання Nvidia із заявленою повною підтримкою технології CUDA наведено на офіційному сайті Nvidia: CUDA-Enabled GPU Products (англ.).

Фактично ж, в даний час на ринку апаратних засобів для ПК підтримку технології CUDA забезпечують наступні периферійні пристрої:

Версія специфікації	GPU	Відеокарти
1.0	G80, G92, G92b, G94, G94b	GeForce 8800GTX/Ultra, 9400GT, 9600GT, 9800GT, Tesla C/D/S870, FX4/5600, 360M, GT 420
1.1	G86, G84, G98, G96, G96b, G94b, G94b, G92b, G92b	GeForce 8400GS/GT, 8600GT/GTS, 8800GT/GTS, 9600 GSO, 9800GTX/GX2, GTS 250, GT 120/30/40, FX 4/570, 3/580, 17/18/30 /370M, 3/5/770M, 16/17/27/28/36/37/3800M, NVS420/50
1.2	GT218, GT216, GT215	GeForce 210, GT 220/40, FX380 LP, 1800M, 370/380M, NVS 2/3100M
1.3	GT200, GT200b	GeForce GTX 260, GTX 275, GTX 280, GTX 285, GTX 295, Tesla C/M1060, S1070, Quadro CX, FX 3/4/5800
2.0	GF100, GF110	GeForce (GF100) GTX 465, GTX 470, GTX 480, GTX 100, 50
2.1	GF104, GF114, GF116, GF108, GF106	GeForce 610M, GT 430, GT 440, GTS 450, GTX 460, GTX 550 Ti, GTX 560, GTX 560 Ti, 500M, Quadro 600, 2000
3.0	GK104, GK106, GK107	GeForce GTX 690, GTX 680, GTX 670, GTX 660 Ti, GTX 660, GForce GTX 680MX, GeForce GTX 675MX, GeFor 6 GeForce GT 645M, GeForce GT 640M
3.5	GK110

Nvidia GeForce для настільних комп'ютерів GeForce GTX 590 GeForce GTX 580 GeForce GTX 570 GeForce GTX 560 Ti GeForce GTX 560 GeForce GTX 550 Ti GeForce GTX 520 GeForce GTX 480 GeForce GTX 470 GeForce GTX 465 GeForce GTX 460 GeForce GTS 450 GeForce GTX 295 GeForce GTX 285 GeForce GTX 280 GeForce GTX 275 GeForce GTX 260 GeForce GTS 250 GeForce GT 240 GeForce GT 220 GeForce 210 GeForce GTS 150 GeForce GT 130 GeForce GT 120 GeForce G100 GeForce 9800 GX2 GeForce 9800 GTX+ GeForce 9800 GTX GeForce 9800 GT GeForce 9600 GSO GeForce 9600 GT GeForce 9500 GT GeForce 9400 GT GeForce 9400 mGPU GeForce 9300 mGPU GeForce 8800 GTS 512 GeForce 8800 GT GeForce 8600 GTS GeForce 8600 GT GeForce 8500 GT GeForce 8400 GS

Nvidia GeForce для мобільних комп'ютерів GeForce GTX 580M GeForce GTX 570M GeForce GTX 560M GeForce GT 555M GeForce GT 540M GeForce GT 525M GeForce GT 520M GeForce GTX 485M GeForce GTX 480M GeForce GTX 470M GeForce GTX 460M GeForce GT 445M GeForce GT 435M GeForce GT 425M GeForce GT 420M GeForce GT 415M GeForce GTX 285M GeForce GTX 280M GeForce GTX 260M GeForce GTS 360M GeForce GTS 350M GeForce GTS 160M GeForce GTS 150M GeForce GT 335M GeForce GT 330M GeForce GT 325M GeForce GT 240M GeForce GT 130M GeForce G210M GeForce G110M GeForce G105M GeForce 310M GeForce 305M GeForce 9800M GTX GeForce 9800M GT GeForce 9800M GTS GeForce 9700M GTS GeForce 9700M GT GeForce 9650M GS GeForce 9600M GT GeForce 9600M GS GeForce 9500M GS GeForce 9500M G GeForce 9300M GS GeForce 9300M G GeForce 9200M GS GeForce 9100M G GeForce 8800M GTS GeForce 8700M GT GeForce 8600M GT GeForce 8600M GS GeForce 8400M GT GeForce 8400M GS

Nvidia Tesla * Tesla C2050/C2070 Tesla M2050/M2070/M2090 Tesla S2050 Tesla S1070 Tesla M1060 Tesla C1060 Tesla C870 Tesla D870 Tesla S870

Nvidia Quadro для настільних комп'ютерів Quadro 6000 Quadro 5000 Quadro 4000 Quadro 2000 Quadro 600 Quadro FX 5800 Quadro FX 5600 Quadro FX 4800 Quadro FX 4700 X2 Quadro FX 4600 Quadro FX 3700 Quadro FX 1700 Quadro FX 570 Quadro FX 470 Quadro FX 380 Low Profile Quadro FX 370 Quadro FX 370 Low Profile Quadro CX Quadro NVS 450 Quadro NVS 420 Quadro NVS 290 Quadro Plex 2100 D4 Quadro Plex 2200 D2 Quadro Plex 2100 S4 Quadro Plex 1000 Model IV

Nvidia Quadro для мобільних комп'ютерів Quadro 5010M Quadro 5000M Quadro 4000M Quadro 3000M Quadro 2000M Quadro 1000M Quadro FX 3800M Quadro FX 3700M Quadro FX 3600M Quadro FX 2800M Quadro FX 2700M Quadro FX 1800M Quadro FX 1700M Quadro FX 1600M Quadro FX 880M Quadro FX 770M Quadro FX 570M Quadro FX 380M Quadro FX 370M Quadro FX 360M Quadro NVS 5100M Quadro NVS 4200M Quadro NVS 3100M Quadro NVS 2100M Quadro NVS 320M Quadro NVS 160M Quadro NVS 150M Quadro NVS 140M Quadro NVS 135M Quadro NVS 130M

• Моделі Tesla C1060, Tesla S1070, Tesla C2050/C2070, Tesla M2050/M2070, Tesla S2050 дозволяють проводити обчислення на GPU з подвійною точністю.

Особливості та специфікації різних версій

Feature support (unlisted features are supported for all compute capabilities)	Compute capability (version)
	1.0	1.1	1.2	1.3	2.x
32-bit words in global memory	Ні	Так
floating point values ​​in global memory
Integer atomic functions operating on 32-bit words in shared memory	Ні		Так
atomicExch() operating on 32-bit floating point values ​​in shared memory
Integer atomic functions operating on 64-bit words in global memory
Warp vote functions
Double-precision floating-point operations	Ні			Так
Atomic functions operating on 64-bit integer values ​​in shared memory	Ні				Так
Floating-point atomic addition operating на 32-bit words in global and shared memory
_ballot()
_threadfence_system()
_syncthreads_count(), _syncthreads_and(), _syncthreads_or()
Surface functions
3D grid of thread block

Technical specifications	Compute capability (version)
	1.0	1.1	1.2	1.3	2.x
Maximum dimensionality of grid of thread blocks	2				3
Maximum x-, y-, або z-dimension of grid of thread blocks	65535
Maximum dimensionality of thread block	3
Maximum x- або y-dimension of a block	512				1024
Maximum z-dimension of a block	64
Maximum number of threads per block	512				1024
Warp size	32
Maximum number of resident blocks per multiprocessor	8
Maximum number of resident warps multiprocessor	24		32		48
Maximum number of resident threads per multiprocessor	768		1024		1536
Номер 32-бітних регістрів для multiprocessor	8 K		16 K		32 K
Maximum amount of shared memory per multiprocessor	16 KB				48 KB
Number of shared memory banks	16				32
Amount of local memory per thread	16 KB				512 KB
Constant memory size	64 KB
Cache working set per multiprocessor for constant memory	8 KB
Cache working set per multiprocessor for texture memory	Device dependent, між 6 KB та 8 KB
Maximum width for 1D texture	8192				32768
Maximum width for 1D texture reference bound to linear memory	2 27
Maximum width and number of layers for a 1D layered texture reference	8192 x 512				16384 x 2048
Maximum width and height for 2D texture reference bound to linear memory or a CUDA array	65536 x 32768				65536 x 65535
Maximum width, height, and number of layers for a 2D layered texture reference	8192 x 8192 x 512				16384 x 16384 x 2048
Maximum width, height and depth для 3D texture reference bound to linear memory or a CUDA array	2048 x 2048 x 2048
Maximum number of textures that can be bound to a kernel	128
Maximum width for 1D surface reference bound to a CUDA array	Not supported				8192
Maximum width and height for a 2D surface reference bound to a CUDA array					8192 x 8192
Maximum number of surfaces that can be bound to a kernel					8
Maximum number of instructions per kernel	2 million

Приклад

CudaArray* cu_array; texture< float , 2 >tex; // Allocate array cudaMalloc( & cu_array, cudaCreateChannelDesc< float>(), width, height); // Copy image data to array cudaMemcpy( cu_array, image, width* height, cudaMemcpyHostToDevice) ; // Bind the array to the texture cudaBindTexture( tex, cu_array) ; // Run kernel dim3 blockDim (16, 16, 1); dim3 gridDim (width / blockDim.x, height / blockDim.y, 1); kernel<<< gridDim, blockDim, 0 >>> (d_odata, width, height); cudaUnbindTexture(tex); __global__ void kernel(float * odata, int height, int width) ( unsigned int x = blockIdx.x * blockDim.x + threadIdx.x ; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y ; float texfetch(tex, x, y);odata[y* width+ x] = c;

Import pycuda.driver як drv import numpy drv.init () dev = drv.Device (0 ) ctx = dev.make_context () mod = drv.SourceModule ( """ __global__ void multiply_them(float *dest, float *a, float *b) ( const int i = threadIdx.x; dest[i] = a[i] * b[i]; ) """) multiply_them = mod.get_function ("multiply_them") a = numpy.random .randn (400 ) .astype (numpy.float32 ) b = numpy.random .randn (400 ) .astype (numpy.float32 (a) multiply_them( drv.Out (dest) , drv.In (a) , drv.In (b) , block= (400 , 1 , 1 ) ) print dest-a*b

CUDA як предмет у вузах

Станом на грудень 2009 року програмна модель CUDA викладається у 269 університетах по всьому світу. У Росії навчальні курси з CUDA читаються в Санкт-Петербурзькому політехнічному університеті, Ярославському державному університеті ім. П. Г. Демидова, Московському, Нижегородському, Санкт-Петербурзькому, Тверському, Казанському, Новосибірському, Новосибірському державному технічному університеті Омському та Пермському державних університетах, Міжнародному університеті природи суспільства і людини «Дубна», Іванівському державному енергетичному університеті ім. Баумана, РХТУ ім. Менделєєва, Міжрегіональному суперкомп'ютерному центрі РАН, . Крім того, у грудні 2009 року було оголошено про початок роботи першого в Росії науково-освітнього центру «Паралельні обчислення», розташованого в місті Дубна, до завдань якого входять навчання та консультації щодо вирішення складних обчислювальних завдань на GPU.

В Україні курси з CUDA читаються у Київському інституті системного аналізу.

Посилання

Офіційні ресурси

• CUDA Zone (рус.) - офіційний сайт CUDA
• CUDA GPU Computing (англ.) – офіційні веб-форуми, присвячені обчисленням CUDA

Неофіційні ресурси

Tom's Hardware

• Дмитро Чеканов. nVidia CUDA: обчислення на відеокарті чи смерть CPU? . Tom's Hardware (22 червня 2008 р.).
• Дмитро Чеканов. nVidia CUDA: тести програм на GPU для масового ринку . Tom"s Hardware (19 травня 2009 р.). Архівовано з першоджерела 4 березня 2012 року. Перевірено 19 травня 2009 року.

iXBT.com

• Олексій Берілло. NVIDIA CUDA – неграфічні обчислення на графічних процесорах. Частина 1 . iXBT.com (23 вересня 2008 р.). Архівовано з першоджерела 4 березня 2012 року. Перевірено 20 січня 2009 року.
• Олексій Берілло. NVIDIA CUDA – неграфічні обчислення на графічних процесорах. Частина 2 . iXBT.com (22 жовтня 2008 р.). - Приклади застосування NVIDIA CUDA. Архівовано з першоджерела 4 березня 2012 року. Перевірено 20 січня 2009 року.

Інші ресурси

• Боресков Олексій Вікторович.Основи CUDA (20 січня 2009). Архівовано з першоджерела 4 березня 2012 року. Перевірено 20 січня 2009 року.
• Володимир Фролов.Введення в технологію CUDA. Мережевий журнал «Комп'ютерна графіка та мультимедіа» (19 грудня 2008 р.). Архівовано з першоджерела 4 березня 2012 року. Перевірено 28 жовтня 2009 року.
• Ігор Осколков. NVIDIA CUDA – доступний квиток у світ великих обчислень. Комп'ютери (30 квітня 2009 р.). Перевірено 3 травня 2009 року.
• Володимир Фролов.Введення у технологію CUDA (1 серпня 2009 р.). Архівовано з першоджерела 4 березня 2012 року. Перевірено 3 квітня 2010 року.
• GPGPU.ru. Використання відеокарт для обчислень
• . Центр Паралельних Обчислень

Примітки

Див. також

Nvidia
Графічні процесори	Ранні NV1 NV2 Сімейство RIVA TNT TNT2 Сімейство GeForce GeForce 256 GeForce 2 GeForce 3 GeForce 4 GeForce 6 GeForce 6 GeForce 6 GeForce 6

Пристрої перетворення персональних комп'ютерів на маленькі суперкомп'ютери відомі досить давно. Ще у 80-х роках минулого століття на ринку пропонувалися так звані трансп'ютери, які вставлялися у поширені тоді слоти розширення ISA. Спочатку їхня продуктивність у відповідних завданнях вражала, але потім зростання швидкодії універсальних процесорів прискорилося, вони посилили свої позиції в паралельних обчисленнях, і сенсу в трансп'ютерах не залишилося. Хоча подібні пристрої існують і зараз, це різноманітні спеціалізовані прискорювачі. Але найчастіше сфера їх застосування вузького та особливого поширення такі прискорювачі не отримали.

Але останнім часом естафета паралельних обчислень перейшла до масового ринку, так чи інакше пов'язаного із тривимірними іграми. Універсальні пристрої з багатоядерними процесорами для паралельних векторних обчислень, що використовуються в 3D-графіці, досягають високої пікової продуктивності, яка універсальним процесорам не під силу. Звичайно, максимальна швидкість досягається лише в ряді зручних завдань і має деякі обмеження, але такі пристрої вже почали досить широко застосовувати у сферах, для яких вони не призначалися. Відмінним прикладом такого паралельного процесора є процесор Cell, розроблений альянсом Sony-Toshiba-IBM і застосовується в ігровій приставці Sony PlayStation 3, а також всі сучасні відеокарти від лідерів ринку - компаній Nvidia і AMD.

Cell ми сьогодні не чіпатимемо, хоч він і з'явився раніше і є універсальним процесором з додатковими векторними можливостями, мова сьогодні не про нього. Для 3D-видеоприскорювачів ще кілька років тому з'явилися перші технології неграфічних розрахунків загального призначення GPGPU (General-Purpose computation on GPUs). Адже сучасні відеочіпи містять сотні математичних виконавчих блоків, і ця міць може використовуватися для значного прискорення безлічі обчислювально-інтенсивних додатків. І нинішні покоління GPU мають досить гнучку архітектуру, що разом з високорівневими мовами програмування і програмно-апаратними архітектурами, подібними до цієї статті, розкриває ці можливості і робить їх значно доступнішими.

На створення GPCPU розробників спонукало появу досить швидких та гнучких шейдерних програм, які здатні виконувати сучасні відеочіпи. Розробники задумали зробити так, щоб GPU розраховували не тільки зображення в 3D-додатках, але й застосовувалися в інших паралельних розрахунках. У GPGPU для цього використовувалися графічні API: OpenGL та Direct3D, коли дані до відеочіпа передавались у вигляді текстур, а розрахункові програми завантажувалися у вигляді шейдерів. Недоліками такого методу є порівняно висока складність програмування, низька швидкість обміну даними між CPU та GPU та інші обмеження, про які ми поговоримо далі.

Обчислення на GPU розвивалися та розвиваються дуже швидко. І надалі, два основних виробника відеочіпів, Nvidia та AMD, розробили та анонсували відповідні платформи під назвою CUDA (Compute Unified Device Architecture) та CTM (Close To Metal або AMD Stream Computing) відповідно. На відміну від попередніх моделей програмування GPU, ці були виконані з урахуванням прямого доступу до апаратних можливостей відеокарт. Платформи не сумісні між собою, CUDA – це розширення мови програмування C, а CTM – віртуальна машина, яка виконує асемблерний код. Натомість обидві платформи ліквідували деякі з важливих обмежень попередніх моделей GPGPU, які використовують традиційний графічний конвеєр та відповідні інтерфейси Direct3D або OpenGL.

Звичайно ж, відкриті стандарти, що використовують OpenGL, здаються найбільш портованими і універсальними, вони дозволяють використовувати один і той же код для відеочіпів різних виробників. Але такі методи мають масу недоліків, вони значно менш гнучкі і не такі зручні у використанні. Крім того, вони не дають використовувати специфічні можливості певних відеокарт, такі, як швидка загальна пам'ять, що розділяється, присутня в сучасних обчислювальних процесорах.

Саме тому компанія Nvidia випустила платформу CUDA – C-подібну мову програмування зі своїм компілятором та бібліотеками для обчислень на GPU. Звичайно ж, написання оптимального коду для відеочіпів зовсім не таке просте і це завдання потребує тривалої ручної роботи, але CUDA якраз і розкриває всі можливості та дає програмісту більший контроль над апаратними можливостями GPU. Важливо, що підтримка Nvidia CUDA є у чіпів G8x, G9x і GT2xx, які застосовуються у відеокартах Geforce серій 8, 9 та 200, які дуже поширені. В даний час випущена фінальна версія CUDA 2.0, в якій з'явилися нові можливості, наприклад, підтримка розрахунків з подвійною точністю. CUDA доступна на 32-бітних та 64-бітних операційних системах Linux, Windows та MacOS X.

Різниця між CPU та GPU у паралельних розрахунках

Зростання частот універсальних процесорів уперся у фізичні обмеження та високе енергоспоживання, і збільшення їхньої продуктивності все частіше відбувається за рахунок розміщення декількох ядер в одному чіпі. Процесори, що продаються зараз, містять лише до чотирьох ядер (подальше зростання не буде швидким) і вони призначені для звичайних додатків, використовують MIMD - множинний потік команд і даних. Кожне ядро ​​працює окремо від інших, виконуючи різні інструкції щодо різних процесів.

Спеціалізовані векторні можливості (SSE2 та SSE3) для чотирьохкомпонентних (одинарна точність обчислень з плаваючою точкою) та двокомпонентних (подвійна точність) векторів з'явилися в універсальних процесорах через збільшені вимоги графічних додатків, насамперед. Саме тому для певних завдань застосування GPU вигідніше, адже вони спочатку зроблено для них.

Наприклад, у відеочіпах Nvidia основний блок - це мультипроцесор з вісьма-десятьма ядрами і сотнями ALU в цілому, декількома тисячами регістрів і невеликою кількістю загальної пам'яті, що розділяється. Крім того, відеокарта містить швидку глобальну пам'ять із доступом до неї всіх мультипроцесорів, локальну пам'ять у кожному мультипроцесорі, а також спеціальну пам'ять для констант.

Найголовніше – ці кілька ядер мультипроцесора в GPU є SIMD (одинний потік команд, безліч потоків даних) ядрами. І ці ядра виконують одні й самі інструкції одночасно, такий стиль програмування є звичайним для графічних алгоритмів та багатьох наукових завдань, але потребує специфічного програмування. Проте такий підхід дозволяє збільшити кількість виконавчих блоків за рахунок їх спрощення.

Отже, перерахуємо основні різницю між архітектурами CPU і GPU. Ядра CPU створені для виконання одного потоку послідовних інструкцій з максимальною продуктивністю, а GPU проектуються для швидкого виконання великої кількості потоків інструкцій, що паралельно виконуються. Універсальні процесори оптимізовані для досягнення високої продуктивності єдиного потоку команд, що обробляє цілі числа і числа з плаваючою точкою. При цьому доступ до пам'яті є випадковим.

Розробники CPU намагаються домогтися виконання якомога більшої кількості інструкцій паралельно для збільшення продуктивності. Для цього починаючи з процесорів Intel Pentium з'явилося суперскалярне виконання, що забезпечує виконання двох інструкцій за такт, а Pentium Pro відзначився позачерговим виконанням інструкцій. Але в паралельного виконання послідовного потоку інструкцій є певні базові обмеження та збільшенням кількості виконавчих блоків кратного збільшення швидкості не досягти.

У відеочіпів робота проста і розпаралелена спочатку. Відеочіп приймає на вході групу полігонів, проводить всі необхідні операції, і на виході видає пікселі. Обробка полігонів та пікселів незалежна, їх можна обробляти паралельно, окремо один від одного. Тому, через початково паралельну організацію роботи в GPU використовується велика кількість виконавчих блоків, які легко завантажити, на відміну від послідовного потоку інструкцій для CPU. Крім того, сучасні GPU також можуть виконувати більше однієї інструкції за такт (dual issue). Так, архітектура Tesla в деяких умовах запускає виконання операції MAD+MUL або MAD+SFU одночасно.

GPU відрізняється від CPU ще й за принципами доступу до пам'яті. У GPU він пов'язаний і легко передбачуваний - якщо з пам'яті читається текстура текстури, то через деякий час прийде час і для сусідніх текселів. Та й при записі те саме - піксель записується у фреймбуфер, і через кілька тактів записуватиметься розташований поруч із ним. Тому організація пам'яті відрізняється від тієї, що використовується CPU. І відеочіпу, на відміну від універсальних процесорів, просто не потрібна кеш-пам'ять великого розміру, а для текстур потрібно лише кілька (до 128-256 в нинішніх GPU) кілобайт.

Та й сама по собі робота з пам'яттю у GPU та CPU дещо відрізняється. Так, не всі центральні процесори мають вбудовані контролери пам'яті, а у всіх GPU зазвичай є кілька контролерів, аж до восьми 64-бітних каналів в чіпі Nvidia GT200. Крім того, на відеокартах застосовується швидша пам'ять, і в результаті відеочіп доступна в рази велика пропускна здатність пам'яті, що також дуже важливо для паралельних розрахунків, що оперують з величезними потоками даних.

У універсальних процесорах великі кількості транзисторів і площа чіпа йдуть на буфери команд, апаратне передбачення розгалуження та величезні об'єми кеш-пам'яті. Всі ці апаратні блоки необхідні для прискорення виконання нечисленних потоків команд. Відеочіпи витрачають транзистори на масиви виконавчих блоків, що управляють потоками блоки, пам'ять, що розділяється, невеликого об'єму і контролери пам'яті на кілька каналів. Перераховане вище не прискорює виконання окремих потоків, воно дозволяє чіпу обробляти декількох тисяч потоків, одночасно виконуються чіпом і вимагають високої пропускної здатності пам'яті.

На відміну від кешування. Універсальні центральні процесори використовують кеш-пам'ять для збільшення продуктивності за рахунок зниження затримок доступу до пам'яті, а GPU використовують кеш або загальну пам'ять для збільшення смуги пропускання. CPU знижують затримки доступу до пам'яті за допомогою кеш-пам'яті великого розміру, а також передбачення розгалужень коду. Ці апаратні частини займають більшу частину площі чіпа і споживають багато енергії. Відеочіпи обходять проблему затримок доступу до пам'яті за допомогою одночасного виконання тисяч потоків - у той час, коли один з потоків очікує даних із пам'яті, відеочіп може виконувати обчислення іншого потоку без очікування та затримок.

Є безліч відмінностей і підтримки багатопоточності. CPU виконує 1-2 потоки обчислень одне процесорне ядро, а відеочіпи можуть підтримувати до 1024 потоків за кожен мультипроцесор, яких у чіпі кілька штук. І якщо перемикання з одного потоку на інший для CPU коштує сотні тактів, GPU перемикає кілька потоків за один такт.

Крім того, центральні процесори використовують SIMD (одна інструкція виконується над численними даними) блоки для векторних обчислень, а відеочіпи застосовують SIMT (одна інструкція та кілька потоків) для обробки скалярних потоків. SIMT не вимагає, щоб розробник перетворював дані на вектори, і допускає довільні розгалуження в потоках.

Коротко можна сказати, що на відміну від сучасних універсальних CPU відеочіпи призначені для паралельних обчислень з великою кількістю арифметичних операцій. І значно більше транзисторів GPU працює за прямим призначенням - обробці масивів даних, а чи не управляє виконанням (flow control) нечисленних послідовних обчислювальних потоків. Це схема того, скільки місця в CPU та GPU займає різноманітна логіка:

У результаті, основою для ефективного використання потужності GPU у наукових та інших неграфічних розрахунках є розпаралелювання алгоритмів на сотні виконавчих блоків, що є у відеочіпах. Наприклад, безліч додатків з молекулярного моделювання добре пристосовано для розрахунків на відеочіпах, вони вимагають високих обчислювальних потужностей і тому зручні для паралельних обчислень. А використання кількох GPU дає ще більше обчислювальних потужностей для вирішення таких завдань.

Виконання розрахунків на GPU показує відмінні результати алгоритмах, що використовують паралельну обробку даних. Тобто, коли ту саму послідовність математичних операцій застосовують до великого обсягу даних. При цьому кращі результати досягаються, якщо відношення числа арифметичних інструкцій до звернень до пам'яті досить велике. Це пред'являє менші вимоги до управління виконанням (flow control), а висока щільність математики та великий обсяг даних скасовує потребу у великих кешах, як у CPU.

В результаті всіх описаних вище відмінностей, теоретична продуктивність відеочіпів значно перевершує продуктивність CPU. Компанія Nvidia наводить такий графік зростання продуктивності CPU та GPU за останні кілька років:

Звичайно, ці дані не без частки лукавства. Адже на CPU набагато простіше на практиці досягти теоретичних цифр, та й цифри наведені для одинарної точності у разі GPU, і для подвійної – у разі CPU. У будь-якому випадку, для частини паралельних завдань одинарної точності вистачає, а різниця у швидкості між універсальними та графічними процесорами дуже велика, і тому шкурка коштує вичинки.

Перші спроби застосування розрахунків на GPU

Відеочіпи в паралельних математичних розрахунках намагалися використати досить давно. Найперші спроби такого застосування були вкрай примітивними та обмежувалися використанням деяких апаратних функцій, таких як растеризація та Z-буферизація. Але в нинішньому столітті з появою шейдерів почали прискорювати обчислення матриць. У 2003 році на SIGGRAPH окрема секція була виділена під обчислення на GPU і отримала назву GPGPU (General-Purpose computation on GPU) - універсальні обчислення на GPU).

Найбільш відомий BrookGPU - компілятор потокової мови програмування Brook, призначений для виконання неграфічних обчислень на GPU. До появи розробники, які використовують можливості відеочипів для обчислень, вибирали одне із двох поширених API: Direct3D чи OpenGL. Це серйозно обмежувало застосування GPU, адже в 3D графіку використовуються шейдери та текстури, про які фахівці з паралельного програмування знати не зобов'язані, вони використовують потоки та ядра. Brook зміг допомогти у полегшенні їхнього завдання. Ці потокові розширення до мови C, розроблені в Стендфордському університеті, приховували від програмістів тривимірний API і представляли відеочіп у вигляді паралельного співпроцесора. Компілятор обробляв файл.br із кодом C++ та розширеннями, виробляючи код, прив'язаний до бібліотеки з підтримкою DirectX, OpenGL або x86.

Звичайно, у Brook було безліч недоліків, на яких ми зупинялися, і про які ще докладніше поговоримо далі. Але навіть його поява викликала значний приплив уваги тих же Nvidia і ATI до ініціативи обчислень на GPU, оскільки розвиток цих можливостей серйозно змінило ринок надалі, відкривши цілий новий його сектор - паралельні обчислювачі на основі відеочіпів.

В подальшому деякі дослідники з проекту Brook влилися в команду розробників Nvidia, щоб представити програмно-апаратну стратегію паралельних обчислень, відкривши нову частку ринку. І головною перевагою цієї ініціативи Nvidia стало те, що розробники чудово знають усі можливості своїх GPU до дрібниць, і у використанні графічного API немає необхідності, а працювати з апаратним забезпеченням можна за допомогою драйвера. Результатом зусиль цієї команди стала Nvidia CUDA (Compute Unified Device Architecture) – нова програмно-апаратна архітектура для паралельних обчислень на Nvidia GPU, якій присвячено цю статтю.

Області застосування паралельних розрахунків на GPU

Щоб зрозуміти, які переваги приносить перенесення розрахунків на відеочіпи, наведемо середні цифри, отримані дослідниками у всьому світі. У середньому, при перенесенні обчислень на GPU, у багатьох завданнях досягається прискорення у 5-30 разів у порівнянні зі швидкими універсальними процесорами. Найбільші цифри (близько 100-кратного прискорення і навіть більше!) досягаються на коді, який дуже добре підходить для розрахунків за допомогою блоків SSE, але цілком зручний для GPU.

Це лише деякі приклади прискорень синтетичного коду на GPU проти SSE-векторизованого коду на CPU (за даними Nvidia):

• Флуоресцентна мікроскопія: 12x;
• Молекулярна динаміка (non-bonded force calc): 8-16x;
• Електростатика (пряме та багаторівневе підсумовування Кулону): 40-120x та 7x.

А це табличка, яку дуже любить Nvidia, показуючи її на всіх презентаціях, на якій ми докладніше зупинимося у другій частині статті, присвяченій конкретним прикладам практичних застосувань CUDA обчислень:

Як бачите, цифри дуже привабливі, особливо вражають 100-150-кратні прирости. У статті, присвяченій CUDA, ми докладно розберемо деякі з цих цифр. А зараз перерахуємо основні додатки, в яких зараз застосовуються обчислення на GPU: аналіз та обробка зображень та сигналів, симуляція фізики, обчислювальна математика, обчислювальна біологія, фінансові розрахунки, бази даних, динаміка газів та рідин, криптографія, адаптивна променева терапія, астрономія, обробка звуку, біоінформатика, біологічні симуляції, комп'ютерний зір, аналіз даних (data mining), цифрове кіно та телебачення, електромагнітні симуляції, геоінформаційні системи, військові застосування, гірниче планування, молекулярна динаміка, магнітно-резонансна томографія (MRI), нейромережі, океанографічні дослідження, фізика частинок, симуляція згортання молекул білка, квантова хімія, трасування променів, візуалізація, радари, гідродинамічне моделювання (reservoir simulation), штучний інтелект, аналіз супутникових даних, сейсмічна розвідка, хірургія, ультразвук, відеоконференції.

Подробиці про багато застосування можна знайти на сайті компанії Nvidia в розділі . Як бачите, список досить великий, але це ще не все! Його можна продовжувати, і напевно можна припустити, що в майбутньому будуть знайдені інші області застосування паралельних розрахунків на відеочіпах, про які ми поки не здогадуємося.

Можливості Nvidia CUDA

Технологія CUDA - це програмно-апаратна обчислювальна архітектура Nvidia, заснована на розширенні мови Сі, яка дає можливість організації доступу до набору інструкцій графічного прискорювача та управління пам'яттю при організації паралельних обчислень. CUDA допомагає реалізовувати алгоритми, виконані на графічних процесорах відеоприскорювачів Geforce восьмого покоління та старші (серії Geforce 8, Geforce 9, Geforce 200), а також Quadro та Tesla.

Хоча трудомісткість програмування GPU за допомогою CUDA досить велика, вона нижча, ніж із ранніми GPGPU рішеннями. Такі програми вимагають розбиття програми між декількома мультипроцесорами подібно до MPI програмування, але без поділу даних, які зберігаються у спільній відеопам'яті. І так як CUDA програмування для кожного мультипроцесора подібно до OpenMP програмування, воно вимагає хорошого розуміння організації пам'яті. Але, звичайно ж, складність розробки та перенесення на CUDA сильно залежить від програми.

Набір розробників містить безліч прикладів коду і добре документований. Процес навчання вимагатиме близько двох-чотирьох тижнів для тих, хто вже знайомий з OpenMP та MPI. В основі API лежить розширена мова Сі, а для трансляції коду цієї мови до складу CUDA SDK входить компілятор командного рядка nvcc, створений на основі відкритого компілятора Open64.

Перерахуємо основні характеристики CUDA:

• уніфіковане програмно-апаратне рішення для паралельних обчислень на відеочіпах Nvidia;
• великий набір рішень, що підтримуються, від мобільних до мультичипових
• стандартна мова програмування Сі;
• стандартні бібліотеки чисельного аналізу FFT (швидке перетворення Фур'є) та BLAS (лінійна алгебра);
• оптимізований обмін даними між CPU та GPU;
• взаємодія з графічними API OpenGL та DirectX;
• підтримка 32- та 64-бітних операційних систем: Windows XP, Windows Vista, Linux та MacOS X;
• можливість розробки на низький рівень.

Щодо підтримки операційних систем потрібно додати, що офіційно підтримуються всі основні дистрибутиви Linux (Red Hat Enterprise Linux 3.x/4.x/5.x, SUSE Linux 10.x), але, судячи з даних ентузіастів, CUDA чудово працює і на інших збірках: Fedora Core, Ubuntu, Gentoo та ін.

Середовище розробки CUDA (CUDA Toolkit) включає:

• компілятор nvcc;
• бібліотеки FFT та BLAS;
• профільувальник;
• налагоджувач gdb для GPU;
• CUDA runtime драйвер у комплекті стандартних драйверів Nvidia
• посібник із програмування;
• CUDA Developer SDK (початковий код, утиліти та документація).

У прикладах вихідного коду: паралельне бітонне сортування (bitonic sort), транспонування матриць, паралельне префіксне підсумовування великих масивів, згортка зображень, дискретне вейвлет-перетворення, приклад взаємодії з OpenGL і Direct3D, використання бібліотек CUBLAS та CUFFT, обчислення ціни опціону ( Шоулза, біноміальна модель, метод Монте-Карло), паралельний генератор випадкових чисел Mersenne Twister, обчислення гістограми великого масиву, шумозаглушення, фільтр Собеля (знаходження кордонів).

Переваги та обмеження CUDA

З погляду програміста графічний конвеєр є набором стадій обробки. Блок геометрії генерує трикутники, а блок растеризації - пікселі, що відображаються на моніторі. Традиційна модель програмування GPGPU виглядає так:

Щоб перенести обчислення на GPU у рамках такої моделі, потрібний спеціальний підхід. Навіть поелементне складання двох векторів вимагатиме відмальовування фігури на екрані або у позаекранний буфер. Фігура розтеризується, колір кожного пікселя визначається за заданою програмою (піксельного шейдера). Програма зчитує вхідні дані з текстур кожного пікселя, складає їх і записує у вихідний буфер. І всі ці численні операції потрібні у тому, що у звичайній мові програмування записується одним оператором!

Тому, застосування GPGPU для обчислень загального призначення має обмеження як занадто великий складності навчання розробників. Та й інших обмежень достатньо, адже піксельний шейдер — це лише формула залежності підсумкового кольору пікселя від його координати, а мова піксельних шейдерів — мова запису цих формул із Сі-подібним синтаксисом. Ранні методи GPGPU є хитрим трюком, що дозволяє використовувати потужність GPU, але без будь-якої зручності. Дані представлені зображеннями (текстурами), а алгоритм — процесом растеризации. Потрібно особливо відзначити і дуже специфічну модель пам'яті та виконання.

Програмно-апаратна архітектура для обчислень на GPU компанії Nvidia відрізняється від попередніх моделей GPGPU тим, що дозволяє писати програми для GPU справжньою мовою Сі зі стандартним синтаксисом, покажчиками та необхідністю в мінімумі розширень для доступу до обчислювальних ресурсів відеочіпів. CUDA не залежить від графічних API, і має деякі особливості, призначені спеціально для обчислень загального призначення.

Переваги CUDA перед традиційним підходом до GPGPU обчислень:

• інтерфейс програмування додатків CUDA заснований на стандартній мові програмування Сі з розширеннями, що спрощує процес вивчення та впровадження архітектури CUDA;
• CUDA забезпечує доступ до розділяється між потоками пам'яті розміром 16 Кб на мультипроцесор, яка може бути використана для організації кеша з широкою смугою пропускання, в порівнянні з текстурними вибірками;
• більш ефективна передача даних між системною та відеопам'яттю
• відсутність необхідності в графічних API з надмірністю та накладними витратами;
• лінійна адресація пам'яті, і gather і scatter; можливість запису за довільними адресами;
• апаратна підтримка цілісних та бітових операцій.

Основні обмеження CUDA:

• відсутність підтримки рекурсії для виконуваних функцій;
• мінімальна ширина блоку 32 потоку;
• закрита архітектура CUDA, що належить Nvidia.

Слабкими місцями програмування за допомогою попередніх методів GPGPU є те, що ці методи не використовують блоки виконання вершинних шейдерів у попередніх неуніфікованих архітектурах, дані зберігаються в текстурах, а виводяться у позаекранний буфер, а багатопрохідні алгоритми використовують піксельні шейдерні блоки. До обмежень GPGPU можна включити: недостатньо ефективне використання апаратних можливостей, обмеження смугою пропускання пам'яті, відсутність операції scatter (тільки gather), обов'язкове використання графічного API.

Основні переваги CUDA у порівнянні з попередніми методами GPGPU випливають з того, що ця архітектура спроектована для ефективного використання неграфічних обчислень на GPU та використовує мову програмування C, не вимагаючи перенесення алгоритмів у зручний для концепції графічного конвеєра вигляд. CUDA пропонує новий шлях обчислень на GPU, що не використовує графічні API, пропонує довільний доступ до пам'яті (scatter або gather). Така архітектура позбавлена ​​недоліків GPGPU та використовує всі виконавчі блоки, а також розширює можливості за рахунок цілісної математики та операцій бітового зсуву.

Крім того, CUDA відкриває деякі апаратні можливості, недоступні з графічних API, такі як пам'ять, що розділяється. Це пам'ять невеликого об'єму (16 кілобайтів на мультипроцесор), до якої мають доступ блоки потоків. Вона дозволяє кешувати дані, що найчастіше використовуються, і може забезпечити більш високу швидкість, в порівнянні з використанням текстурних вибірок для цього завдання. Що, своєю чергою, знижує чутливість до пропускної спроможності паралельних алгоритмів у багатьох додатках. Наприклад, це корисно для лінійної алгебри, швидкого перетворення Фур'є та фільтрів обробки зображень.

Зручніше в CUDA та доступ до пам'яті. Програмний код у графічних API виводить дані у вигляді 32-х значень з плаваючою точкою одинарної точності (RGBA значення одночасно у вісім render target) у наперед визначені області, а CUDA підтримує scatter запис - необмежену кількість записів за будь-якою адресою. Такі переваги уможливлюють виконання на GPU деяких алгоритмів, які неможливо ефективно реалізувати за допомогою методів GPGPU, заснованих на графічних API.

Також, графічні API обов'язково зберігають дані в текстурах, що вимагає попередньої упаковки великих масивів у текстури, що ускладнює алгоритм і змушує використовувати спеціальну адресацію. А CUDA дозволяє читати дані на будь-яку адресу. Ще однією перевагою CUDA є оптимізований обмін даними між CPU та GPU. А для розробників, які бажають отримати доступ до низького рівня (наприклад, під час написання іншої мови програмування), CUDA пропонує можливість низькорівневого програмування на асемблері.

Історія розвитку CUDA

Розробка CUDA була анонсована разом із чіпом G80 у листопаді 2006, а реліз публічної бета-версії CUDA SDK відбувся у лютому 2007 року. Версія 1.0 вийшла у червні 2007 року під запуск у продаж рішень Tesla, заснованих на чіпі G80, та призначених для ринку високопродуктивних обчислень. Потім наприкінці року вийшла бета-версія CUDA 1.1, яка, незважаючи на малозначне збільшення номера версії, ввела досить багато нового.

З CUDA 1.1, що з'явився, можна відзначити включення CUDA-функціональності в звичайні відеодрайвери Nvidia. Це означало, що у вимогах до будь-якої програми CUDA достатньо було вказати відеокарту серії Geforce 8 і вище, а також мінімальну версію драйверів 169.xx. Це дуже важливо для розробників, за дотримання цих умов CUDA програми працюватимуть у будь-якого користувача. Також було додано асинхронне виконання разом із копіюванням даних (тільки для чіпів G84, G86, G92 і вище), асинхронне пересилання даних у відеопам'ять, атомарні операції доступу до пам'яті, підтримка 64-бітних версій Windows та можливість мультичіпової роботи CUDA у режимі SLI.

На даний момент актуальною є версія для рішень на основі GT200 – CUDA 2.0, що вийшла разом із лінійкою Geforce GTX 200. Бета-версія була випущена ще навесні 2008 року. У другій версії з'явилися: підтримка обчислень подвійної точності (апаратна підтримка тільки у GT200), підтримується Windows Vista (32 і 64-бітні версії) і Mac OS X, додані засоби налагодження та профілювання, підтримуються 3D текстури, оптимізована пересилання даних.

Що стосується обчислень з подвійною точністю, то їх швидкість на поточному апаратному поколінні нижче одинарної точності в кілька разів. Причини розглянуті в нашій. Реалізація в GT200 цієї підтримки полягає в тому, блоки FP32 не використовуються для отримання результату в чотири рази меншому темпі, для підтримки FP64 обчислень у Nvidia вирішили зробити виділені обчислювальні блоки. І в GT200 їх удесятеро менше, ніж блоків FP32 (по одному блоку подвійної точності на кожен мультипроцесор).

Реально продуктивність може бути навіть ще меншою, тому що архітектура оптимізована для 32-бітного читання з пам'яті та регістрів, крім того, подвійна точність не потрібна в графічних додатках, і в GT200 вона зроблена швидше, щоб просто була. Та й сучасні чотириядерні процесори показують не набагато меншу реальну продуктивність. Але навіть у 10 разів повільніше, ніж одинарна точність, така підтримка корисна для схем зі змішаною точністю. Одна з найпоширеніших технік - отримати спочатку наближені результати в одинарній точності, а потім їх уточнити в подвійній. Тепер це можна зробити безпосередньо на відеокарті, без пересилання проміжних даних до CPU.

Ще одна корисна особливість CUDA 2.0 не має відношення до GPU, як не дивно. Просто тепер можна компілювати код CUDA у високоефективний багатопотоковий SSE код для швидкого виконання на центральному процесорі. Тобто тепер ця можливість годиться не тільки для налагодження, але і реального використання на системах без відеокарти Nvidia. Адже використання CUDA у звичайному коді стримується тим, що відеокарти Nvidia хоч і найпопулярніші серед виділених відеорішень, але є не у всіх системах. І до версії 2.0 у таких випадках довелося б робити два різні коди: для CUDA та окремо для CPU. А тепер можна виконувати будь-яку програму CUDA на CPU з високою ефективністю, нехай і з меншою швидкістю, ніж на відеочіпах.

Рішення з підтримкою Nvidia CUDA

Всі відеокарти, що володіють підтримкою CUDA, можуть допомогти прискорити більшість вимогливих завдань, починаючи від аудіо- та відеообробки, і закінчуючи медициною та науковими дослідженнями. Єдине реальне обмеження полягає в тому, що багато CUDA програми вимагають мінімум 256 мегабайт відеопам'яті, і це одна з найважливіших технічних характеристик для CUDA-додатків.

Актуальний список продуктів, що підтримують CUDA, можна отримати на . На момент написання статті розрахунки CUDA підтримували всі продукти серій Geforce 200, Geforce 9 і Geforce 8, у тому числі і мобільні продукти, починаючи з Geforce 8400M, а також і чіпсети Geforce 8100, 8200 і 8300. Також підтримкою всі Tesla: S1070, C1060, C870, D870 та S870.

Особливо зазначимо, що разом з новими відеокартами Geforce GTX 260 та 280, були анонсовані та відповідні рішення для високопродуктивних обчислень: Tesla C1060 та S1070 (подано на фото вище), які будуть доступні для придбання восени цього року. GPU в них застосований той же - GT200, C1060 він один, в S1070 - чотири. Натомість, на відміну від ігрових рішень, у них використовується по чотири гігабайти пам'яті на кожен чіп. З мінусів хіба що менша частота пам'яті та ПСП, ніж у ігрових карток, що забезпечує по 102 гігабайт/с на чіп.

Склад Nvidia CUDA

CUDA включає два API: високого рівня (CUDA Runtime API) та низького (CUDA Driver API), хоча в одній програмі одночасне використання обох неможливе, потрібно використовувати або один чи інший. Високорівневий працює "зверху" низькорівневого, всі виклики runtime транслюються в прості інструкції, що обробляються низькорівневим Driver API. Але навіть «високрівневий» API передбачає знання про пристрій та роботу відеочіпів Nvidia, надто високого рівня абстракції там немає.

Є ще один рівень, навіть вищий — дві бібліотеки:

CUBLAS- CUDA варіант BLAS (Basic Linear Algebra Subprograms), призначений для обчислень задач лінійної алгебри та використовує прямий доступ до ресурсів GPU;

CUFFT— CUDA варіант бібліотеки Fast Fourier Transform для розрахунку швидкого перетворення Фур'є, що широко використовується для обробки сигналів. Підтримуються такі типи перетворень: complex-complex (C2C), real-complex (R2C) та complex-real (C2R).

Розглянемо ці бібліотеки докладніше. CUBLAS — це перекладені мовою CUDA стандартні алгоритми лінійної алгебри, на даний момент підтримується лише певний набір основних функцій CUBLAS. Бібліотеку дуже легко використовувати: потрібно створити матрицю та векторні об'єкти у пам'яті відеокарти, заповнити їх даними, викликати необхідні функції CUBLAS, та завантажити результати з відеопам'яті назад у системну. CUBLAS містить спеціальні функції для створення та знищення об'єктів у пам'яті GPU, а також для читання та запису даних у цю пам'ять. Функції BLAS, що підтримуються: рівні 1, 2 і 3 для дійсних чисел, рівень 1 CGEMM для комплексних. Рівень 1 – це векторно-векторні операції, рівень 2 – векторно-матричні операції, рівень 3 – матрично-матричні операції.

CUFFT - CUDA варіант функції швидкого перетворення Фур'є - широко використовується і дуже важлива при аналізі сигналів, фільтрації і т.п. CUFFT надає простий інтерфейс для ефективного обчислення FFT на відеочіпах виробництва Nvidia без необхідності розробки власного варіанту FFT для GPU. CUDA варіант FFT підтримує 1D, 2D, та 3D перетворення комплексних та дійсних даних, пакетне виконання для декількох 1D трансформацій у паралелі, розміри 2D та 3D трансформацій можуть бути в межах , для 1D підтримується розмір до 8 мільйонів елементів.

Основи створення програм на CUDA

Для розуміння подальшого тексту слід розумітися на базових архітектурних особливостях відеочіпів Nvidia. GPU складається з кількох кластерів текстурних блоків (Texture Processing Cluster). Кожен кластер складається з укрупненого блоку текстурних вибірок та двох-трьох потокових мультипроцесорів, кожен з яких складається з восьми обчислювальних пристроїв та двох суперфункціональних блоків. Всі інструкції виконуються за принципом SIMD, коли одна інструкція застосовується до всіх потоків у warp (термін із текстильної промисловості, у CUDA це група з 32 потоків - мінімальний обсяг даних, що обробляються мультипроцесорами). Цей спосіб виконання назвали SIMT (single instruction multiple threads — одна інструкція та багато потоків).

Кожен із мультипроцесорів має певні ресурси. Так, є спеціальна пам'ять, що розділяється, обсягом 16 кілобайт на мультипроцесор. Але це не кеш, так як програміст може використовувати її для будь-яких потреб, подібно до Local Store в SPU процесорів Cell. Ця пам'ять, що розділяється, дозволяє обмінюватися інформацією між потоками одного блоку. Важливо, що всі потоки одного блоку завжди виконуються тим самим мультипроцесором. А потоки з різних блоків обмінюватись даними не можуть, і треба пам'ятати це обмеження. Пам'ять, що розділяється, часто буває корисною, крім тих випадків, коли кілька потоків звертаються до одного банку пам'яті. Мультипроцесори можуть звертатися і до відеопам'яті, але з великими затримками та гіршою пропускною здатністю. Для прискорення доступу та зниження частоти звернення до відеопам'яті, у мультипроцесорів є по 8 кілобайт кешу на константи та текстурні дані.

Мультипроцесор використовує 8192-16384 (для G8x/G9x і GT2xx, відповідно) регістру, загальні всім потоків всіх блоків, виконуваних у ньому. Максимальне число блоків однією мультипроцесор для G8x/G9x дорівнює восьми, а число warp — 24 (768 потоків однією мультипроцессор). Усього топові відеокарти серій Geforce 8 та 9 можуть обробляти до 12288 потоків одночасно. Geforce GTX 280 на основі GT200 пропонує до 1024 потоків на мультипроцесор, в ньому є 10 кластерів по три мультипроцесори, що обробляють до 30720 потоків. Знання цих обмежень дає змогу оптимізувати алгоритми під доступні ресурси.

Першим кроком при перенесенні існуючого додатка на CUDA є його профільування та визначення ділянок коду, що є «пляшковим шийкою», що гальмує роботу. Якщо серед таких ділянок є придатні для швидкого паралельного виконання, ці функції переносяться на Cі розширення CUDA для виконання на GPU. Програма компілюється за допомогою компілятора Nvidia, який генерує код і для CPU, і для GPU. При виконанні програми центральний процесор виконує свої порції коду, а GPU виконує CUDA код з найбільш важкими паралельними обчисленнями. Ця частина призначена для GPU називається ядром (kernel). У ядрі визначаються операції, які будуть виконані над даними.

Відеочип отримує ядро ​​та створює копії для кожного елемента даних. Ці копії називаються потоками (thread). Потік містить лічильник, регістри та стан. Для великих обсягів даних, таких як обробка зображень, запускаються мільйони потоків. Потоки виконуються групами по 32 штуки, званими warp"и. Warp"ам призначається виконання на певних потокових мультипроцесорах. Кожен мультипроцесор складається із восьми ядер - потокових процесорів, які виконують одну інструкцію MAD за один такт. Для виконання одного 32-потокового warp"а потрібно чотири такти роботи мультипроцесора (мова про частоту shader domain, яка дорівнює 1.5 ГГц і вище).

Мультипроцесор не є традиційним багатоядерним процесором, він відмінно пристосований для багатопоточності, підтримуючи до 32 warp"ів одночасно. Кожен такт апаратне забезпечення вибирає, який з warp"ів виконувати, і перемикається від одного до іншого без втрат у тактах. Якщо проводити аналогію з центральним процесором, це схоже на одночасне виконання 32 програм та перемикання між ними кожен такт без втрат на перемикання контексту. Реально ядра CPU підтримують одноразове виконання однієї програми та перемикаються на інші із затримкою в сотні тактів.

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

Повторимося, що CUDA використовує паралельну модель обчислень, коли кожен із SIMD процесорів виконує ту ж інструкцію над різними елементами даних паралельно. GPU є обчислювальним пристроєм, співпроцесором (device) для центрального процесора (host), що має власну пам'ять і обробляє паралельно велику кількість потоків. Ядром (kernel) називається функція для GPU, що виконується потоками (аналогія з 3D графіки – шейдер).

Ми говорили вище, що відеочіп відрізняється від CPU тим, що може обробляти одночасно десятки тисяч потоків, що зазвичай для графіки, яка добре розпаралелюється. Кожен потік скалярний не вимагає упаковки даних у 4-компонентні вектори, що зручніше для більшості завдань. Кількість логічних потоків та блоків потоків перевищує кількість фізичних виконавчих пристроїв, що дає хорошу масштабованість для всього модельного ряду рішень компанії.

Модель програмування в CUDA передбачає групування потоків. Потоки об'єднуються в блоки потоків (thread block) - одномірні або двомірні сітки потоків, що взаємодіють між собою за допомогою пам'яті, що розділяється, і точок синхронізації. Програма (ядро, kernel) виконується над сіткою (grid) блоків потоків (thread blocks), див. малюнок нижче. Одночасно виконується одна сітка. Кожен блок може бути одно-, дво- або тривимірним формою, і може складатися з 512 потоків на поточному апаратному забезпеченні.

Блоки потоків виконуються як невеликих груп, званих варп (warp), розмір яких — 32 потоку. Це мінімальний обсяг даних, які можуть бути оброблені в мультипроцессорах. І оскільки це не завжди зручно, CUDA дозволяє працювати з блоками, що містять від 64 до 512 потоків.

Угруповання блоків в сітки дозволяє уникнути обмежень і застосувати ядро ​​до більшого числа потоків за один виклик. Це допомагає і за масштабування. Якщо GPU недостатньо ресурсів, він буде виконувати послідовно блоки. У протилежному випадку блоки можуть виконуватися паралельно, що важливо для оптимального розподілу роботи на відеочіпах різного рівня, починаючи від мобільних та інтегрованих.

Модель пам'яті CUDA

Модель пам'яті в CUDA відрізняється можливістю побайтної адресації, підтримкою як gather, і scatter. Доступно досить велику кількість регістрів на кожен потоковий процесор, до 1024 штук. Доступ до них дуже швидкий, зберігати в них можна 32-бітові цілі чи числа з плаваючою точкою.

Кожен потік має доступ до наступних типів пам'яті:

Глобальна пам'ять- Найбільший обсяг пам'яті, доступний для всіх мультипроцесорів на відеочіпі, розмір становить від 256 мегабайт до 1.5 гігабайт на поточних рішеннях (і до 4 Гбайт на Tesla). Має високу пропускну здатність, більше 100 гігабайт/с для топових рішень Nvidia, але дуже великими затримками в кілька сотень тактів. Не кешується, підтримує узагальнені вказівки load і store, і звичайні покажчики на згадку.

Локальна пам'ять— це невеликий обсяг пам'яті, якого має доступ лише один потоковий процесор. Вона відносно повільна — така сама, як і глобальна.

Пам'ять, що розділяється– це 16-кілобайтний (у відеочіпах нинішньої архітектури) блок пам'яті із загальним доступом для всіх потокових процесорів у мультипроцесорі. Ця пам'ять дуже швидка, така сама, як регістри. Вона забезпечує взаємодію потоків, управляється розробником безпосередньо та має низькі затримки. Переваги пам'яті, що розділяється: використання у вигляді керованого програмістом кеша першого рівня, зниження затримок при доступі виконавчих блоків (ALU) до даних, скорочення кількості звернень до глобальної пам'яті.

Пам'ять констант- Область пам'яті об'ємом 64 кілобайти (те ж - для нинішніх GPU), доступна тільки для читання всіма мультипроцесорами. Вона кешується по 8 кілобайт на кожен процесор. Досить повільна - затримка кілька сотень тактів за відсутності потрібних даних у кеші.

Текстурна пам'ять— блок пам'яті, доступний читання усіма мультипроцессорами. Вибірка даних здійснюється за допомогою текстурних блоків відеочіпа, тому надаються можливості лінійної інтерполяції без додаткових витрат. Кешується по 8 кілобайт на кожен мультипроцесор. Повільна, як глобальна – сотні тактів затримки за відсутності даних у кеші.

Природно, що глобальна, локальна, текстурна та пам'ять констант - це фізично та сама пам'ять, відома як локальна відеопам'ять відеокарти. Їх відмінності у різних алгоритмах кешування та моделях доступу. Центральний процесор може оновлювати та запитувати лише зовнішню пам'ять: глобальну, константну та текстурну.

З написаного вище зрозуміло, що CUDA передбачає спеціальний підхід до розробки не зовсім такий, як прийнятий у програмах для CPU. Потрібно пам'ятати про різні типи пам'яті, у тому, що локальна і світова пам'ять не кешується і затримки при доступі до неї набагато вище, ніж у регістрової пам'яті, оскільки вона фізично перебуває у окремих мікросхемах.

Типовий, але не обов'язковий шаблон вирішення задач:

• завдання розбивається на підзавдання;
• вхідні дані діляться на блоки, які вміщуються в пам'ять, що розділяється;
• кожен блок обробляється блоком потоків;
• підблок підвантажується в пам'ять, що розділяється, з глобальної;
• над даними в пам'яті, що розділяється, проводяться відповідні обчислення;
• результати копіюються з пам'яті, що розділяється, назад у глобальну.

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

До складу CUDA входять runtime бібліотеки:

• загальна частина, що надає вбудовані векторні типи та підмножини викликів RTL, що підтримуються на CPU та GPU;
• CPU-компонента, для керування одним або декількома GPU;
• GPU-компонент, що надає специфічні функції для GPU.

Основний процес CUDA працює на універсальному процесорі (host), він запускає кілька копій процесів kernel на відеокарті. Код для CPU робить наступне: ініціалізує GPU, розподіляє пам'ять на відеокарті та системі, копіює константи у пам'ять відеокарти, запускає кілька копій процесів kernel на відеокарті, копіює отриманий результат із відеопам'яті, звільняє пам'ять та завершує роботу.

Як приклад для розуміння наведемо CPU код для складання векторів, представлений у CUDA:

Функції, що виконуються відеочіпом, мають такі обмеження: відсутня рекурсія, немає статичних змінних усередині функцій та змінного числа аргументів. Підтримується два види управління пам'яттю: лінійна пам'ять з доступом за 32-бітними покажчиками, і CUDA-масиви з доступом лише через функції текстурної вибірки.

Програми на CUDA можуть взаємодіяти з графічними API: для рендерингу даних, згенерованих у програмі, зчитування результатів рендерингу та їх обробки засобами CUDA (наприклад, при реалізації фільтрів постобробки). Для цього ресурси графічних API можуть відображатися (з отриманням адреси ресурсу) в простір глобальної пам'яті CUDA. Підтримуються такі типи ресурсів графічних API: Buffer Objects (PBO/VBO) в OpenGL, вершинні буфери та текстури (2D, 3D та кубічні карти) Direct3D9.

Стадії компіляції CUDA-програми:

Файли вихідного коду на CUDA C компілюються за допомогою програми NVCC, яка є оболонкою над іншими інструментами, і викликає їх: cudacc, g++, cl та ін. Сі, та об'єктний код PTX для відеочіпа. Здійснювані файли з кодом на CUDA обов'язково вимагають наявності бібліотек CUDA runtime library (cudart) і CUDA core library (cuda).

Оптимізація програм на CUDA

Звичайно, в рамках оглядової статті неможливо розглянути серйозні питання оптимізації в програмуванні CUDA. Тому просто коротко розповімо про базові речі. Для ефективного використання можливостей CUDA потрібно забути про звичайні методи написання програм для CPU, і використовувати алгоритми, які добре розпаралелюються на тисячі потоків. Також важливо знайти оптимальне місце для зберігання даних (реєстри, пам'ять, що розділяється тощо), мінімізувати передачу даних між CPU і GPU, використовувати буферизацію.

Загалом, при оптимізації програми CUDA потрібно постаратися досягти оптимального балансу між розміром і кількістю блоків. Більша кількість потоків у блоці знизить вплив затримок пам'яті, але знизить доступне число регістрів. Крім того, блок з 512 потоків неефективний, сама Nvidia рекомендує використовувати блоки по 128 або 256 потоків як компромісне значення для досягнення оптимальних затримок і кількості регістрів.

Серед основних моментів оптимізації програм CUDA: якомога більш активне використання пам'яті, що розділяється, так як вона значно швидше глобальної відеопам'яті відеокарти; Операції читання та записи з глобальної пам'яті повинні бути об'єднані (coalesced) по можливості. Для цього потрібно використовувати спеціальні типи даних для читання та запису відразу по 32/64/128 біта даних однією операцією. Якщо операції читання важко поєднати, можна спробувати використовувати текстурні вибірки.

Висновки

Представлена ​​компанією Nvidia програмно-апаратна архітектура для розрахунків на відеочіпах CUDA добре підходить для вирішення широкого кола завдань із високим паралелізмом. CUDA працює на великій кількості відеочіпів Nvidia, і покращує модель програмування GPU, значно спрощуючи її і додаючи велику кількість можливостей, таких як пам'ять, що розділяється, можливість синхронізації потоків, обчислення з подвійною точністю і цілочисленні операції.

CUDA — це доступна кожному розробнику програмного забезпечення технологія, її може використовувати будь-який програміст, який знає мову Сі. Прийдеться тільки звикнути до іншої парадигми програмування, властивої паралельним обчисленням. Але якщо алгоритм у принципі добре розпаралелюється, то вивчення та витрати часу на програмування на CUDA повернуться у багаторазовому розмірі.

Цілком імовірно, що через широке поширення відеокарт у світі, розвиток паралельних обчислень на GPU сильно вплине на індустрію високопродуктивних обчислень. Ці можливості вже викликали великий інтерес у наукових колах, та й не лише у них. Адже потенційні можливості прискорення алгоритмів, що добре піддаються розпаралелюванню (на доступному апаратному забезпеченні, що не менш важливо) відразу в десятки разів бувають не так часто.

Універсальні процесори розвиваються досить повільно, вони не мають таких стрибків продуктивності. По суті, нехай це і звучить занадто голосно, всі, хто потребує швидких обчислювачів, тепер можуть отримати недорогий персональний суперкомп'ютер на своєму столі, іноді навіть не вкладаючи додаткових коштів, так як відеокарти Nvidia широко поширені. Не кажучи вже про підвищення ефективності в термінах GFLOPS/$ і GFLOPS/Вт, які подобаються виробникам GPU.

Майбутнє безлічі обчислень явно за паралельними алгоритмами, майже всі нові рішення та ініціативи направлені в цей бік. Поки що, втім, розвиток нових парадигм знаходиться на початковому етапі, доводиться вручну створювати потоки та планувати доступ до пам'яті, що ускладнює завдання порівняно із звичним програмуванням. Але технологія CUDA зробила крок у правильному напрямку і в ній явно проглядається успішне рішення, особливо якщо Nvidia вдасться переконати якомога розробників у його користі та перспективах.

Але, звісно, ​​GPU не замінять CPU. У їхньому нинішньому вигляді вони й не призначені для цього. Тепер що відеочіпи рухаються поступово в бік CPU, стаючи все більш універсальними (розрахунки з плаваючою точкою одинарної і подвійної точності, цілочисленні обчислення), так і CPU стають все більш «паралельними», обзаводячись великою кількістю ядер, технологіями багатопоточності, не кажучи про появу блоків. SIMD та проектів гетерогенних процесорів. Швидше за все, GPU та CPU в майбутньому просто зіллються. Відомо, що багато компаній, у тому числі Intel та AMD працюють над подібними проектами. І неважливо, чи будуть GPU поглинені CPU, чи навпаки.

У статті ми здебільшого говорили про переваги CUDA. Але є й ложечка дьогтю. Один з нечисленних недоліків CUDA – слабка переносимість. Ця архітектура працює тільки на відеочіпах цієї компанії, та ще й не на всіх, а починаючи із серії Geforce 8 та 9 та відповідних Quadro та Tesla. Так, таких рішень у світі дуже багато, Nvidia наводить цифру в 90 мільйонів CUDA-сумісних відеочіпів. Це просто чудово, але конкуренти пропонують свої рішення, відмінні від CUDA. Так, у AMD є Stream Computing, у Intel у майбутньому буде Ct.

Яка з технологій переможе, стане поширеною і проживе довше за інших - покаже лише час. Але у CUDA є непогані шанси, тому що в порівнянні з Stream Computing, наприклад, вона представляє більш розвинене та зручне для використання середовище програмування звичайною мовою Сі. Можливо, у визначенні допоможе третя сторона, випустивши загальне рішення. Наприклад, в наступному оновленні DirectX під версією 11, компанією Microsoft обіцяні обчислювальні шейдери, які можуть стати якимось усередненим рішенням, що влаштовує всіх, або багатьох.

Судячи з попередніх даних, цей новий тип шейдерів запозичує багато моделей CUDA. І програмуючи в цьому середовищі вже зараз, можна отримати переваги одразу та необхідні навички для майбутнього. З точки зору високопродуктивних обчислень, DirectX також має явний недолік у вигляді поганої переносимості, оскільки цей API обмежений платформою Windows. Втім, розробляється ще один стандарт - відкрита мультиплатформенна ініціатива OpenCL, яка підтримується більшістю компаній, серед яких Nvidia, AMD, Intel, IBM та багато інших.

Не забувайте, що в наступній статті з CUDA на вас чекає дослідження конкретних практичних застосувань наукових та інших неграфічних обчислень, виконаних розробниками з різних куточків нашої планети за допомогою Nvidia CUDA.

Протягом десятиліть діяв закон Мура, який свідчить, що кожні два роки кількість транзисторів на кристалі подвоюватиметься. Однак це було в далекому 1965 році, а останні 5 років стала бурхливо розвиватися ідея фізичної багатоядерності в процесорах споживчого класу: в 2005 Intel представила Pentium D, а AMD - Athlon X2. Тоді додатків, які використовують 2 ядра, можна було перерахувати на пальцях однієї руки. Однак наступне покоління процесорів Intel, що здійснило революцію, мало саме 2 фізичні ядра. Більше того, у січні 2007 року з'явилася серія Quad, тоді ж і сам Мур зізнався, що незабаром його закон не діятиме.

Що ж зараз? Двоядерні процесори навіть у бюджетних офісних системах, а 4 фізичні ядра стало нормою і це всього за 2-3 роки. Частота процесорів не нарощується, а покращується архітектура, збільшується кількість фізичних та віртуальних ядер. Однак ідея використання відеоадаптерів, наділених десятками, а то й сотнями обчислювальних блоків витала давно.

І хоча перспективи обчислень силами GPU величезні, найбільш популярне рішення - Nvidia CUDA безкоштовно, має безліч документацій і в цілому вельми нескладне в реалізації додатків, що використовують цю технологію не так багато. Здебільшого це всілякі спеціалізовані розрахунки, яких рядовому користувачеві здебільшого немає справи. Але є і програми, розраховані на масового користувача, про них ми поговоримо в цій статті.

Спочатку трохи про саму технологію і з чим її їдять. Т.к. при написанні статті я орієнтуюсь на широке коло читачів, то й пояснити постараюся доступною мовою без складних термінів і трохи коротенько.

CUDA(англ. Compute Unified Device Architecture) - програмно-апаратна архітектура, що дозволяє проводити обчислення з використанням графічних процесорів NVIDIA, що підтримують технологію GPGPU (довільних обчислень на відеокартах). Архітектура CUDA вперше з'явилися на ринку з виходом чіпа NVIDIA восьмого покоління - G80 і є у всіх наступних серіях графічних чіпів, які використовуються в сімействах прискорювачів GeForce, Quadro і Tesla. (С) Wikipedia.org

Вхідні потоки обробляються незалежно друг від друга, тобто. паралельно.

При цьому існує поділ на 3 рівні:

Grid- Ядро. Містить один/двох/тривимірний масив блоків.

Block- Містить у собі безліч потоків (thread). Потоки різних блоків між собою не можуть взаємодіяти. Навіщо потрібно було вводити блоки? Кожен блок по суті відповідає за своє завдання. Наприклад, велике зображення (яке є матрицею) можна розбити на кілька дрібніших частин (матриць) і паралельно працювати з кожною частиною зображення.

Thread- Потік. Потоки всередині одного блоку можуть взаємодіяти або через загальну (shared) пам'ять, яка, до речі, набагато швидше за глобальну (global) пам'ять, або через засоби синхронізації потоків.

Warp– це об'єднання взаємодіючих між собою потоків, для всіх сучасних GPU розмір Warp'а дорівнює 32. Далі йде half-warp, Що є половинкою warp'a, т.к. звернення до пам'яті зазвичай йде окремо для першої та другої половини warp'a.

Як можна помітити, дана архітектура чудово підходить для розпаралелювання завдань. І хоча програмування ведеться мовою Сі з деякими обмеженнями, насправді все так просто, т.к. не все можна розпаралелити. Немає ж стандартних функцій для генерації випадкових чисел (або ініціалізації), все це доводиться реалізовувати окремо. І хоча готових варіантів є достатньо, радості все це не приносить. Можливість використання рекурсії виникла порівняно недавно.

Для наочності була написана невелика консольна (для мінімізації коду) програма, яка здійснює операції із двома масивами типу float, тобто. з нецілочисленними значеннями. З зазначених вище причин ініціалізація (заповнення масиву різними довільними значеннями) здійснювалося силами CPU. Далі з відповідними елементами кожного масиву вироблялося 25 різноманітних операцій, проміжні результати записувалися в третій масив. Змінювався розмір масиву, результати наступні:

Усього було проведено 4 тести:

1024 елементи в кожному масиві:

Наочно видно, що з такому малому кількості елементів користі від паралельних обчислень небагато, т.к. самі обчислення проходять значно швидше, ніж їх підготовка.

4096 елементів у кожному масиві:

І вже видно, що відеокарта втричі швидше проводить операції над масивами, ніж процесор. Більше того, час виконання цього тесту на відеокарті не збільшився (незначне зменшення часу можна заслати на похибку).

Тепер 12288 елементів у кожному масиві:

Відрив відеокарти збільшився ще вдвічі. Знову ж таки варто звернути увагу, що час виконання на відеокарті збільшився
незначно, тоді як на процесорі більш ніж 3 разу, тобто. пропорційно до ускладнення завдання.

І останній тест – 36864 елементи у кожному масиві:

У цьому випадку прискорення досягає значних значень - майже в 22 рази швидше на відеокарті. І знову ж таки час виконання на відеокарті збільшився незначно, а на процесорі – покладені 3 рази, що знову ж таки пропорційно до ускладнення завдання.

Якщо і далі ускладнювати обчислення, то відеокарта виграє дедалі більше. Хоч і приклад дещо утрирований, але загалом ситуацію показує наочно. Але як згадувалося вище, не все можна розпаралелити. Наприклад, обчислення числа Пі. Існують лише приклади, написані у вигляді методу Monte Carlo, але точність обчислень становить 7 знаків після коми, тобто. звичайний float. А, щоб збільшити точність обчислень необхідна довга арифметика, тоді як тут і настають проблеми, т.к. ефективно це реалізувати дуже складно. В інтернеті знайти прикладів, що використовують CUDA та розраховують число Пі до 1 мільйона знаків після коми мені не вдалося. Були спроби написати таку програму, але найпростіший і найефективніший метод розрахунку числа Пі - це алгоритм Брента - Саламіна або формула Гауса. У відомому SuperPI швидше за все (судячи зі швидкості роботи та кількості ітерацій) використовується формула Гауса. І, судячи з
тому, що SuperPI однопоточний, відсутність прикладів під CUDA і провал моїх спроб, ефективно розпаралелити підрахунок Pi неможливо.

До речі, можна помітити, як у процесі виконання обчислень підвищується навантаження на GPU, а також відбувається виділення пам'яті.

Тепер перейдемо до більш практичної користі від CUDA, а саме існуючі на даний момент програми, що використовують цю технологію. Здебільшого це всілякі аудіо/відео конвертери та редактори.

У тестуванні використовувалися 3 різні відеофайли:

* Історія створення фільму Аватар - 1920x1080, MPEG4, h.264.

*Серія "Lie to me" - 1280х720, MPEG4, h.264.

*Серія "У Філадельфії завжди сонячно" - 624х464, xvid.

Контейнер і розмір перших двох файлів був .mkv та 1,55 гб, а останнього - .avi та 272 мб.

Почнемо з дуже гучного та популярного продукту – Badaboom. Використовувалася версія 1.2.1.74 . Вартість програми складає $29.90 .

Інтерфейс програми простий і наочний – ліворуч вибираємо вихідний файл або диск, а праворуч – необхідний пристрій, для якого кодуватимемо. Є і режим користувача, в якому вручну задаються параметри, він і використовувався.

Для початку розглянемо наскільки швидко та якісно кодується відео «в собі», тобто. у той же дозвіл і приблизно той самий розмір. Швидкість вимірюватимемо у fps, а не в витраченому часі – так зручніше і порівнювати, і підраховувати скільки стискатиметься відео довільної тривалості. Т.к. сьогодні ми розглядаємо технологію «зелених», то й графіки будуть відповідні -)

Швидкість кодування залежить від якості, це очевидно. Варто відзначити, що легка роздільна здатність (назвемо його традиційно – SD) – не проблема для Badaboom – швидкість кодування в 5,5 разів перевищила вихідний (24 fps) фреймрейт відео. Та й навіть важкий 1080p відеоролик програма перетворює у реальному часі. Слід зазначити, що якість підсумкового відео дуже близьке до вихідного відеоматеріалу, тобто. кодує Badaboom дуже і дуже якісно.

Але зазвичай переганяють відео в нижчу роздільну здатність, подивимося як справи в цьому режимі. При зниженні роздільної здатності знижувався і бітрейт відео. Він становив 9500 кбіт/с для 1080p вихідного файлу, 4100 кбіт/с для 720 p та 2400 кбіт/с для 720х404. Вибір зроблений виходячи з розумного співвідношення розмір/якість.

Коментарі зайві. Якщо робити з 720p ріп до звичайної якості SD, то на перекодування фільму тривалістю 2 години піде близько 30 хвилин. І при цьому завантаження процесора буде незначним, можна займатися своїми справами, не відчуваючи дискомфорту.

А якщо перегнати відео у формат для мобільного пристрою? Для цього виберемо профіль iPhone (бітрейт 1 мбіт/с, 480х320) та подивимося на швидкість кодування:

Чи треба щось казати? Двогодинний фільм у звичайній якості для iPhone перекодується менш ніж за 15 хвилин. З HD якістю складніше, але все одно дуже швидко. Головне, що якість вихідного відео залишається на досить високому рівні при перегляді на дисплеї телефону.

Загалом враження від Badaboom позитивні, швидкість роботи радує, інтерфейс простий та зрозумілий. Різні баги ранніх версій (користувався ще бетою в 2008-му році) виліковані. Крім одного – шлях до вихідного файлу, а також до папки, в яку зберігається готове відео, не повинен містити російських букв. Але на тлі переваг програми цей недолік незначний.

Наступним на черзі у нас буде Super LoiLoScope. За звичайну його версію просять 3280 рублів, А за touch версію, що підтримує сенсорне керування в Windows 7, просять аж 4440 рублів. Спробуємо розібратися, за що розробник хоче таких грошей і навіщо відеоредактору підтримка multitouch. Використовувалася остання версія – 1.8.3.3 .

Описати інтерфейс програми словами досить складно, тому вирішив зняти невеликий відеоролик. Відразу скажу, що, як і всі відеоконвертери під CUDA, прискорення засобами GPU підтримується тільки для виведення відео MPEG4 з кодеком h.264.

Під час кодування завантаження процесора становить 100%, проте дискомфорту це не викликає. Браузер та інші не важкі програми не гальмують.

Тепер перейдемо до продуктивності. Для початку все те саме, що і з Badaboom - перекодування відео в аналогічне за якістю.

Результати набагато кращі, ніж у Badaboom. Якість на висоті, різницю з оригіналом можна помітити тільки порівнюючи попарно кадри під лупою.

Ого, а ось тут LoiloScope обходить Badaboom у 2,5 рази. При цьому можна запросто паралельно різати та кодувати інше відео, читати новини та навіть дивитися кіно, причому навіть FullHD програються без проблем, хоч завантаження процесора та максимальне.

Тепер спробуємо зробити відео для мобільного пристрою, профіль назвемо так, як він називався в Badaboom - iPhone (480x320, 1 мбіт/с):

Жодної помилки немає. Все перевірявся ще раз кілька разів, кожен раз результат був аналогічним. Швидше за все, це відбувається з тієї простої причини, що файл SD записаний з іншим кодеком і в іншому контейнері. Під час перекодування відео спочатку декодується, розбивається на матриці певного розміру, стискається. ASP декодер, що використовується у разі xvid, повільніше, ніж AVC (для h.264) при паралельному декодуванні. Однак і 192 fps - це у 8 разів швидше, ніж швидкість вихідного відео, серія тривалістю 23 хвилини стискається менш ніж за 4 хвилини. Ситуація повторювалася й іншими файлами, перетиснутими в xvid/DivX.

LoiloScopeзалишив про себе тільки приємні враження - інтерфейс, незважаючи на свою незвичність, зручний і функціональний, а швидкість роботи вища за будь-які похвали. Дещо засмучує відносно бідний функціонал, але найчастіше при простому монтажі потрібно лише трохи підкоригувати кольори, зробити плавні переходи, накласти текст, а з цим LoiloScope чудово справляється. Дещо лякає і ціна – понад $100 за звичайну версію нормально для зарубіжжя, але нам такі цифри поки що здаються дещо дикими. Хоча, зізнаюся, що якби я, наприклад, часто знімав та монтував домашнє відео, то, можливо, і задумався над покупкою. Заодно, до речі, перевірив можливість редагування HD (а точніше AVCHD) контенту прямо з відеокамери без попереднього конвертування в інший формат, у LoiloScope жодних проблем із файлами типу mts не виявлено.

Дозвольте звернутися до історії - повернутися в 2003 рік, коли Intel та AMD брали участь у спільних перегонах за найпотужніший процесор. Усього за кілька років в результаті цих перегонів тактові частоти суттєво зросли, особливо після виходу Intel Pentium 4.

Але гонка швидко наближалася до краю. Після хвилі величезного приросту тактових частот (між 2001 і 2003 роками тактова частота Pentium 4 подвоїлася з 1,5 до 3 ГГц), користувачам довелося задовольнятися десятими частками гігагерц, які змогли вичавити виробники (з 2003 до 2005 8 ГГц).

Навіть архітектури, оптимізовані під високі тактові частоти, та сама Prescott, стали відчувати труднощі, причому цього разу не тільки виробничі. Виробники чіпів просто вперлися у закони фізики. Деякі аналітики навіть пророкували, що закон Мура перестане діяти. Але цього не сталося. Оригінальний сенс закону часто спотворюють, проте стосується кількості транзисторів на поверхні кремнієвого ядра. Довгий час підвищення числа транзисторів у CPU супроводжувалося відповідним зростанням продуктивності - що призвело до спотворення сенсу. Але потім ситуація ускладнилася. Розробники архітектури CPU підійшли до закону скорочення приросту: число транзисторів, яке потрібно додати для необхідного збільшення продуктивності, ставало все більшим, заводячи в глухий кут.

Поки виробники CPU рвали на голові останнє волосся, намагаючись знайти вирішення своїх проблем, виробники GPU продовжували чудово вигравати від переваг закону Мура.

Чому ж вони не зайшли в той же глухий кут, як розробники архітектури CPU? Причина дуже проста: центральні процесори розробляються для отримання максимальної продуктивності на потоці інструкцій, які обробляють різні дані (як цілі числа, так і числа плаваючою комою), виробляють випадковий доступ до пам'яті і т.д. До цього часу розробники намагаються забезпечити більший паралелізм інструкцій - тобто виконувати якомога більше інструкцій паралельно. Так, наприклад, із Pentium з'явилося суперскалярне виконання, коли за деяких умов можна було виконувати дві інструкції за такт. Pentium Pro отримав позачергове виконання інструкцій, що дозволило оптимізувати роботу обчислювальних блоків. Проблема полягає в тому, що паралельне виконання послідовного потоку інструкцій має очевидні обмеження, тому сліпе підвищення числа обчислювальних блоків не дає виграшу, оскільки більшу частину часу вони все одно простоюватимуть.

Навпаки, робота GPU є відносно простою. Вона полягає у прийнятті групи полігонів з одного боку та генерації групи пікселів з іншого. Полігони та пікселі незалежні один від одного, тому їх можна обробляти паралельно. Таким чином, GPU можна виділити велику частину кристала на обчислювальні блоки, які, на відміну від CPU, будуть реально використовуватися.

Натисніть на зображення для збільшення.

GPU відрізняється від CPU не лише цим. Доступ до пам'яті в GPU дуже пов'язаний - якщо зчитується тексель, через кілька тактів буде зчитуватися сусідній тексель; коли записується піксель, то через кілька тактів записуватиметься сусідній. Розумно організуючи пам'ять, можна отримати продуктивність, близьку до теоретичної пропускної спроможності. Це означає, що GPU, на відміну від CPU, не вимагає величезного кешу, оскільки його роль полягає в прискоренні операцій текстурування. Все, що потрібно, це кілька кілобайт, що містять кілька текселів, що використовуються в білінійних та трилінійних фільтрах.

Натисніть на зображення для збільшення.

Хай живе GeForce FX!

Два світи тривалий час залишалися розділеними. Ми використовували CPU (або навіть кілька CPU) для офісних завдань та інтернет-застосунків, а GPU добре підходили лише для прискорення візуалізації. Але одна особливість змінила все: зокрема, поява програмованих GPU. Спочатку центральним процесорам не було чого боятися. Перші так звані програмовані GPU (NV20 та R200) ​​навряд чи становили загрозу. Число інструкцій у програмі залишалося обмеженим близько 10, вони працювали над екзотичними типами даних, такими як 9 або 12-бітними числами з фіксованою комою.

Натисніть на зображення для збільшення.

Але закон Мура знову показав себе з найкращого боку. Збільшення числа транзисторів як дозволило підвищити кількість обчислювальних блоків, а й поліпшило їх гнучкість. Поява NV30 можна вважати важливим кроком вперед з кількох причин. Звичайно, геймерам карти NV30 не дуже сподобалися, проте нові графічні процесори стали спиратися на дві особливості, які мали змінити сприйняття GPU вже не тільки як графічних акселераторів.

• Підтримка обчислень з плаваючою комою одинарної точності (нехай це навіть не відповідало стандарту IEEE754);
• підтримка числа інструкцій понад тисячу.

Ось ми і отримали всі умови, які здатні залучити дослідників-першопрохідців, які завжди бажають отримати додаткову обчислювальну потужність.

Ідея використання графічних акселераторів для математичних розрахунків не є новою. Перші спроби було зроблено ще у 90-х роках минулого століття. Звичайно, вони були дуже примітивними - обмежуючись, здебільшого, використанням деяких апаратно закладених функцій, наприклад, растеризації та Z-буферів для прискорення таких завдань, як пошук маршруту чи виведення діаграм Вороного .

Натисніть на зображення для збільшення.

У 2003 році, з появою шейдерів, що еволюціонували, була досягнута нова планка - цього разу виконання матричних обчислень. Це був рік, коли цілу секцію SIGGRAPH ("Computations on GPUs/Обчислення на GPU") було виділено під нову область ІТ. Ця рання ініціатива дістала назву GPGPU (General-Purpose computation on GPU, універсальні обчислення на GPU). І раннім поворотним моментом стала поява.

Щоб зрозуміти роль BrookGPU, необхідно розібратися, як все відбувалося до появи. Єдиним способом отримати ресурси GPU у 2003 році було використання одного з двох графічних API – Direct3D або OpenGL. Отже, розробникам, які хотіли отримати можливості GPU для своїх обчислень, доводилося спиратися на два згадані API. Проблема в тому, що вони не завжди були експертами у програмуванні відеокарт, а це серйозно ускладнювало доступ до технологій. Якщо 3D-програмісти оперують шейдерами, текстурами та фрагментами, то фахівці в галузі паралельного програмування спираються на потоки, ядра, розкиди тощо. Тому спочатку треба було навести аналогії між двома світами.

• Потік (stream)є потік елементів одного типу, в GPU він може бути представлений текстурою. У принципі, у класичному програмуванні є такий аналог як масив.
• Ядро (Kernel)- функція, яка застосовуватиметься незалежно до кожного елемента потоку; є еквівалентом піксельного шейдера. У класичному програмуванні можна навести аналогію циклу - він застосовується до великої кількості елементів.
• Щоб зчитувати результати застосування ядра до потоку, має бути створена текстура. На CPU еквівалента немає, оскільки є повний доступ до пам'яті.
• Управління місцезнаходженням у пам'яті, куди буде робитися запис (в операціях розкиду/scatter), здійснюється через вершинний шейдер, оскільки піксельний шейдер не може змінювати координати пікселя, що обробляється.

Як можна бачити, навіть з урахуванням наведених аналогій, завдання не виглядає простим. І на допомогу прийшов Brook. Під цією назвою маються на увазі розширення до мови C ("C with streams", "C з потоками"), як назвали їх розробники у Стенфорді. По суті, завдання Brook зводилася до приховування від програміста всіх складових 3D API, що дозволяло уявити GPU як співпроцесор для паралельних обчислень. Для цього компілятор Brook обробляв файл.br із кодом C++ та розширеннями, після чого генерував код C++, який прив'язувався до бібліотеки з підтримкою різних виходів (DirectX, OpenGL ARB, OpenGL NV3x, x86).

Натисніть на зображення для збільшення.

Brook має кілька заслуг, перша з яких полягає у виведенні GPGPU з тіні, щоб з цією технологією могли знайомитися і широкі маси. Хоча після оголошення про проект ряд ІТ-сайтів дуже оптимістично повідомив про те, що вихід Brook ставить під сумнів існування CPU, які незабаром будуть замінені на більш потужні GPU. Але, як бачимо, і за п'ять років цього не сталося. Чесно кажучи, ми не думаємо, що це взагалі колись станеться. З іншого боку, дивлячись на успішну еволюцію CPU, які все більше орієнтуються у бік паралелізму (більше ядер, технологія багатопоточності SMT, розширення блоків SIMD), а також і на GPU, які навпаки стають все більш універсальними (підтримка розрахунків з плаваючою комою одинарної точності, цілочисленні обчислення, підтримка розрахунків з подвійною точністю), схоже, що GPU та CPU незабаром просто зіллються. Що ж тоді станеться? Чи будуть GPU поглинені CPU, як свого часу сталося з математичними співпроцесорами? Цілком можливо. Intel та AMD сьогодні працюють над подібними проектами. Але ще дуже багато може змінитися.

Але повернемось до нашої теми. Перевага Brook полягала в популяризації концепції GPGPU, він спростив доступ до ресурсів GPU, що дозволило все більшим користувачам освоювати нову модель програмування. З іншого боку, незважаючи на всі якості Brook, чекав ще довгий шлях, перш ніж ресурси GPU можна буде використовувати для обчислень.

Одна з проблем пов'язана з різними рівнями абстракції, а також, зокрема, з надмірним додатковим навантаженням, яке створюється 3D API, яке може бути дуже відчутним. Але серйознішою можна вважати проблему сумісності, з якою розробники Brook нічого не могли зробити. Між виробниками GPU існує жорстка конкуренція, тому вони часто оптимізують свої драйвери. Якщо подібні оптимізації, здебільшого, хороші для геймерів, вони можуть відразу покінчити з сумісністю Brook. Тому складно уявити використання цього API у промисловому коді, який десь працюватиме. І довгий час Brook залишався долею дослідників-аматорів та програмістів.

Однак успіху Brook виявилося достатньо, щоб привернути увагу ATI та nVidia, у них зародився інтерес до подібної ініціативи, оскільки вона могла б розширити ринок, відкривши для компаній новий важливий сектор.

Дослідники, які спочатку залучені до проекту Brook, швидко приєдналися до команд розробників у Санта-Кларі, щоб представити глобальну стратегію для розвитку нового ринку. Ідея полягала у створенні комбінації апаратного та програмного забезпечення, що підходить для завдань GPGPU. Оскільки розробники nVidia знають усі секрети своїх GPU, то графічне API можна було й не спиратися, а зв'язуватися з графічним процесором через драйвер. Хоча, звісно, ​​у своїй виникають свої проблеми. Команда розробників CUDA (Compute Unified Device Architecture) створила набір програмних рівнів для роботи з GPU.

Натисніть на зображення для збільшення.

Як можна побачити на діаграмі, CUDA забезпечує два API.

• Високорівневий API: CUDA Runtime API;
• низькорівневий API: CUDA Driver API.

Оскільки високорівневий API реалізований над низькорівневим, кожен виклик функції рівня Runtime розбивається на простіші вказівки, які обробляє Driver API. Зверніть увагу, що два API взаємно виключають один одного: програміст може використовувати один або інший API, але змішувати дзвінки функцій двох API не вдасться. Взагалі термін "високорівневий API" відносний. Навіть Runtime API такий, що багато хто вважає його низькорівневим; втім, він все ж таки надає функції, дуже зручні для ініціалізації або управління контекстом. Але не чекайте особливо високого рівня абстракції - вам все одно потрібно мати гарний набір знань про nVidia GPU і про те, як вони працюють.

З Driver API працювати ще складніше; для запуску обробки на GPU вам знадобиться більше зусиль. З іншого боку, низькорівневий API більш гнучкий, надаючи програмісту додатковий контроль, якщо потрібно. Два API здатні працювати з ресурсами OpenGL або Direct3D (тільки дев'ята версія на сьогодні). Користь від такої можливості очевидна - CUDA може використовуватися для створення ресурсів (геометрія, процедурні текстури тощо), які можна передати на графічне API або, навпаки, можна зробити так, що 3D API надсилатиме результати рендерингу програмі CUDA, яка, у свою чергу, виконуватиме пост-обробку. Є багато прикладів таких взаємодій і перевага полягає в тому, що ресурси продовжують зберігатися в пам'яті GPU, їх не потрібно передавати через шину PCI Express, яка, як і раніше, залишається "вузьким місцем".

Втім, слід зазначити, що спільне використання ресурсів у відеопам'яті не завжди проходить ідеально і може призвести до деяких головних болів. Наприклад, при зміні роздільної здатності або глибини кольору, графічні дані є пріоритетними. Тому якщо потрібно збільшити ресурси в кадровому буфері, то драйвер без проблем зробить це рахунок ресурсів додатків CUDA, які просто "вилетять" з помилкою. Звичайно, не дуже елегантно, але така ситуація не повинна траплятися дуже часто. І якщо вже ми почали говорити про недоліки: якщо ви хочете використовувати кілька GPU для програм CUDA, то вам потрібно спочатку відключити режим SLI, інакше програми CUDA зможуть "бачити" лише один GPU.

Нарешті, третій програмний рівень відданий бібліотекам - двом, якщо точним.

• CUBLAS де є необхідні блоки для обчислень лінійної алгебри на GPU;
• CUFFT, яка підтримує розрахунок перетворень Фур'є - алгоритм, який широко використовується в області обробки сигналів.

Перед тим, як ми поринемо в CUDA, дозвольте визначити ряд термінів, розкиданих за документацією nVidia. Компанія вибрала вельми специфічну термінологію, до якої важко звикнути. Насамперед, зазначимо, що потік (thread)в CUDA має далеко не таке ж значення, як потік CPU, а також не є еквівалентом потоку в наших статтях про GPU. Потік GPU у разі є базовий набір даних, які потрібно обробити. На відміну від потоків CPU, потоки CUDA дуже "легкі", тобто перемикання контексту між двома потоками - аж ніяк не ресурсомістка операція.

Другий термін, який часто зустрічається в документації CUDA - варп (warp). Тут плутанини немає, оскільки в російській аналога не існує (хіба що ви не є фанатом Start Trek або ігри Warhammer). Насправді термін взятий із текстильної промисловості, де через основну пряжу (warp yarn), яка розтягнута на верстаті, простягається уточна пряжа (weft yarn). Варп в CUDA є групою з 32 потоків і є мінімальним обсягом даних, оброблюваних SIMD-способом в мультипроцесорах CUDA.

Але подібна "зернистість" який завжди зручна для програміста. Тому в CUDA, замість роботи з варпами безпосередньо, можна працювати з блоками/block, що містять від 64 до 512 потоків

Нарешті, ці блоки збираються разом у сітки/grid. Перевага подібного угруповання полягає в тому, що число блоків, що одночасно обробляються GPU, тісно пов'язане з апаратними ресурсами, як ми побачимо нижче. Угруповання блоків у сітки дозволяє повністю абстрагуватися від цього обмеження та застосувати ядро/kernel до більшого числа потоків за один виклик, не думаючи про фіксовані ресурси. За це відповідають бібліотеки CUDA. Крім того, подібна модель добре масштабується. Якщо GPU має мало ресурсів, він буде виконувати блоки послідовно. Якщо число обчислювальних процесорів велике, блоки можуть виконуватися паралельно. Тобто, той самий код може працювати на GPU як початкового рівня, так і на топових і навіть майбутніх моделях.

Є ще пара термінів у CUDA API, які позначають CPU ( хост/host) та GPU ( пристрій/device). Якщо це невелике введення вас не злякало, настав час ближче познайомитися з CUDA.

Якщо ви регулярно читаєте Tom"s Hardware Guide, то архітектура останніх GPU від nVidia вам знайома. Якщо ні, ми рекомендуємо ознайомитись зі статтею " nVidia GeForce GTX 260 та 280: нове покоління відеокартЩо стосується CUDA, то nVidia представляє архітектуру дещо по-іншому, демонструючи деякі деталі, що раніше залишалися прихованими.

Як можна бачити по ілюстрації вище, ядро ​​шейдерів nVidia складається з кількох кластерів текстурних процесорів (Texture Processor Cluster, TPC). Відеокарта 8800 GTX, наприклад, використала вісім кластерів, 8800 GTS – шість і т.д. Кожен кластер, по суті, складається з текстурного блоку та двох потокових мультипроцесорів (streaming multiprocessor). Останні включають початок конвеєра (front end), що виконує читання та декодування інструкцій, а також відсилання їх на виконання, і кінець конвеєра (back end), що складається з восьми обчислювальних пристроїв та двох суперфункціональних пристроїв SFU (Super Function Unit)де інструкції виконуються за принципом SIMD, тобто одна інструкція застосовується до всіх потоків у варпі. nVidia називає такий спосіб виконання SIMT(Single instruction multiple threads, одна інструкція, багато потоків). Важливо відзначити, що кінець конвеєра працює на частоті вдвічі перевершує його початок. Насправді це означає, що ця частина виглядає вдвічі "ширше", ніж вона є насправді (тобто як 16-канальний блок SIMD замість восьмиканального). Потокові мультипроцесори працюють так: кожен такт початок конвеєра вибирає варп, готовий до виконання, і запускає виконання інструкції. Щоб інструкція застосувалася до всіх 32 потоків у варпі, кінцю конвеєра знадобиться чотири такти, але оскільки він працює на подвоєній частоті в порівнянні з початком, знадобиться лише два такти (з точки зору початку конвеєра). Тому, щоб початок конвеєра не простоювало такт, а апаратне забезпечення було максимально завантажено, в ідеальному випадку можна чергувати інструкції кожен такт – класична інструкція в один такт та інструкція для SFU – в інший.

Кожен мультипроцесор має певний набір ресурсів, у яких варто розібратися. Є невелика область пам'яті під назвою "Спільна пам'ять/Shared Memory", по 16 кбайт на мультипроцесор Це аж ніяк не кеш-пам'ять: програміст може використовувати її на власний розсуд. Тобто перед нами щось близьке до Local Store у SPU на процесорах Cell. Ця деталь дуже цікава, оскільки вона підкреслює, що CUDA - це комбінація програмних та апаратних технологій. Ця область пам'яті не використовується для піксельних шейдерів, що nVidia дотепно підкреслює "нам не подобається, коли пікселі розмовляють один з одним".

Ця область пам'яті відкриває можливість обміну інформацією між потоками в одному блоці. Важливо наголосити на цьому обмеженні: всі потоки в блоці гарантовано виконуються одним мультипроцесором. Навпаки, прив'язка блоків до різних мультипроцесорів взагалі не обговорюється, і два потоки з різних блоків не можуть обмінюватися інформацією між собою під час виконання. Тобто користуватись спільною пам'яттю не так і просто. Втім, загальна пам'ять все ж таки виправдана за винятком випадків, коли кілька потоків спробують звернутися до одного банку пам'яті, викликаючи конфлікт. В інших ситуаціях доступ до загальної пам'яті такий самий швидкий, як і до регістрів.

Загальна пам'ять – не єдина, до якої можуть звертатися мультипроцесори. Вони можуть використовувати відеопам'ять, але з меншою пропускною здатністю та більшими затримками. Тому, щоб знизити частоту звернення до цієї пам'яті, nVidia оснастила мультипроцесори кешем (приблизно 8 кбайт на мультипроцесор), що зберігає константи та текстури.

Мультипроцесор має 8 192 регістри, які загальні всім потоків всіх блоків, активних на мультипроцессоре. Число активних блоків на мультипроцесор не може перевищувати восьми, а кількість активних варпів обмежена 24 (768 потоків). Тому 8800 GTX може обробляти до 12288 потоків в один момент часу. Всі ці обмеження варто згадати, оскільки вони дозволяють оптимізувати алгоритм залежно від доступних ресурсів.

Оптимізація програми CUDA, таким чином, полягає у отриманні оптимального балансу між кількістю блоків та їх розміром. Більше потоків на блок будуть корисні зниження затримок роботи з пам'яттю, а й кількість регістрів, доступних потік, зменшується. Більш того, блок з 512 потоків буде неефективним, оскільки на мультипроцесорі може бути активним лише один блок, що призведе до втрати 256 потоків. Тому nVidia рекомендує використовувати блоки по 128 або 256 потоків, що дає оптимальний компроміс між зниженням затримок та кількістю регістрів більшості ядер/kernel.

З програмної точки зору CUDA складається із набору розширень до мови C, що нагадує BrookGPU, а також кількох специфічних викликів API. Серед розширень присутні специфікатори типу, що належать до функцій та змінних. Важливо запам'ятати ключове слово __global__, яке, будучи наведеним перед функцією, показує, що остання відноситься до ядра/kernel - цю функцію викликатиме CPU, а виконуватиметься вона на GPU. Префікс __device__вказує, що функція буде виконуватися на GPU (який, до речі, CUDA і називає "пристрій/device"), але вона може бути викликана тільки з GPU (іншими словами, з іншої функції __device__ або з функції __global__). Зрештою, префікс __host__опціональний, він позначає функцію, яка викликається CPU і виконується CPU - тобто звичайну функцію.

Є ряд обмежень, пов'язаних з функціями __device__ і __global__: вони не можуть бути рекурсивними (тобто викликати себе), і не можуть мати змінну кількість аргументів. Нарешті, оскільки функції __device__ розташовуються у просторі пам'яті GPU, цілком логічно, що отримати їхню адресу не вдасться. Змінні теж мають низку кваліфікаторів, які вказують на область пам'яті, де вони зберігатимуться. Змінна з префіксом __shared__означає, що вона зберігатиметься у спільній пам'яті потокового мультипроцесора. Виклик функції __global__ трохи відрізняється. Справа в тому, при виклику потрібно задати конфігурацію виконання - більш конкретно розмір сітки/grid, до якої буде застосовано ядро/kernel, а також розмір кожного блоку. Візьмемо, наприклад, ядро ​​з наступним підписом.

__global__ void Func(float* parameter);

Воно буде викликатись у вигляді

Func<<< Dg, Db >>> (Parameter);

де Dg є розміром сітки, а Db – розміром блоку. Дві ці змінні відносяться до нового типу вектора, що з'явився з CUDA.

API CUDA містить функції для роботи з пам'яттю у VRAM: cudaMalloc для виділення пам'яті, cudaFree для звільнення та cudaMemcpy для копіювання пам'яті між RAM та VRAM та навпаки.

Ми закінчимо цей огляд дуже цікавим способом, яким компілюється програма CUDA: компіляція виконується у кілька етапів. Спочатку виймається код, що відноситься до CPU, що передається стандартному компілятору. Код, призначений для GPU, спочатку перетворюється на проміжну мову PTX. Він подібний до асемблеру і дозволяє вивчати код у пошуках потенційних неефективних ділянок. Нарешті, остання фаза полягає у трансляції проміжної мови в специфічні команди GPU та створення двійкового файлу.

Переглянувши документацію nVidia, так хочеться спробувати CUDA на тижні. Справді, що може бути краще за оцінку API шляхом створення власної програми? Саме тоді більшість проблем повинні виплисти на поверхню, нехай навіть на папері виглядає ідеально. Крім того, практика найкраще покаже, як добре ви зрозуміли всі принципи, викладені в документації CUDA.

У подібний проект поринути досить легко. Сьогодні для завантаження доступна велика кількість безкоштовних, але якісних інструментів. Для нашого тесту ми використовували Visual C++ Express 2005 де є все необхідне. Найскладніше полягати в тому, щоб знайти програму, портування якої на GPU не зайняло б кілька тижнів, і водночас вона була б досить цікавою, щоб наші зусилля не пропали даремно. Зрештою, ми вибрали відрізок коду, який бере карту висот та розраховує відповідну карту нормалей. Ми не детально заглиблюватимемося в цю функцію, оскільки в даній статті це навряд чи цікаво. Якщо бути коротким, то програма займається викривленням ділянок: для кожного пікселя початкового зображення ми накладаємо матрицю, що визначає колір результуючого пікселя в зображенні прилеглих пікселів, використовуючи більш-менш складну формулу. Перевага цієї функції в тому, що її дуже легко розпаралелити, тому цей тест чудово показує можливості CUDA.

Ще одна перевага полягає в тому, що у нас вже є реалізація на CPU, тому ми можемо порівнювати її результат із версією CUDA – і не винаходити колесо наново.

Ще раз повторимо, що метою тесту було знайомство з утилітами CUDA SDK, а не порівняльне тестування версій під CPU та GPU. Оскільки це була перша наша спроба створення програми CUDA, ми не дуже сподівалися отримати високу продуктивність. Так як дана частина коду не є критичною, то версія під CPU не оптимізована, тому пряме порівняння результатів навряд чи цікаво.

Продуктивність

Однак ми заміряли час виконання, щоб подивитися, чи є перевага у використанні CUDA навіть з найбрутальнішою реалізацією, чи нам буде потрібно тривала та стомлююча практика, щоб отримати якийсь виграш при використанні GPU. Тестова машина була взята з нашої лабораторії розробки – ноутбук з процесором Core 2 Duo T5450 та відеокартою GeForce 8600M GT, що працює під Vista. Це далеко не суперкомп'ютер, але результати дуже цікаві, оскільки тест не заточений під GPU. Завжди приємно бачити, коли nVidia демонструє величезний приріст на системах з монстроподібними GPU та чималою пропускною спроможністю, але на практиці багато хто з 70 мільйонів GPU з підтримкою CUDA на сучасному ринку ПК далеко не такі потужні, тому і наш тест має право на життя.

Для зображення 2 048 x 2 048 пікселів ми отримали такі результати.

• CPU 1 потік: 1419 мс;
• CPU 2 потоку: 749 мс;
• CPU 4 потоки: 593 мс
• GPU (8600M GT) блоки по 256 потоків: 109 мс;
• GPU (8600M GT) блоки по 128 потоків: 94 мс;
• GPU (8800 GTX) блоки по 128 потоків/256 потоків: 31 мс.

За наслідками можна зробити кілька висновків. Почнемо з того, що, незважаючи на розмови про очевидну лінь програмістів, ми модифікували початкову версію CPU під кілька потоків. Як ми вже згадували, код ідеальний для цієї ситуації - все, що потрібно, це розбити початкове зображення на стільки зон, скільки потоків. Зверніть увагу, що від переходу від одного потоку на два на нашому двоядерному CPU прискорення вийшло майже лінійним, що також вказує на паралельну природу тестової програми. Дуже несподівано, але версія із чотирма потоками теж виявилася швидше, хоча на нашому процесорі це дуже дивно - можна було, навпаки, очікувати падіння ефективності через накладні витрати на управління додатковими потоками. Як можна пояснити такий результат? Важко сказати, але, можливо, винний планувальник потоків під Windows; у будь-якому разі, результат повторюємо. З текстурами меншого розміру (512x512) приріст від поділу на потоки був не такий виражений (приблизно 35% проти 100%), і поведінка версії з чотирма потоками була логічніше, без приросту порівняно з версією на два потоки. GPU працював все ще швидше, але вже не так виражено (8600M GT була втричі швидше ніж версія з двома потоками).

Натисніть на зображення для збільшення.

Друге значуще спостереження - навіть найповільніша реалізація GPU виявилася майже вшестеро швидше, ніж найпродуктивніша версія CPU. Для першої програми та неоптимізованої версії алгоритму результат дуже навіть підбадьорливий. Зверніть увагу, що ми отримали відчутно найкращий результат на невеликих блоках, хоча інтуїція може підказувати про інше. Пояснення просте - наша програма використовує 14 регістрів на потік, і з 256-потоковими блоками потрібно 3584 регістру на блок, а для повного навантаження процесора потрібно 768 потоків, як ми показували. У нашому випадку це становить три блоки або 10572 регістру. Але мультипроцесор має всього 8192 регістру, тому він може підтримувати активними тільки два блоки. Навпаки, з блоками по 128 потоків нам потрібно 1792 регістра на блок; якщо 8192 поділити на 1792 і округлити до найближчого цілого, то ми отримаємо чотири блоки. На практиці число потоків буде таким самим (512 на мультипроцесор, хоча для повного навантаження теоретично потрібно 768), але збільшення числа блоків дає GPU перевагу гнучкості по доступу до пам'яті - коли йде операція з великими затримками, то можна запустити виконання інструкцій іншого блоку, чекаючи надходження результатів. Чотири блоки явно знижують затримки, особливо з огляду на те, що наша програма використовує кілька доступів на згадку.

Аналіз

Нарешті, незважаючи на те, що ми сказали вище, ми не змогли встояти перед спокусою і запустили програму на 8800 GTX, яка виявилася втричі швидше за 8600, незалежно від розміру блоків. Можна подумати, що на практиці на відповідних архітектурах результат буде в чотири або більше разів вищий: 128 АЛУ/шейдерних процесорів проти 32 і вища тактова частота (1,35 ГГц проти 950 МГц), але так не вийшло. Швидше за все, обмежуючим фактором виявився доступ до пам'яті. Якщо бути точнішим, доступ до початкового зображення здійснюється як до багатовимірного масиву CUDA - дуже складний термін для того, що є не більш ніж текстурою. Але є кілька переваг.

• доступи виграють від кешу текстури;
• ми використовуємо wrapping mode, в якому не потрібно обробляти межі зображення на відміну від версії CPU.

Крім того, ми можемо отримати перевагу від "безкоштовної" фільтрації з нормалізованою адресацією між замість , але в нашому випадку це навряд чи корисно. Як ви знаєте, 8600 оснащений 16 текстурними блоками порівняно з 32 у 8800 GTX. Тому між двома архітектурами співвідношення лише два до одного. Додайте до цього різницю в частотах, і ми отримаємо співвідношення (32 x 0,575)/(16 x 0,475) = 2,4 - близько до "трьома до одного", що ми отримали насправді. Ця теорія також пояснює, чому розмір блоків багато на G80 не змінює, оскільки АЛУ все одно упирається в текстурні блоки.

Натисніть на зображення для збільшення.

Окрім перспективних результатів, наше перше знайомство з CUDA пройшло дуже добре, враховуючи не найсприятливіші вибрані умови. Розробка на ноутбуці під Vista має на увазі, що доведеться використовувати CUDA SDK 2.0, що все ще знаходиться в стані бета-версії, з драйвером 174.55, який теж бета-версія. Незважаючи на це ми не можемо повідомити про якісь неприємні сюрпризи - тільки початкові помилки під час першої налагодження, коли наша програма, все ще дуже "глючна", спробувала адресувати пам'ять за межами виділеного простору.

Монітор почав дико мерехтіти, потім екран почорнів... доки Vista не запустила службу відновлення драйвера, і все стало гаразд. Але все ж таки трохи дивно це спостерігати, якщо ви звикли бачити типову помилку Segmentation Fault на стандартних програмах, подібно до нашої. Нарешті, невелика критика у бік nVidia: у всій документації, доступної для CUDA, немає невеликого керівництва, яке крок за кроком розповідало про те, як налаштувати оточення розробки під Visual Studio. Власне, проблема невелика, оскільки у SDK є повний набір прикладів, які можна вивчити для розуміння каркасу для додатків CUDA, але керівництво для новачків не завадило б.

Натисніть на зображення для збільшення.

nVidia представила CUDA з випуском GeForce 8800. І в той час обіцянки здавались дуже спокусливими, але ми притримали свій ентузіазм до реальної перевірки. Справді, на той час це здавалося більше розміткою території, щоб залишатися на хвилі GPGPU. Без доступного SDK складно сказати, що перед нами не чергова маркетингова пустушка, з якої нічого не вийде. Вже не вперше хороша ініціатива була оголошена дуже рано і на той час не вийшла на світ через брак підтримки - особливо в такому конкурентному секторі. Тепер, через півтора роки після оголошення, ми з упевненістю можемо сказати, що nVidia дотрималася слова.

SDK досить швидко з'явився в бета-версії на початку 2007 року, з того часу він швидко оновлювався, що доводить важливість цього проекту для nVidia. Сьогодні CUDA дуже приємно розвивається: SDK доступний вже у бета-версії 2.0 для основних операційних систем (Windows XP і Vista, Linux, а також 1.1 для Mac OS X), а для розробників nVidia виділила цілий розділ сайту.

На більш професійному рівні враження від перших кроків з CUDA виявилося дуже позитивним. Якщо ви навіть знайомі з архітектурою GPU, ви легко розберетеся. Коли API виглядає зрозумілим з першого погляду, то відразу починаєш вважати, що отримаєш переконливі результати. Але чи не втрачатиметься обчислювальний час від численних передач із CPU на GPU? І як використовувати ці тисячі потоків майже без примітива синхронізації? Ми розпочинали наші експерименти з усіма цими побоюваннями в умі. Але вони швидко розвіялися, коли перша версія нашого алгоритму, нехай і досить тривіального, виявилася значно швидшою, ніж на CPU.

Так що CUDA – це не "паличка-виручалочка" для дослідників, які хочуть переконати керівництво університету купити ним GeForce. CUDA – вже повністю доступна технологія, яку може використовувати будь-який програміст зі знанням C, якщо він готовий витратити час та зусилля на звикання до нової парадигми програмування. Ці зусилля не будуть втрачені даремно, якщо ваші алгоритми добре розпаралелюються. Також ми хотіли б подякувати nVidia за надання повної та якісної документації, де знайдуть відповіді програмісти CUDA-початківці.

Що ж потрібно CUDA, щоб стати відомим API? Якщо казати одним словом: переносимість. Ми знаємо, що майбутнє ІТ криється в паралельних обчисленнях – сьогодні вже кожен готується до подібних змін, і всі ініціативи, як програмні, так і апаратні, спрямовані у цьому напрямі. Однак на даний момент, якщо дивитися на розвиток парадигм, ми ще на початковому етапі: ми створюємо потоки вручну і намагаємося спланувати доступ до загальних ресурсів; з усім цим ще якось можна впоратися, якщо кількість ядер можна перерахувати на пальцях однієї руки. Але через кілька років, коли кількість процесорів обчислюватиметься сотнями, такої можливості вже не буде. З випуском CUDA nVidia зробила перший крок у вирішенні цієї проблеми – але, звичайно, це рішення підходить тільки для GPU від цієї компанії, та й то не для всіх. Тільки GF8 та 9 (і їх похідні Quadro/Tesla) сьогодні можуть працювати з програмами CUDA. І нова лінійка 260/280, звісно.

Натисніть на зображення для збільшення.

nVidia може хвалитися тим, що продала 70 мільйонів CUDA-сумісних GPU у всьому світі, але цього все одно мало, щоб стати стандартом де-факто. З урахуванням того, що конкуренти не сидять, склавши руки. AMD пропонує свій SDK (Stream Computing), та й Intel оголосила про рішення (Ct), хоча воно ще не доступне. Настає війна стандартів, і на ринку явно не буде місця для трьох конкурентів, поки інший гравець, наприклад Microsoft, не вийде з пропозицією загального API, що, звичайно, полегшить життя розробникам.

Тому nVidia має чимало труднощів на шляху затвердження CUDA. Хоча технологічно перед нами, безперечно, успішне рішення, ще залишається переконати розробників у його перспективах – і це буде нелегко. Втім, судячи з багатьох недавніх оголошень та новин з приводу API, майбутнє виглядає аж ніяк не сумним.

- Набір низькорівневих програмних інтерфейсів ( API) для створення ігор та інших високопродуктивних мультимедіа-додатків. Включає підтримку високопродуктивної 2D- І 3D-графіки, звуку та пристроїв введення.

Direct3D (D3D) - інтерфейс виведення тривимірних примітивів(Геометричних тіл). Входить в .

OpenGL(Від англ. Open Graphics Library, Дослівно - відкрита графічна бібліотека) - специфікація, що визначає незалежний від мови програмування крос-платформний програмний інтерфейс для написання додатків, що використовують двовимірну та тривимірну комп'ютерну графіку. Включає понад 250 функцій для малювання складних тривимірних сцен із простих примітивів. Використовується для створення відеоігор, віртуальної реальності, візуалізації в наукових дослідженнях. На платформі Windowsконкурує з .

OpenCL(Від англ. Open Computing Language, дослівно – відкрита мова обчислень) – фреймворк(Каркас програмної системи) для написання комп'ютерних програм, пов'язаних з паралельними обчисленнями на різних графічних ( GPU) та ( ). У фреймворк OpenCLвходять мову програмування та інтерфейс програмування додатків ( API). OpenCLзабезпечує паралелізм на рівні інструкцій та на рівні даних та є реалізацією техніки GPGPU.

GPGPU(скор. від англ. General-P urpose G raphics P rokussing U nits, дослівно – GPUзагального призначення) - техніка використання графічного процесора відеокарти для загальних обчислень, які зазвичай проводить .

Шейдер(Англ. shader) – програма побудови тіней на синтезованих зображеннях, використовують у тривимірної графіці визначення остаточних параметрів об'єкта чи зображення. Як правило, включає довільної складності опис поглинання та розсіювання світла, накладання текстури, відображення та заломлення, затінювання, зміщення поверхні та ефекти пост-обробки. Складні поверхні можуть бути візуалізовані за допомогою простих геометричних форм.

Рендеринг(Англ. rendering) – візуалізація, у комп'ютерній графіці процес отримання зображення за моделлю за допомогою програмного .

SDK(скор. від англ. Software Development Kit) - Набір інструментальних засобів розробки програмного забезпечення.

CPU(скор. від англ. Central Processing Unit, дослівно - центральний/основний/головний обчислювальний пристрій) - центральний (мікро); пристрій, що виконує машинні інструкції; частина апаратного забезпечення, що відповідає за виконання обчислювальних операцій (заданих операційною системою та прикладним програмним) і координує роботу всіх пристроїв.

GPU(скор. від англ. Graphic Processing Unit, дослівно - графічний обчислювальний пристрій) - графічний процесор; окремий пристрій або ігрової приставки, що виконує графічний рендеринг (візуалізацію). Сучасні графічні процесори дуже ефективно обробляють і реалістично відображають комп'ютерну графіку. Графічний процесор у сучасних відеоадаптерах застосовується як прискорювач тривимірної графіки, проте його можна використовувати в деяких випадках і для обчислень ( GPGPU).

Проблеми CPU

Довгий час підвищення продуктивності традиційних переважно відбувалося рахунок послідовного збільшення тактової частоти (близько 80% продуктивності визначала саме тактова частота) з одночасним збільшенням кількості транзисторів однією кристалі. Однак подальше підвищення тактової частоти (при тактовій частоті більше 3,8 ГГц чіпи просто перегріваються!) впирається в ряд фундаментальних фізичних бар'єрів (оскільки технологічний процес майже впритул наблизився до розмірів атома: , А розміри атома кремнію – приблизно 0,543 нм):

По-перше, зі зменшенням розмірів кристала і підвищенням тактової частоти зростає струм витоку транзисторів. Це веде до підвищення споживаної потужності та збільшення викиду тепла;

По-друге, переваги вищої тактової частоти частково зводяться нанівець через затримки при зверненні до пам'яті, так як час доступу до пам'яті не відповідає зростаючим тактовим частотам;

По-третє, для деяких додатків традиційні послідовні архітектури стають неефективними зі зростанням тактової частоти через так зване «фон-нейманівське вузьке місце» – обмеження продуктивності в результаті послідовного потоку обчислень. У цьому зростають резистивно-емкостные затримки передачі сигналів, що є додатковим вузьким місцем, що з підвищенням тактової частоти.

Розвиток GPU

Паралельно з йшло (і йде!) розвиток GPU:

…

Листопад 2008 р. – Intelпредставила лінійку 4-ядерних Intel Core i7, в основу яких покладено мікроархітектуру нового покоління Nehalem. Процесори працюють на тактовій частоті 26-32 ГГц. Виконані за 45-нм техпроцесом.

Грудень 2008 р. – розпочалися поставки 4-ядерного AMD Phenom II 940(кодова назва – Deneb). Працює на частоті 3 ГГц, випускається за техпроцесом 45-нм.

…

Травень 2009 р. – компанія AMDпредставила версію графічного процесора ATI Radeon HD 4890із тактовою частотою ядра, збільшеною з 850 МГц до 1 ГГц. Це перший графічнийпроцесор, працюючий на частоті 1 ГГц. Обчислювальна потужність чіпа завдяки збільшенню частоти зросла з 1,36 до 1,6 терафлоп. Процесор містить 800 (!) обчислювальних ядер, підтримує відеопам'ять GDDR5, DirectX 10.1, ATI CrossFireXта всі інші технології, властиві сучасним моделям відеокарт. Чіп виготовлений на базі 55-нм технології.

Основні відмінності GPU

Відмінними рисами GPU(порівняно з ) є:

- архітектура, максимально націлена на збільшення швидкості розрахунку текстур та складних графічних об'єктів;

- пікова потужність типового GPUнабагато вище, ніж у ;

– завдяки спеціалізованій конвеєрній архітектурі, GPUнабагато ефективніше у обробці графічної інформації, ніж .

«Криза жанру»

«Криза жанру» для назріло до 2005 р., – саме тоді з'явилися . Але, незважаючи на розвиток технології, зростання продуктивності звичайних помітно знизився. Водночас продуктивність GPUпродовжує зростати. Так, до 2003 р. і кристалізувалась ця революційна ідея – використовувати для потреб обчислювальну міць графічного. Графічні процесори стали активно використовуватися для «неграфічних» обчислень (симуляція фізики, обробка сигналів, обчислювальна математика/геометрія, операції з базами даних, обчислювальна біологія, обчислювальна економіка, комп'ютерний зір тощо).

Головна проблема в тому, що не було жодного стандартного інтерфейсу для програмування GPU. Розробники використовували OpenGLабо Direct3Dале це було дуже зручно. Корпорація NVIDIA(один з найбільших виробників графічних, медіа- та комунікаційних процесорів, а також бездротових медіа-процесорів; заснована у 1993 р.) зайнялася розробкою якогось єдиного та зручного стандарту, – і представила технологію CUDA.

Як це починалося

2006 р. – NVIDIAдемонструє CUDA™; початок революції у обчисленнях на GPU.

2007 р. – NVIDIAвипускає архітектуру CUDA(Початкова версія CUDA SDKбула представлена ​​15 лютого 2007 р.); номінація «Найкраща новинка» від журналу Popular Scienceта «Вибір читачів» від видання HPCWire.

2008 р. – технологія NVIDIA CUDAперемогла у номінації «Технічна перевага» від PC Magazine.

Що таке CUDA

CUDA(скор. від англ. Compute Unified Device Architecture, дослівно - уніфікована обчислювальна архітектура пристроїв) - архітектура (сукупність програмних та апаратних засобів), що дозволяє виробляти GPUобчислення загального призначення, у своїй GPUПрактично виступає у ролі потужного співпроцесора.

Технологія NVIDIA CUDA™– це єдине середовище розробки мовою програмування C, Що дозволяє розробникам створювати програмне вирішення складних обчислювальних завдань менший час, завдяки обчислювальної потужності графічних процесорів. У світі вже працюють мільйони GPUз підтримкою CUDA, та тисячі програмістів вже користуються (безкоштовно!) інструментами CUDAдля прискорення додатків та для вирішення найскладніших ресурсомістких завдань – від кодування відео- та аудіо- до пошуків нафти та газу, моделювання продуктів, виведення медичних зображень та наукових досліджень.

CUDAдає розробнику можливість на власний розсуд організовувати доступом до набору інструкцій графічного прискорювача і керувати його пам'яттю, організовувати у ньому складні паралельні обчислення. Графічний прискорювач із підтримкою CUDAстає потужною програмованою відкритою архітектурою, подібно до сьогоднішніх. Все це надає розпоряднику низькорівневий, розподільний і високошвидкісний доступ до обладнання, роблячи CUDAнеобхідною основою для побудови серйозних високорівневих інструментів, таких як компілятори, відладники, математичні бібліотеки, програмні платформи.

Уральський, провідний спеціаліст з технологій NVIDIA, порівнюючи GPUі , каже так: « - Це позашляховик. Він їздить завжди та скрізь, але не дуже швидко. А GPU- Це спорткар. На поганій дорозі він просто нікуди не поїде, але дайте хороше покриття – і він покаже всю свою швидкість, яка позашляховику і не снилася!..».

Можливості технології CUDA