Дослідження засобів організації паралельних обчислень на графічних процесорах

1. НТУУ «КПІ» ННК «ІПСА» Дослідження засобів організації паралельних обчислень на графічних процесорах Кобилінський Олександр Анатолійович Корначевський Ярослав Ілліч

2. GPGPU У наш час вже нічим не здивувати. У нас з'явилися телефони, які з легкістю вміщуються в долоні і керуються лише дотиками пальця до дисплея, абсолютно крихітні аудіо-плеєри, високошвидкісні канали зв'язку, за допомогою яких без затримок можна обмінюватися великими обсягами різної інформації, високопродуктивні процесори і тонкі ноутбуки все з тими ж продуктивними процесорами. Але є одна велика проблема - обробка інформації і даних у встановлені терміни. Буває, що навіть найсучасніших серверних залів з найсучаснішими процесорами недостатньо і економічно невигідно використовувати для рішень поставлених завдань. Тут вступає в гру, не так давно отримавши масове визнання, напрямок GPGPU - загальні обчислення на графічних процесорах. З кожним днем графічні процесори стають все більш продуктивними. То чому б не скористатися даною продуктивністю, і не розвантажити центральний процесор? Схоже таке запитання задали собі всі великі компанії і створили відповідні технології. У AMD технологія GPGPU має назву ATI Stream, а у NVIDIA - CUDA.

3. Сучасний графічний процесор Загальна схема архітектури G80. Тут: ВПВ – вибірка з потоку вершин (Vertex Thread Issue); ВПГ – вибірка з потоку геометрії (Geometry Thread Issue); ВПП – вибірка з потоку пікселів (Pixel Thread Issue); ПУП – процесор управління потоками (Thread Processor); SP – потоковий процесор (Streaming Processor); TF – блок фильтрації текстур (Texture Filtering Unit); L1 – кеш-память первого рівня; L2 – кеш-память другого рівня; FB – шина обміну даними (Front Bus)

4. Архітектура R600 Загальна схема архітектури R600:

5. Fermi – революція від NVIDIA

6. Fermi та CUDA Нововведення в архітектурі Fermi: • підтримує всі чотири методи округлення за стандартом IEEE 754-2008 (до найближчого ненульового значення, до нуля, до позитивної та негативної безкінечностей), що робить її повністю підходящою для наукових обчислень з 32-бітної точністю. • кардинально перероблений механізм обчислень з подвійною точністю, тепер кожен мультипроцесор може обробити до 16 операцій додавання-множення з подвійною точністю за такт. • введення підтримки корекції помилок ECC • нове покоління GigaThread Engine здатне запускати на виконання відразу кілька так званих ядер CUDA • зручність розробки ПЗ. По-перше, тепер використовується загальний адресний простір пам'яті для CUDA, що дозволяє NVIDIA говорити про повну підтримку об'єктно-орієнтованих мов програмування C і С++ • замість 16-кілобайтної колективної пам'яті в структурі SM тепер є 64 КБ, що розподіляються за необхідності на загальну пам'ять і традиційний L1-кеш

7. Відмінності між CPU та GPU

8. Сучасні терафлопи Динаміка зростання потужностей сучасних процесорів

9. CUDA CUDA - це скорочення від Compute Unified Device Architecture, тобто уніфікована архітектура комп'ютерних обчислень. Це комплекс апаратно-програмних засобів, який дозволяє запускати програмний код на графічному процесорі (GPU). Протягом останніх 10-15 років традиційні CPU нарощували продуктивність при виконанні послідовного коду шляхом ускладнення архітектури. Аналізуючи код, вони намагалися збільшити його ефективність на етапі виконання. Продовжувалося збільшення розмірів кеш-пам'яті, нарощування тактової частоти. Однак сьогодні можливості подальшого нарощування швидкості виконання послідовного коду практично підійшли до межі. Єдиним виходом є паралелізм. Тому GPU стає все більш і більш популярним як обчислювач. За умови правильного написання коду під GPU можна отримати велику продуктивність на квадратний міліметр кристалу, одиницю витраченої енергії, долар і т.п. GPU використовують інший підхід до нарощування продуктивності: графічний процесор використовує масиви, або «море» обчислювачів, які працюють паралельно, але, як правило, на більш низькій частоті. Така архітектура кардинально відрізняється, і CUDA - програмно-апаратна платформа, яка дозволяє розробникам використовувати потенціал цих масивів, при цьому застосовуючи стандартні засоби розробки.

10. CUDA потоки Рис. Умовне зображення архітектури GT200 Візьмемо реалізацію CUDA в графічному процесорі GT200. Він має 30 мультипроцесорів, один мультипроцесор працює на частоті приблизно 1,4 GHz, має 8 виконавчих пристроїв загального призначення, які виконують за такт типові операції, як то: додавання і множення дійсних чисел типу float, порівняння, умовні переходи, логічні операції, операції з регістрами.

11. Один мультипроцесор може одночасно виконувати 1024 програмних потоків(ниток). Але виконавчих пристроїв всього 8. Тому інструкції всіх 1024 потоків в один момент виконуватися не можуть. Нитки розбиті на групи по 32 штуки, так звані warp. В один момент виконується один warp, він виконується 4 такти. Але, з усіх 32 ниток, за один захід реально виконуються тільки ті, які виконують одну й ту ж інструкцію програми. Кожна нитка CUDA-додатку виконує одну і ту ж саму програму, але алгоритму доступний номер нитки, і тому алгоритм може довільно змінюватися залежно від її номера. Можна запустити хоч 10000 різних підпрограм, для кожної нитки вибравши свою. Але це буде неефективно із-за вищенаведеної особливості виконання warp. Коли всі нитки усередині кожного warp мають однаковий шлях виконання, досягається максимальна продуктивність. CUDA потоки

12. Організація потоків CUDA

13. Типи пам`яті в CUDA

14. Компіляція програм на CUDA Для того, щоб використовувати CUDA на комп'ютері, необхідно наступне: 1. CUDA-сумісний GPU 2. GPU драйвер, що містить в собі CUDA драйвер 3. СUDA software Для написання CUDA-додатків можно використовувати будь-який текстовий редактор. Зазвичай, для написання CUDA-програм під операційну систему Microsoft Windows використовують інструментальне середовище Microsoft Visual Studio. Варто відзначити появу вкрай корисного інструментарію NVIDIA Nexus - надбудови для Microsoft Visual Studio, що дозволяє писати код, відслідковувати помилки, компілювати і тестувати ПЗ в цьому середовищі розробки з використанням звичних інструментів.

15. CUDA позначення Позначення функцій, перед функціями в .сu файлі можуть стояти наступні «модифікатори»: • __device__ • __global__ • __host__ Позначення даних: • __device__ - означає що змінна знаходиться в глобальній пам'яті відеокарти (тобто якої там 512Мб, тобто). Дуже повільна пам'ять за мірками обчислень на відеокарті (хоча і швидше пам'яті центрального процесора в кілька разів), рекомендується використовувати як можна рідше. cudaMemcpy (device_variable, host_variable, size, cudaMemcpyHostToDevice); cudaMemcpyDeviceToHost – читання або запис даних у зворотній бік зворотний бік • __constant__ - задає змінну в константній пам'яті. • __shared__ - задає змінну в колективній пам'яті блоку потоків (тобто і значення буде спільне на всіх). Потрібно використовувати __syncthreads (); щоб дані гарантовано записалися.

16. Використання CUDA Зазвичай використання CUDA зводиться до наступних етапів: • виділяємо пам`ять на GPU • копіюємо дані з пам`яті CPU у виділену пам`ять GPU • здійснюємо запуск ядра(або послідовно запускаємо декілька ядер) • копіюємо результати обчислень назад в пам`ять CPU • звільнюємо виділену пам`ять GPU

17. Множення матриць // Parallel matrix multiplication __global__ void matrixMultiplicationParallel(float *matrixA, float *matrixB, float *matrixC, int numColsA, int numColsB) { // Block index int bx = blockIdx.x; int by = blockIdx.y; // Thread index int tx = threadIdx.x; int ty = threadIdx.y; int row = bx * blockDim.x + tx; int col = by * blockDim.y + ty; unsigned int k = 0; unsigned int idA = 0; unsigned int idB = 0; matrixC[row * numColsB + col] = 0; for(k = 0; k < numColsA; k++) { idA = row * numColsA + k; idB = k * numColsB + col; matrixC[row * numColsB + col] += matrixA[idA] * matrixB[idB]; } }

18. ATI Stream Технологія ATI Stream являє собою набір апаратних і програмних технологій, які дозволяють графічним процесорам AMD (GPU), що працюють у взаємодії з центральними процесорами системи (CPU), для прискорення додатків за межами традиційної графіки і обробки відео. Це дозволяє виконувати складні обчислювальні завдання більш ефективно. Для того щоб використовувати ATIStream потрібно лише встановити ATI Stream SDK. ATI Steam SDK 2.1 забезпечує програмування з використанням OpenCL для GPU і x86-сумісних CPU AMD. Інакше кажучи, ATI Stream SDK допомагає розробникам підвищити продуктивність додатків за рахунок використання об'єднаних ресурсів CPU і GPU, доступних в системі.

19. OpenCL OpenCL - фреймворк для створення комп'ютерних програм, пов'язаних з паралельними обчисленнями на різних графічних і центральних процесорах. У фреймворк OpenCL входять мова програмування, яка базується на стандарті C99, та інтерфейс програмування комп'ютерних програм (API). OpenCL - повністю відкритий стандарт, його використання доступне на базі вільних ліцензій. Мета OpenCL полягає в тому, щоб доповнити OpenGL і OpenAL, які є відкритими галузевими стандартами для тривимірної комп'ютерної графіки і звуку, користуючись можливостями GPU.

20. Моделі OpenCL Платформа OpenCL складається з хоста поєднаного з пристроями, що підтримують OpenCL. Модель платформи OpenCL Виконання OpenCL-програми складається з двох частин: хостової частини програми і kernel(ядра) Модель виконання OpenCL

21. Моделі OpenCL Модель пам`яті OpenCL Програмна модель OpenCL

22. Використання OpenCL Позначення функцій OpenCL: • __kernel декларує функцію як ядро OpenCL Позначення даних: • __global • __constant • __local • __private // OpenCL Kernel Function for matrix addition __kernel void addMatrix(__global float *matrixA, __global float *matrixB, __global float *matrixC, int HC, int WC) { unsigned int row = get_group_id(0) * get_local_size(0) + get_local_id(0); unsigned int col = get_group_id(1) * get_local_size(1) + get_local_id(1); unsigned int index = row *WC + col; if ( row < HC && col < WC ) { matrixC[index] = matrixA[index] + matrixB[index]; } }

23. Результати досліджень

24. Результати досліджень Транспонування матриці

25. Результати досліджень Порівняння OpenCL та CUDA

26. Результати досліджень Послідовний пошук

27. Висновки З кожним днем графічні процесори стають все більш продуктивними. Сьогодні їх використовують не лише для розрахунку графічних задач, але й для більш широкого колу задач. На графічних процесорах добре вирішувати задачі з високою ступінню паралелізму. Але потрібно звичайно ж зважати і на аспекти роботи з GPU. Провідні гравці на ринку графічних процесорів швидко усвідомили перспективність напрямку GPGPU та випустили власні технології для розпаралелювання задач на своїх графічних процесорах: у NVIDIA це CUDA, у AMD – це ATI Stream. Але ж не переписувати кожного разу додаток з однієї платформи на іншу? Схоже це запитання задала собі компанія Apple та ініціювала розробку такого стандарту як OpenCL. Консорціум Khronos Group, в який входить багато компаній,розробляє стандарт OpenCL. Стандарт OpenCL на даний момент підтримують і NVIDIA і AMD, а також Apple у своїй операційній системі Mac OS X. Вже сьогодні технологіїGPGPU знаходять своє застосування у багатьох галузях науки. Окрім цього, щодня удосконалюються архітектури графічних процесорів та програмні засоби для розробки додатків під ці архітектури. Технологія CUDA дозволвяє у сотні разів збільшити швидкодію додатків. Швидше за все, в майбутньому CPU та GPU зіллються в один пристрій.