CUDA з місця у кар’єр

NB: Стаття - короткий вступ, покрити всі нюанси програмування під CUDA в одній статті навряд чи можливо :-)

З чого почати?

NB: Якщо ви додасте параметр -keep, то після компіляції зможете знайти багато цікавих проміжних файлів.

Визначення функцій

Визначення даних

__constant__ - задає змінну в константній пам'яті. Слід звернути увагу, що значення констант потрібно завантажувати функцією cudaMemcpyToSymbol. Константи доступні з усіх тредів, швидкість роботи можна порівняти з регістрами (коли в кеш потрапляє).__shared__ - задає змінну в загальній пам'яті блоку тредів (тобто значення буде спільне на всіх). Тут потрібно підходити з обережністю - компілятор агресивно оптимізує доступ сюди (можна придушити модифікатором volatile), можна отримувати race condition, потрібно використовувати __syncthreads(); щоб дані гарантовано записалися. Shared memory розділена на банки, і коли 2 потоки одночасно намагаються звернутися до одного банку, виникає bank conflict і падає швидкість.

Всілокальні змінніякі ви визначили в ядрі (__device__) - у регістрах, найвища швидкість доступу.

Як потік дізнається над чим йому працювати

Основна ідея CUDA в тому, що для вирішення вашого завдання ви запускаєте тисячі та тисячі потоків, тому не варто лякатися того, що тут буде далі написано :-) Припустимо, треба зробити якусь операцію над картинкою 200x200. Картинка розбивається на шматки 10x10, і кожен піксел такого шматочка запускаємо по потоку. Виглядатиме це так: dim3 threads(10, 10);//розмір квардатика, 10*10 dim3 grid(20, 20);//скільки квадратиків потрібно щоб покрити все зображення

your_kernel >>(image,200,200); // Цей рядок запустить 40'000 потоків (не одночасно, одночасно працювати буде 200-2000 потоків приблизно).

На відміну від Brook+ від AMD, де ми відразу визначаємо якому потоку над якими даними працювати, в CUDA все не так: передавані kernel-у параметри однакові для всіх потоків, і потік повинен сам отримати дані для себе, щоб зробити це, потоку потрібно обчислити , де зображення він знаходиться. У цьому допомагають магічні змінні blockDim, blockIdx. const int ix = blockDim.x * blockIdx.x + threadIdx.x; const int iy = blockDim.y * blockIdx.y + threadIdx.y;

У ix та iy — координати, за допомогою яких можна отримати вихідні дані з масиву image, та записати результат роботи.

Оптимізація

Пару слів про те, як не зробити вашу програму дуже повільною (написати програму працюючу повільніше ніж CPU набагато простіше, ніж працюючу в 10 разів швидше :-))

  • Якнайменше використовуйте __global__ пам'ять.
  • При роботі з пам'яттю __shared__ уникайте конфліктів банків (втім багато завдань можуть бути вирішені без shared пам'яті).
  • Якнайменше розгалужень у коді, де різні потоки йдуть різними шляхами. Такий код не виконується паралельно.
  • Використовуйте якнайменше пам'яті. Чим менше пам'яті ви використовуєте, тим агресивніше компілятор і залізо зможуть запускати ваш kernel (наприклад, він може взяти 100 тредів, і використовуючи в 100 більше регістрів запустити одночасно на одному MP, радикально зменшуючи затримки)

Не виходить?

Насамперед слід прочитати документацію разом із SDK (NVIDIA_CUDA_Programming_Guide, CudaReferenceManual, ptx_isa), після цього можна запитати на офіційному форумі — там навіть девелопери nVidia часто відписуються, та й взагалі багаторозумних людей. Українською можна запитати у мене на форумі, наприклад, де відповім я :-) Також багато людей мешкає на gpgpu.ru.

Хардкорна конфа за С++. Ми запрошуємо лише профі.