Технологія nvidia cuda. NVidia CUDA: обчислення на відеокарті чи смерть CPU? Мінімізація переміщень даних Host Device

Та інші. Однак, пошук комбінації «CUDA scan» видав лише 2 статті ніяк не пов'язані з, власне, алгоритмом scan на GPU – а це один із самих базових алгоритмів. Тому, надихнувшись щойно переглянутим курсом на Udacity - Intro to Parallel Programming, я і наважився написати повнішу серію статей про CUDA. Відразу скажу, що серія буде ґрунтуватися саме на цьому курсі, і якщо у вас є час – набагато корисніше буде пройти його. На даний момент плануються такі статті:
Частина 1: Запровадження.
Частина 2: Апаратне забезпечення GPU та шаблони паралельної комунікації.
Частина 3: Фундаментальні алгоритми GPU: згортка (reduce), сканування (scan) та гістограма (histogram).
Частина 4: Фундаментальні алгоритми GPU: ущільнення (compact), сегментоване сканування (segmented scan), сортування. Практичне застосування деяких алгоритмів.
Частина 5: Оптимізація програм GPU.
Частина 6: Приклад паралелізації послідовних алгоритмів.
Частина 7: Додаткові теми паралельного програмування, динамічний паралелізм.

Затримка vs пропускна здатність

Перше питання, яке має поставити кожен перед застосуванням GPU для вирішення своїх завдань - а для яких цілей хороший GPU, коли варто його застосовувати? Для відповіді слід визначити 2 поняття:
Затримка(latency) - час, що витрачається виконання однієї інструкції/операції.
Пропускна спроможність- кількість інструкцій/операцій, які виконуються за одиницю часу.
Простий приклад: маємо легковий автомобіль зі швидкістю 90 км/год та місткістю 4 особи, і автобус зі швидкістю 60 км/год та місткістю 20 осіб. Якщо за операцію прийняти переміщення 1 особи на 1 кілометр, то затримка легкового автомобіля – 3600/90=40с – за стільки секунд 1 людина подолає відстань у 1 кілометр, пропускна спроможність автомобіля – 4/40=0.1 операцій/секунду; затримка автобуса - 3600/60=60с, пропускна спроможність автобуса - 20/60=0.3(3) операцій/секунду.
Так от, CPU - це автомобіль, GPU - автобус: він має велику затримку, але також і велику пропускну здатність. Якщо для вашого завдання затримка кожної конкретної операції не така важлива як кількість цих операцій за секунду - варто розглянути застосування GPU.

Базові поняття та терміни CUDA

• Пристрій (device)- GPU. Виконує роль «підлеглого» – робить лише те, що йому каже CPU.
• Хост (host)- CPU. Виконує керуючу роль - запускає завдання на пристрої, виділяє пам'ять на пристрої, переміщує пам'ять/з пристрою. І так, використання CUDA передбачає, що як пристрій, так і хост мають свою окрему пам'ять.
• Ядро (Kernel)- Завдання, що запускається хостом на пристрої.

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

1. Хост виділяє необхідну кількість пам'яті на пристрої.
2. Хост копіює дані зі своєї пам'яті в пам'ять пристрою.
3. Хост стартує виконання певних ядер на пристрої.
4. Пристрій виконує ядра.
5. Хост копіює результати з пам'яті пристрою на свою пам'ять.

Ядра

1. Усередині ядер ви маєте можливість дізнатися «ідентифікатор» або, простіше кажучи, позицію потоку, який зараз виконується - використовуючи цю позицію ми добиваємося того, що одне ядро ​​працюватиме з різними даними залежно від потоку, в якому воно запущено. До речі, така організація паралельних обчислень називається SIMD (Single Instruction Multiple Data) - коли кілька процесорів виконують одночасно одну й ту саму операцію, але на різних даних.
2. У деяких випадках у коді ядра необхідно використовувати різні способи синхронізації.

• Спочатку задаються розміри так званої сітки (grid), в 3D координатах: grid_x, grid_y, grid_z. В результаті, сітка складатиметься з grid_x*grid_y*grid_zблоків.
• Потім задаються розміри блоку в 3D координатах: block_x, block_y, block_z. В результаті блок буде складатися з block_x*block_y*block_zпотоків. Отже, маємо grid_x*grid_y*grid_z*block_x*block_y*block_zпотоків. Важливе зауваження - максимальна кількість потоків в одному блоці обмежена і залежить від моделі GPU - типові значення 512 (старіші моделі) і 1024 (новіші моделі).
• Усередині ядра доступні змінні threadIdxі blockIdxз полями x, y, z- вони містять 3D координати потоку в блоці та блоку в сітці відповідно. Також доступні змінні blockDimі gridDimз тими ж полями – розміри блоку та сітки відповідно.

Пишемо першу програму на CUDA

main.cpp

#include #include #include #include #include #include #include #include #include "openMP.hpp" #include "CUDA_wrappers.hpp" #include "common/image_helpers.hpp" using namespace cv; using namespace std; int main(int argc, char** argv) ( using namespace std::chrono; if(argc != 2) ( cout<<" Usage: convert_to_grayscale imagefile" << endl; return -1; } Mat image, imageGray; uchar4 *imageArray; unsigned char *imageGrayArray; prepareImagePointers(argv, image, &imageArray, imageGray, &imageGrayArray, CV_8UC1); int numRows = image.rows, numCols = image.cols; auto start = system_clock::now(); RGBtoGrayscaleOpenMP(imageArray, imageGrayArray, numRows, numCols); auto duration = duration_cast(system_clock::now() - start); cout<<"OpenMP time (ms):" << duration.count() << endl; memset(imageGrayArray, 0, sizeof(unsigned char)*numRows*numCols); RGBtoGrayscaleCUDA(imageArray, imageGrayArray, numRows, numCols); return 0; }

prepareImagePointers

template void prepareImagePointers(const char * const inputImageFileName, cv::Mat& inputImage, T1** inputImageArray, cv::Mat& outputImage, T2** outputImageArray, const int outputImageType) ( using namespace std; using namespace cv; inputImage = imread IMREAD_COLOR), if (inputImage.empty()) ( cerr<< "Couldn"t open input file." << endl; exit(1); } //allocate memory for the output outputImage.create(inputImage.rows, inputImage.cols, outputImageType); cvtColor(inputImage, inputImage, cv::COLOR_BGR2BGRA); *inputImageArray = (T1*)inputImage.ptr(0); *outputImageArray = (T2*)outputImage.ptr (0); }

openMP.hpp

#include #include #include void RGBtoGrayscaleOpenMP(uchar4 *imageArray, unsigned char *imageGrayArray, int numRows, int numCols) ( #pragma omp parallel for collapse(2) for (int i = 0; i< numRows; ++i) { for (int j = 0; j < numCols; ++j) { const uchar4 pixel = imageArray; imageGrayArray = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z; } } }

Прихований текст

void RGBtoGrayscaleCUDA(const uchar4 * const h_imageRGBA, unsigned char* const h_imageGray, size_t numRows, size_t numCols) ( uchar4 *d_imageRGBA; unsigned char *d_imageGray; const size_t num //облікувальна пам'ять на пристрої для запуску і запуску checkCudaErrors(cudaMalloc(&d_imageRGBA, sizeof(uchar4) * numPixels)); the GPU checkCudaErrors(cudaMemcpy(d_imageRGBA, h_imageRGBA, sizeof(uchar4) * numPixels, cudaMemcpyHostToDevice));

Прихований текст

#include #define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__) template void check (T err, const char * const func, const char * const file, const int line) ( if (err! = cudaSuccess) ( std::cerr<< "CUDA error at: " << file << ":" << line << std::endl; std::cerr << cudaGetErrorString(err) << " " << func << std::endl; exit(1); } }

Прихований текст

dim3 blockSize; dim3 gridSize; int threadNum; cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); threadNum = 1024; blockSize = dim3(threadNum, 1, 1); gridSize = dim3(numCols/threadNum+1, numRows, 1); cudaEventRecord(start); rgba_to_grayscale_simple<<>>(d_imageRGBA, d_imageGray, numRows, numCols); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); float milliseconds = 0; cudaEventElapsedTime(&milliseconds, start, stop); std::cout<< "CUDA time simple (ms): " << milliseconds << std::endl;

rgba_to_grayscale_simple

Global__ void rgba_to_grayscale_simple(const uchar4* const d_imageRGBA, unsigned char* const d_imageGray, int numRows, int numCols) ( int y = blockDim.y*blockIdx.y + threadIdx.x; .x, if (x>=numCols || y>=numRows) return, const int offset = y*numCols+x; const uchar4 pixel = d_imageRGBA; 0.114f*pixel.z;

Прихований текст

checkCudaErrors(cudaMemcpy(h_imageGray, d_imageGray, sizeof(unsigned char) * numPixels, cudaMemcpyDeviceToHost)); cudaFree(d_imageGray); cudaFree(d_imageRGBA);

Прихований текст

Процесор: Intel Core (TM) i7-3615QM CPU @ 2.30GHz.
GPU: NVIDIA GeForce GT 650M, 1024 МБ, 900 МГц.
RAM: DD3, 2x4GB, 1600 МГц.
OS: OS X 10.9.5.
Компілятор: g ++ (GCC) 4.9.2 20141029.
CUDA компілятор: Cuda compilation tools, release 6.0, V6.0.1.
Версія OpenMP, що підтримується: OpenMP 4.0.

rgba_to_grayscale_optimized

#define WARP_SIZE 32 __global__ void rgba_to_grayscale_optimized(const uchar4* const d_imageRGBA, unsigned char* const d_imageGray, int numRows, int numCols, int elemsPerThread) ( int yx; x*blockIdx.x + threadIdx.x; const int loop_start = (x/WARP_SIZE * WARP_SIZE)*(elemsPerThread-1)+x;for (int i=loop_start, j=0; j

Прихований текст

threadNum=128; const int elemsPerThread = 16; blockSize = dim3(threadNum, 1, 1); gridSize = dim3(numCols/(threadNum*elemsPerThread) + 1, numRows, 1); cudaEventRecord(start); rgba_to_grayscale_optimized<<>>(d_imageRGBA, d_imageGray, numRows, numCols, elemsPerThread); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); milliseconds = 0; cudaEventElapsedTime(&milliseconds, start, stop); std::cout<< "CUDA time optimized (ms): " << milliseconds << std::endl;

Теги: Додати теги

- Набір низькорівневих програмних інтерфейсів ( 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

Технологія CUDA

Володимир Фролов,[email protected]

Анотація

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

1. Введення

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

Тому всі виробники сучасних обчислювальних систем йдуть у бік збільшення числа процесорів і ядер, а не збільшують частоту одного процесора. Кількість ядер центрального процесора (CPU) у передових системах зараз дорівнює 8.

Інша причина-відносно невисока швидкість роботи оперативної пам'яті. Як би швидко не працював процесор, вузькими місцями, як показує практика, є зовсім не арифметичні операції, а саме невдалі звернення до пам'яті-кеш-промахи.

Проте якщо подивитися у бік графічних процесорів GPU (Graphics Processing Unit), там по шляху паралелізму пішли набагато раніше. У сьогоднішніх відеокартах, наприклад GF8800GTX, число процесорів може досягати 128. Продуктивність подібних систем при вмілому їх програмуванні може бути дуже значною (рис. 1).

Рис. 1. Кількість операцій з плаваючою точкою для CPU та GPU

Коли перші відеокарти тільки з'явилися у продажу, вони були досить прості (порівняно з центральним процесором) вузькоспеціалізовані пристрої, призначені для того, щоб зняти з процесора навантаження по візуалізації двомірних даних. З розвитком ігрової індустрії та появою таких тривимірних ігор як Doom (рис. 2) та Wolfenstein 3D (рис. 3) виникла потреба у 3D візуалізації.

Малюнки 2.3. Ігри Doom та Wolfenstein 3D

З часу створення компанією 3Dfx перших відеокарт Voodoo, (1996) і аж до 2001 року в GPU був реалізований тільки фіксований набір операцій над вхідними даними.

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

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

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

Сьогодні з'явилася тенденція нетрадиційного використання відеокарт для вирішення завдань у галузях квантової механіки, штучного інтелекту, фізичних розрахунків, криптографії, фізично коректної візуалізації, реконструкції з фотографій, розпізнавання тощо. Ці завдання незручно вирішувати у межах графічних API (DirectX, OpenGL), оскільки ці API створювалися зовсім інших застосувань.

Розвиток програмування загального призначення на GPU (General Programming on GPU, GPGPU) логічно призвело до виникнення технологій, націлених на ширше коло завдань, ніж розтеризація. В результаті компанією Nvidia була створена технологія Compute Unified Device Architecture (або скорочено CUDA), а компанією ATI, що конкурує, - технологія STREAM.

Слід зазначити, що на момент написання цієї статті технологія STREAM сильно відставала у розвитку від CUDA, і тому тут вона не розглядатиметься. Ми зосередимося на CUDA – технології GPGPU, що дозволяє писати програми на підмножині мови C++.

2. Принципова різниця між CPU та GPU

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

2.1. Можливості

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

Для GPU це негаразд. Як ви дізнаєтеся, прочитавши цю статтю, CUDA має цілих 6 видів пам'яті. Читати можна з будь-якого осередку, доступного фізично, але записувати – не в усі осередки. Причина полягає в тому, що GPU в будь-якому випадку є специфічним пристроєм, призначеним для конкретних цілей. Це обмеження запроваджено задля збільшення швидкості роботи певних алгоритмів та зниження вартості обладнання.

2.2. Швидкодія пам'яті

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

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

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

На GPU повільні звернення до пам'яті приховують за допомогою паралельних обчислень. Поки одні завдання чекають на дані, працюють інші, готові до обчислень. Це один із основних принципів CUDA, що дозволяють сильно підняти продуктивність системи в цілому.

3. Ядро CUDA

3.1. Поточна модель

Обчислювальна архітектура CUDA базується на концепціїодна команда на безліч даних(Single Instruction Multiple Data, SIMD) та понятті мультипроцесора.

Концепція SIMD має на увазі, що одна інструкція дозволяє одночасно обробити безліч даних. Наприклад, команда addps в процесорі Pentium 3 і новіших моделях Pentium дозволяє складати одночасно 4 числа з плаваючою точкою одинарної точності.

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

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

Під пристроєм (device) у нашій статті ми розумітимемо відеоадаптер, що підтримує драйвер CUDA, або інший спеціалізований пристрій, призначений для виконання програм, що використовують CUDA (таке, наприклад, як NVIDIA Tesla). У нашій статті ми розглянемо GPU лише як логічне обладнання, уникаючи конкретних деталей реалізації.

Хостом (host ) ми називатимемо програму у звичайній оперативної пам'яті комп'ютера, використовує CPU і виконує керуючі функції роботи з пристроєм.

Фактично, та частина вашої програми, яка працює на CPU – цехост, а ваша відеокарта -пристрій. Логічно пристрій можна подати як набір мультипроцесорів (мал. 4) плюс драйвер CUDA.

Рис. 4. Пристрій

Припустимо, що ми хочемо запустити на нашому пристрої деяку процедуру в N потоках (тобто хочемо розпаралелити її роботу). Відповідно до документації CUDA, назвемо цю процедуру ядром.

Особливістю архітектури CUDA є блочно-сіточна організація, незвичайна багатопоточних додатків (рис. 5). Драйвер CUDA самостійно розподіляє ресурси пристрою між потоками.

Рис. 5. Організація потоків

На рис. 5. ядро ​​позначене як Kernel. Всі потоки, що виконують це ядро, об'єднуються в блоки (Block), а блоки, своєю чергою, об'єднуються в сітку (Grid).

Як видно з рис 5, для ідентифікації потоків використовуються двомірні індекси. Розробники CUDA надали можливість працювати з тривимірними, двомірними або простими (одномірними) індексами, залежно від того, як зручніше програмісту.

У загальному випадку індекси є тривимірними векторами. Для кожного потоку будуть відомі: індекс потоку всередині блоку threadIdx та індекс блоку всередині сітки blockIdx. При запуску всі потоки відрізнятимуться лише цими індексами. Фактично саме через ці індекси програміст здійснює управління, визначаючи, яка саме частина його даних обробляється в кожному потоці.

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

Блок завдань (потоків) виконується на мультипроцесорі частинами або пулами, званими warp. Розмір warp на поточний момент у відеокартах із підтримкою CUDA дорівнює 32 потокам. Завдання всередині пулу warp виконуються у SIMD стилі, тобто. у всіх потоках всередині warp одночасно може виконуватись лише одна інструкція.

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

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

3.1.1. Розгалуження

Відразу ж виникає питання: якщо в той самий момент часу всі потоки всередині warp виконують ту саму інструкцію, то як бути з розгалуженнями? Адже якщо код програми зустрічається розгалуження, то інструкції будуть вже різні. Тут застосовується стандартне SIMD програмування рішення (рис 6).

Рис. 6. Організація розгалуження в SIMD

Нехай є наступний код:

if(cond)B;

У випадку SISD (Single Instruction Single Data) ми виконуємо оператор A, перевіряємо умову, потім виконуємо оператори B і D (якщо умова є істинною).

Нехай тепер ми маємо 10 потоків, що виконуються в стилі SIMD. У всіх 10 потоках ми виконуємо оператор A, потім перевіряємо умову cond і виявляється, що в 9 із 10 потоках воно істинне, а в одному потоці - хибно.

Зрозуміло, що ми не можемо запустити 9 потоків для виконання оператора B, а один - для виконання оператора C, тому що одночасно у всіх потоках може виконуватися тільки одна інструкція. У цьому випадку потрібно вчинити так: спочатку «вбиваємо» потік, що відколовся, так, щоб він не псував нічиї дані, і виконуємо 9 потоків, що залишилися. Потім "вбиваємо" 9 потоків, що виконали оператор B, і проходимо один потік з оператором C. Після цього потоки знову об'єднуються і виконують оператор D все одночасно.

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

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

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

3.1.2. Взаємодія між потоками

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

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

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

3.2. Пам'ять

У CUDA виділяють шість видів пам'яті (рис. 7). Це регістри, локальна, глобальна, розділяється, константна та текстурна пам'ять.

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

Рис. 7. Види пам'яті у CUDA

3.2.0. Реєстри

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

При звичайному поділі в 64 потоки на блок виходить всього 128 регістрів (існують деякі об'єктивні критерії, але 64 підходить в середньому для багатьох завдань). Реально, 128 регістрів nvcc ніколи не виділить. Зазвичай він не дає більше 40, а решта змінних попаде в локальну пам'ять. Так відбувається тому, що на одному мультипроцесорі може виконуватися кілька блоків. Компілятор намагається максимізувати кількість одночасно працюючих блоків. Для більшої більшої ефективності треба намагатися займати менше ніж 32 регістри. Тоді теоретично може бути запущено 4 блоки (8 warp-ів, якщо 64 треди в одному блоці) на одному мультипроцесорі. Однак тут ще слід враховувати обсяг розділеної пам'яті, що займається потоками, так як якщо один блок займає всю пам'ять, що розділяється, два таких блоки не можуть виконуватися на мультипроцесорі одночасно .

3.2.1. Локальна пам'ять

У випадках, коли локальні дані процедур займають занадто великий розмір, або компілятор не може обчислити їм певний постійний крок при зверненні, він може помістити їх у локальну пам'ять. Цьому може сприяти, наприклад, приведення покажчиків типів різних розмірів.

Фізично локальна пам'ять є аналогом глобальної пам'яті і працює з тією ж швидкістю. На момент написання статті не було жодних механізмів, що дозволяють явно заборонити компілятор використання локальної пам'яті для конкретних змінних. Оскільки проконтролювати локальну пам'ять досить важко, краще не використовувати її зовсім (див. розділ 4 «Рекомендації оптимізації»).

3.2.2. Глобальна пам'ять

У документації CUDA як одне з основних досягненьТехнологія наводить можливість довільної адресації глобальної пам'яті. Тобто можна читати з будь-якого осередку пам'яті, і писати можна також у довільний осередок (на GPU це зазвичай не так).

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

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

Змінні, оголошені з кваліфікатором __global__, розміщуються у світовій пам'яті. Глобальну пам'ять можна також виділити динамічно, викликавши функцію cudaMalloc(void* mem, int size) на хості. З пристрою цю функцію не можна викликати. Звідси випливає, що розподілом пам'яті має займатися програма-хост, що працює на CPU. Дані з хоста можна надсилати у пристрій викликом функції cudaMemcpy:

cudaMemcpy(void* gpu_mem, void* cpu_mem, int size, cudaMemcpyHostToDevice);

Так само можна зробити і зворотну процедуру:

cudaMemcpy(void* cpu_mem, void* gpu_mem, int size, cudaMemcpyDeviceToHost);

Цей виклик також здійснюється з хоста.

При роботі з глобальною пам'яттю важливо дотримуватися правила коалесингу (coalescing). Основна ідея в тому, що треди повинні звертатися до послідовних осередків пам'яті, причому 4,8 або 16 байт. При цьому найперший тред повинен звертатися за адресою, вирівняною на кордон відповідно 4,8 або 16 байт. Адреси, що повертаються cudaMalloc, вирівняні як мінімум за кордоном 256 байт.

3.2.3. Пам'ять, що розділяється

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

Відмінною рисою пам'яті, що розділяється, є те, що вона адресується однаково для всіх завдань усередині блоку (рис. 7). Звідси випливає, що її можна використовувати для обміну даними між потоками лише одного блоку.

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

Змінні, оголошені з кваліфікатором __shared__, розміщуються в пам'яті, що розділяється.

Shared__ float mem_shared;

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

float x = mem_shared;

Де threadIdx.x – індекс x потоку всередині блоку.

3.2.4. Константна пам'ять

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

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

#define N 100

Constant__ int gpu_buffer[N];

void host_function()

int cpu_buffer[N];

cudaMemcpyToSymbol(gpu_buffer, cpu_buffer, sizeof(int )*N);

// __global__ означає, що device_kernel – ядро, яке може бути запущене на GPU

Global__ void device_kernel()

int a = gpu_buffer;

int b = gpu_buffer + gpu_buffer;

// gpu_buffer = a; ПОМИЛКА! константна пам'ять доступна лише для читання

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

3.2.5. Текстурна пам'ять

Текстурна пам'ять кешується (рис. 4). Для кожного мультипроцесора є лише один кеш, отже, цей кеш загальний всім завдань усередині блоку.

Назва текстурної пам'яті (і, на жаль, функціональність) успадковано від понять «текстура» та «текстурування». Текстурування - це процес накладання текстури (просто картинки) на полігон у процесі растеризації. Текстурна пам'ять оптимізована під вибірку 2D даних і має такі можливості:

швидка вибірка значень фіксованого розміру (байт, слово, подвійне або вчетверне слово) з одномірного або двомірного масиву;

нормалізована адресація числами типу float в інтервалі. Потім їх можна вибирати, використовуючи нормалізовану адресацію. Результуючим значенням буде слово типу float4, відображене інтервал ;

CudaMalloc((void**) &gpu_memory, N*sizeof (uint4 )); //виділимо пам'ять у GPU

// Налаштування параємтрів текстури texture

Texture.addressMode = cudaAddressModeWrap; //режим Wrap

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; //найближче значення

Texture.normalized = false; // не використовувати нормалізовану адресацію

CudaBindTexture (0, texture , gpu _ memory , N ) // відтепер ця пам'ять вважатиметься текстурною


CudaMemcpy (gpu _ memory , cpu _ buffer , N * sizeof (uint 4), cudaMemcpyHostToDevice ); // копіюємо дані наGPU

// __global__ означає, що device_kernel - ядро, яке потрібно розпаралелити

Global__ void device_kernel()

uint4 a = tex1Dfetch(texture,0); // можна вибирати дані лише в такий спосіб!

uint4 b = tex1Dfetch(texture,1);

int c = a.x*b.y;

...

3.3. Простий приклад

Як простий приклад пропонується розглянути програму cppIntegration з CUDA SDK. Вона демонструє прийоми роботи з CUDA, а також використання nvcc (спеціальний компілятор підмножини С++ від Nvidia) у поєднанні з MS Visual Studio, що спрощує розробку програм на CUDA.

4.1. Правильно проводьте розбиття вашого завдання

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

4.2. Вибір типу пам'яті

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

Використовуйте глобальну пам'ять у поєднанні з пам'яттю, що розділяється, якщо всі завдання звертаються безсистемно до різних, далеко розташованих один від одного ділянок пам'яті (з дуже різними адресами або координатами, якщо це 2D/3D дані).

глобальна пам'ять => пам'ять, що розділяється

Syncthreads();

Обробити дані в пам'яті

Syncthreads();

глобальна пам'ять<= разделяемая память

4.3. Увімкніть лічильники пам'яті

Прапор компілятора --ptxas-options=-v дозволяє точно сказати, скільки і якої пам'яті (регістрів, що розділяється, локальної, константної) ви використовуєте. Якщо компілятор використовує локальну пам'ять, ви знаєте про це. Аналіз даних про кількість і типи пам'яті, що використовується, може сильно допомогти вам при оптимізації програми.

4.4. Намагайтеся мінімізувати використання регістрів і пам'яті, що розділяється

Чим більше ядро ​​використовує регістрів або пам'яті, що розділяється, тим менше потоків (вірніше warp-ів) одночасно можуть виконуватися на мультипроцесорі, т.к. ресурси мультипроцесора обмежені. Тому невелике збільшення зайнятості регістрів або пам'яті, що розділяється, може призводити в деяких випадках до падіння продуктивності в два рази - саме через те, що тепер рівно в два рази менше warp-ів одночасно виконуються на мультипроцесорі.

4.5. Пам'ять, що розділяється, замість локальної.

Якщо компілятор Nvidia з якоїсь причини розташував дані в локальній пам'яті (зазвичай це помітно з дуже сильного падіння продуктивності в місцях, де нічого ресурсомісткого немає), з'ясуйте, які саме дані потрапили в локальну пам'ять, і помістіть їх у пам'ять, що розділяється (shared memory ).

Найчастіше компілятор має змінну в локальній пам'яті, якщо вона використовується не часто. Наприклад, це акумулятор, де ви накопичуєте значення, розраховуючи щось у циклі. Якщо цикл великий за обсягом коду (але не за часом виконання!), компілятор може помістити ваш акумулятор в локальну пам'ять, т.к. він використовується відносно рідко, а регістрів мало. Втрата продуктивності у разі може бути помітною.

Якщо ж ви дійсно рідко використовуєте змінну - краще помістити її в глобальну пам'ять.

Хоча автоматичне розміщення компілятором таких змінних у локальній пам'яті може бути зручним, насправді це не так. Непросто знайти вузьке місце при наступних модифікаціях програми, якщо змінна почне використовуватися частіше. Компілятор може перенести таку змінну на регістрову пам'ять, а може й не перенести. Якщо модифікатор __global__ буде вказано явно, програміст швидше зверне на це увагу.

4.6. Розгортання циклів

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

Ось як можна розгорнути цикл знаходження суми масиву (наприклад, цілісного):

int a[N]; int summ;

for (int i=0;i

Зрозуміло, цикли можна розгорнути і вручну (як показано вище), але це малопродуктивна праця. Набагато краще використовувати шаблони С++ у поєднання з функціями, що вбудовуються.

template

class ArraySumm

Device__ static T exec (const T * arr) ( return arr + ArraySumm (arr+1); )

template

class ArraySumm<0,T>

Device__ static T exec(const T* arr) ( return 0; )

for (int i=0;i

summ+= ArraySumm<4,int>::exec(a);

Слід зазначити одну цікаву особливість компілятора nvcc. Компілятор завжди вбудовуватиме функції типу __device__ за замовчуванням (щоб це скасувати, існує спеціальна директива __noinline__) .

Отже, можна бути впевненим у тому, що приклад, подібний до наведеного вище, розгорнеться у просту послідовність операторів, і ні в чому не поступатиметься за ефективністю коду, написаному вручну. Однак у загальному випадку (не nvcc) у цьому впевненим бути не можна, тому що inline є лише вказівкою компілятора, яку він може проігнорувати. Тому не гарантується, що ваші функції вбудовуватимуться.

4.7. Вирівнювання даних та вибірка по 16 байт

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

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

4.8. Конфлікти банків пам'яті, що розділяється

Пам'ять, що розділяється, організована у вигляді 16 (всього-то!) банків пам'яті з кроком в 4 байти. Під час виконання пулу потоків warp на мультипроцесорі, він ділиться на дві половинки (якщо warp-size = 32) по 16 потоків, які здійснюють доступ до пам'яті по черзі.

Завдання в різних половинах warp не конфліктують з пам'яті, що розділяється. Через завдання однієї половинки пулу warp будуть звертатися до однакових банків пам'яті, виникнуть колізії і, як наслідок, падіння продуктивності. Завдання в межах однієї половинки warp можуть звертатися до різних ділянок пам'яті, що розділяється, з певним кроком.

Оптимальні кроки - 4, 12, 28, ..., 2 n-4 байт (рис. 8).



Рис. 8. Оптимальні кроки.

Чи не оптимальні кроки – 1, 8, 16, 32, ..., 2^n байт (рис. 9).



Рис. 9. Неоптимальні кроки

4.9. Мінімізація переміщень даних Host<=>Device

Намагайтеся якнайрідше передавати проміжні результати на host для обробки за допомогою CPU. Реалізуйте якщо не весь алгоритм, то принаймні його основну частину на GPU, залишаючи CPU лише керуючі завдання.

5. CPU/GPU переносима математична бібліотека

Автором цієї статті написана бібліотека MGML_MATH, що переноситься, для роботи з простими просторовими об'єктами, код якої працездатний як на пристрої, так і на хості.

Бібліотека MGML_MATH може бути використана як каркас для написання CPU/GPU переносних (або гібридних) систем розрахунку фізичних, графічних чи інших просторових завдань. Основна її перевага в тому, що один і той же код може використовуватися як на CPU, так і на GPU, і при цьому на чільне місце вимог, що пред'являються до бібліотеки, ставиться швидкість.

6 . Література

Кріс Касперскі. Техніка оптимізації програм. Ефективне використання пам'яті. – Спб.: БХВ-Петербург, 2003. – 464 с.: іл.

CUDA Programming Guide 1.1 ( http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

CUDA Programming Guide 1.1. page 14-15

CUDA Programming Guide 1.1. page 48

Згідно з Дарвінською теорією еволюції, перша людиноподібна мавпа (якщо
бути точним – homo antecessor, людина-попередник) перетворилася згодом
у нас. Багатотонні обчислювальні центри з тисячею і більше радіоламп,
що займають цілі кімнати, змінилися півкілограмовими ноутами, які, до речі,
не поступляться продуктивністю першим. Допотопні друкарські машинки перетворилися
у друкуючі що завгодно та на чому завгодно (навіть на тілі людини)
багатофункціональні пристрої. Процесорні гіганти раптом надумали замурувати
графічне ядро ​​у «камінь». А відеокарти стали не тільки показувати картинку з
прийнятним FPS та якістю графіки, але й проводити всілякі обчислення. Так
ще як робити! Про технологію багатопотокових обчислень засобами GPU, і йтиметься.

Чому GPU?

Цікаво, чому всю обчислювальну міць вирішили перекласти на графічний
адаптер? Як видно, процесори ще в моді, та й навряд чи поступляться своїм теплом
Містечко. Але у GPU є пара козирів у рукаві разом із джокером, та й рукавів
вистачає. Сучасний центральний процесор заточений під отримання максимальної
продуктивності при обробці цілісних даних та даних з плаваючою
коми, особливо не переймаючись при цьому паралельною обробкою інформації. У той же
час архітектура відеокарти дозволяє швидко і без проблем «розпаралелити»
обробку даних. З одного боку йде облік полігонів (за рахунок 3D-конвеєра),
з іншого – піксельна обробка текстур. Видно, що відбувається «злагоджена
розбивка» навантаження у ядрі карти. Крім того, робота пам'яті та відеопроцесора
оптимальніше, ніж зв'язування «ОЗУ-кеш-процесор». В той момент, коли одиниця даних
у відеокарті починає оброблятися одним потоковим процесором GPU, інша
одиниця паралельно завантажується в інший, і, в принципі, легко можна досягти
завантаженості графічного процесора, порівнянної з пропускною здатністю шини,
однак для цього завантаження конвеєрів має здійснюватися одноманітно, без
всяких умовних переходів та розгалужень. Центральний процесор через свою
універсальності вимагає для своїх процесорних потреб кеш, заповнений
інформацією.

Вчені чоловіки замислилися щодо роботи GPU у паралельних обчисленнях та
математики і вивели теорію, що багато наукових розрахунків багато в чому схожі з
обробкою 3D-графіки. Багато експертів вважають, що основним фактором у
розвитку GPGPU (General Purpose computation on GPU – універсальні
розрахунки засобами відеокарти) стала поява у 2003 році проекту Brook GPU.

Творцям проекту зі Стендфордського університету належало вирішити непросту
проблему: апаратно та програмно змусити графічний адаптер виробляти
різнопланові обчислення. І вони це вийшло. Використовуючи універсальну мову C,
американські вчені змусили працювати GPU як процесор, з поправкою на
паралельну обробку. Після Brook з'явився цілий ряд проектів за VGA-розрахунками,
таких як бібліотека Accelerator, бібліотека Brahma, система
метапрограмування GPU++ та інші.

CUDA!

Передчуття перспективності розробки змусило AMDі NVIDIA
вчепитися в Brook GPU, як пітбуль. Якщо опустити маркетингову політику, то,
реалізувавши все правильно, можна закріпитися не лише у графічному секторі
ринку, а й у обчислювальному (дивись на спеціальні обчислювальні карти та
сервери Teslaз сотнями мультипроцесорів), потіснивши звичні для всіх CPU.

Звичайно, «володарі FPS» розійшлися біля каменю спотикання кожен за своєю.
стежці, але основний принцип залишився незмінним – проводити обчислення
засобами GPU. І зараз ми розглянемо технологію «зелених». CUDA
(Compute Unified Device Architecture).

Робота нашої «героїні» полягає у забезпеченні API, причому одразу двох.
Перший – високорівневий, CUDA Runtime, є функціями, які
розбиваються більш прості рівні і передаються нижньому API – CUDA Driver. Так
що фраза «високрівневий» застосовна до процесу з натяжкою. Вся сіль знаходиться
саме у драйвері, і добути її допоможуть бібліотеки, люб'язно створені
розробниками NVIDIA: CUBLAS (засоби для математичних розрахунків) та
FFT (розрахунок у вигляді алгоритму Фур'є). Ну що ж, перейдемо до практичної
частини матеріалу.

Термінологія CUDA

NVIDIAоперує дуже своєрідними визначеннями для CUDA API. Вони
відрізняються від визначень, що застосовуються до роботи з центральним процесором.

Потік (thread)- Набір даних, який необхідно обробити (не
вимагає великих ресурсів під час обробки).

Варп (warp)- Група з 32 потоків. Дані обробляються тільки
варпами, отже варп - це мінімальний обсяг даних.

Блок (block)- Сукупність потоків (від 64 до 512) або сукупність
варпів (від 2 до 16).

Сітка (grid)- Це сукупність блоків. Такий поділ даних
застосовується виключно підвищення продуктивності. Так, якщо число
мультипроцесорів велике, то блоки виконуватимуться паралельно. Якщо ж з
карткою не пощастило (розробники рекомендують для складних розрахунків використовувати
адаптер не нижче рівня GeForce 8800 (GTS 320 Мб), то блоки даних обробляться
послідовно.

Також NVIDIA вводить такі поняття, як ядро (kernel), хост (host)
і девайс (device).

Працюємо!

Для повноцінної роботи з CUDA потрібно:

1. Знати будову шейдерних ядер GPU, оскільки суть програмування
полягає у рівномірному розподілі навантаження між ними.
2. Вміти програмувати серед С, з урахуванням деяких аспектів.

Розробники NVIDIAрозкрили «начинки» відеокарти кілька
інакше, ніж ми звикли бачити. Так що хоч-не-хоч доведеться вивчати все
тонкощі архітектури. Розберемо будову «каменю» G80 легендарною GeForce 8800
GTX.

Шейдерне ядро ​​складається з восьми TPC (Texture Processor Cluster) – кластерів
текстурних процесорів (так, у GeForce GTX 280- 15 ядер, у 8800 GTS
їх шість, у 8600 - Чотири і т.д.). Ті, у свою чергу, складаються з двох
потокових мультипроцесорів (streaming multiprocessor – далі SM). SM (їх всього
16) складається з front end (вирішує завдання читання та декодування інструкцій) та
back end (кінцевий виведення інструкцій) конвеєрів, а також восьми scalar SP (shader
processor) та двома SFU (суперфункціональні блоки). За кожен такт (одиницю
часу) front end вибирає варп та обробляє його. Щоб усі потоки варпа
(Нагадаю, їх 32 штуки) обробилися, потрібно 32/8 = 4 такти в кінці конвеєра.

Кожен мультипроцесор має так звану загальну пам'ять (shared memory).
Її розмір становить 16 кілобайт та надає програмісту повну свободу
дій. Розподіляй як хочеш:). Shared memory забезпечує зв'язок потоків у
одному блоці та не призначена для роботи з піксельними шейдерами.

Також SM можуть звертатися до GDDR. Для цього їм «пришили» по 8 кілобайт
кеш-пам'яті, що зберігають все найголовніше для роботи (наприклад, обчислювальні
константи).

Мультипроцесор має 8192 регістри. Число активних блоків не може бути
більше восьми, а число варпів - не більше 768/32 = 24. З цього видно, що G80
може обробити максимум 32*16*24 = 12 288 потоків за одиницю часу. Не можна не
враховувати ці цифри при оптимізації програми надалі (на одній чаші ваг
- Розмір блоку, на інший - кількість потоків). Баланс параметрів може зіграти
важливу роль надалі, тому NVIDIAрекомендує використовувати блоки
зі 128 або 256 потоками. Блок з 512 потоків неефективний, оскільки має
підвищеними затримками. Враховуючи всі тонкощі будови GPU відеокарти плюс
непогані навички у програмуванні, можна створити вельми продуктивне
засіб для паралельних обчислень. До речі, про програмування...

Програмування

Для «творчості» разом із CUDA потрібно відеокарта GeForce не нижче
восьмий серії. З

офіційного сайту потрібно завантажити три програмні пакети: драйвер з
підтримкою CUDA (для кожної ОС – свій) безпосередньо пакет CUDA SDK (друга
бета-версія) та додаткові бібліотеки (CUDA toolkit). Технологія підтримує
операційні системи Windows (XP та Vista), Linux та Mac OS X. Для вивчення я
вибрав Vista Ultimate Edition x64 (забігаючи вперед, скажу, що система поводилася
просто чудово). У момент написання цих рядків актуальним для роботи був
драйвер для ForceWare 177.35. Як набор інструментів використовувався
програмний пакет Borland C++ 6 Builder (хоча підійде будь-яке середовище, що працює з
мовою C).

Людині, яка знає мову, легко освоїться в новому середовищі. Потрібно лише
запам'ятати основні параметри. Ключове слово _global_ (ставиться перед функцією)
показує, що функція належить до kernel (ядру). Її викликатиме центральний
процесор, а вся робота відбудеться на GPU. Виклик _global_ вимагає більше
конкретних деталей, а саме розмір сітки, розмір блоку та яке ядро ​​буде
застосовано. Наприклад, рядок _global_ void saxpy_parallel<<>>, де X –
розмір сітки, а Y – розмір блоку, визначає ці параметри.

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

Також мова для CUDA має низку функцій для роботи з відеопам'яттю: cudafree
(звільнення пам'яті між GDDR і RAM), cudamemcpy та cudamemcpy2D (копіювання
пам'яті між GDDR і RAM) і cudamalloc (виділення пам'яті).

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

Набір для вивчення

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

Спеціально для новачків розроблено CUDA SDK Browser. Будь-який бажаючий може
відчути силу паралельних обчислень на своїй шкурі (краща перевірка на
стабільність – робота прикладів без артефактів та вильотів). Додаток має
великий ряд показових міні-програм (61 «тест»). До кожного досвіду є
докладна документація програмного коду плюс PDF-файли. Відразу видно, що люди,
присутні зі своїми витворами в браузері, займаються серйозною роботою.
Тут же можна порівняти швидкості роботи процесора та відеокарти під час обробки
даних. Наприклад, сканування багатовимірних масивів відеокартою GeForce 8800
GT 512 Мб з блоком з 256 потоками виробляє за 0.17109 мілісекунди.
Технологія не розпізнає SLI-тандеми, тому якщо у тебе дует або тріо,
відключай функцію «спарювання» перед роботою, інакше CUDA побачить лише один
девайс. Двоядерний AMD Athlon 64 X2(Частота ядра 3000 МГц) той же досвід
проходить за 2.761528 мілісекунди. Виходить, що G92 більш ніж у 16 ​​разів
швидше «каменю» AMD! Як бачиш, далеко не екстремальна система в
тандемі з нелюбимою в масах операційною системою показує непогані
результати.

Крім браузера існує низка корисних суспільству програм. Adobe
адаптувала свої продукти до нової технології. Тепер Photoshop CS4 у повній
мірою використовує ресурси графічних адаптерів (необхідно завантажити спеціальний
плагін). Такими програмами, як Badaboom media converter та RapiHD, можна
зробити декодування відео формату MPEG-2. Для обробки звуку непогано
підійде безкоштовна утиліта Accelero. Кількість софту, заточеного під CUDA API,
безперечно, буде рости.

А в цей час…

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

Творіння «мікродевайсерів», Fusion, складатиметься з кількох ядер під
кодовою назвою Bulldozer та відеочіпа RV710 (Kong). Їхній взаємозв'язок буде
здійснюватись за рахунок покращеної шини HyperTransport. Залежно від
кількості ядер та їх частотних характеристик AMD планує створити цілу цінову
ієрархію «каменів». Також планується виробляти процесори як для ноутбуків (Falcon),
і для мультимедійних гаджетів (Bobcat). Причому саме застосування технології
у портативних пристроях буде первісним завданням для канадців. З розвитком
Паралельних обчислень застосування таких «камінців» має бути дуже популярним.

Intelтрішки відстає за часом зі своєю Larrabee. Продукти AMD,
якщо нічого не трапиться, з'являться на прилавках магазинів наприкінці 2009 – на початку
2010 року. А рішення супротивника вийде на світ божий лише майже через два
року.

Larrabee налічуватиме велику кількість (читай – сотні) ядер. На початку
ж вийдуть товари, розраховані на 8 – 64 ядер. Вони дуже подібні до Pentium, але
досить сильно перероблено. Кожне ядро ​​має 256 кілобайт кешу другого рівня
(Згодом його розмір збільшиться). Взаємозв'язок здійснюватиметься за рахунок
1024-бітної двонаправленої кільцевої шини. Інтел каже, що їхнє «дитя» буде
добре працювати з DirectX і Open GL API (для «яблучників»), тому ніяких
програмних втручань не потрібно.

А чого я все це тобі розповів? Очевидно, що Larrabee та Fusion не витіснять
звичайні, стаціонарні процесори з ринку, так само, як не витіснять з ринку
відеокарти. Для геймерів та екстремалів межею мрій, як і раніше, залишиться
багатоядерний CPU та тандем з декількох топових VGA. Але те, що навіть
процесорні компанії переходять на паралельні обчислення за принципами,
аналогічним GPGPU, говорить вже багато про що. Зокрема про те, що така
технологія, як CUDA, має право на існування і, мабуть, буде
дуже популярна.

Невелике резюме

Паралельні обчислення засобами відеокарти – лише хороший інструмент
в руках роботящого програміста. Навряд чи процесорам на чолі із законом Мура
настане кінець. Компанії NVIDIAналежить пройти ще довгий шлях по
просуванню в маси свого API (те ж можна сказати і про дітище ATI/AMD).
Якою вона буде, покаже майбутнє. Отже, CUDA will be back:).

P.S. Початківцям програмістам і людям, що зацікавилися, рекомендую відвідати
наступні «віртуальні заклади»:

офіційний сайт NVIDIA та сайт
GPGPU.com. Вся
дана інформація – англійською мовою, але, спасибі хоча б, що не
китайською. Так що дерзай! Сподіваюся, що автор хоч трохи допоміг тобі
захоплюючих починаннях пізнання 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, GTX140 500
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 Сімейство