|
|
Введення в CUDA API. Розширення мови C ⇐ ПредыдущаяСтр 6 из 6
- специфікаторів функцій, що показують, де буде виконуватися функція і звідки вона може бути викликана; - специфікаторів змінних, які задають тип пам'яті, що використовується для цих змінних; - директиви, що служить для запуску ядра, яка задає як дані, так і ієрархію ниток; - вбудованих змінних, що містять інформацію про поточну нитку; - ruпtime, що включає в себе додаткові типи даних.(Табл. 6.2)
Таблиця 6.2 — Специфікатори функцій та змінних
Специфікатори __host__ і __device__ можуть бути використані разом (це означає, що відповідна функція може виконуватися як на GPU, так і на CPU - відповідний код для обох платформ буде автоматично згенерований компілятором). Специфікатори __global__ і __host__ не можуть бути використані разом. Специфікатор __global__ позначає ядро, і відповідна функція повинна повертати значення типу void. На функції, що виконуються на GPU (__ device__ і __global__), накладаються такі обмеження: - не можна брати їх адресу (за винятком __global__ функцій); - не підтримується рекурсія; - не підтримуються stаtiс-змінні всередині функції; - не підтримується змінне число вхідних аргументів. Для завдання розміщення в пам'яті GPU змінних використовуються наступні специфікатори __device__, __constant__ і __shared__. На їх використання також накладається ряд обмежень: - ці специфікатор не можуть бути застосовані до полів структури (struct або union); - відповідні змінні можуть використовуватися тільки в межах одного файлу, їх не можна оголошувати як extern; - запис в змінні типу __constant__ може здійснюватися тільки CPU за допомогою спеціальних функцій; - __shared__ змінні не можуть бути ініціалізовані при оголошенні.
Додані типи
Звернення до компонентів вектора йде поіменно: х, у, z і w. Для створення значень-векторів заданого типу служить конструкція виду make_<typeName>.
int2 а = make_int2(1, 7);
Для цих типів не підтримуються векторні покомпонентні операції, тобто не можна просто скласти два вектора за допомогою оператора «+» - це необхідно явно робити для кожної компоненти. Також доданий тип dim3, що використовується для завдання розмірності. Цей тип заснований на типі uint3, але має при цьому нормальний конструктор, що ініціалізує всі не задані компоненти одиницями.
Dim3 blocks(16, 16); // Эквивалентно blocks(16, 16, 1). Dim3 grid(256); // Эквивалентно grid(256, 1, 1).
Додані змінні
- gridDim - розмір сітки (має тип dimз); - blockDim - розмір блоку (має тип dim3); - blockldx - індекс поточного блоку в сітці (має тип uint3); - threadldx - індекс поточної нитки в блоці (має тип uint3); - warpSize - розмір warp'a (має тип int).
Директива виклику ядра
- kenelname - це ім'я відповідної __global__ функції; - через Dg позначена змінна (або значення) типу dim3, що задає розмірність і розмір сітки (в блоках); - змінна (або значення) Db - типу dim3 задає розмірність і розмір блоку (в нитках); - необов'язкова змінна (або значення) Ns типу size_t задає додатковий обсяг пам'яті в байтах, яка повинна бути динамічно виділена кожному блоку (до вже статично виділеної пам'яті, що розділяється), якщо не задано, то використовується значення 0; - змінна (або значення) S типу cudaStream_t задає потік (CUDA stream), в якому повинен відбутися виклик, за замовчуванням використовується потік 0; - через args позначені аргументи виклику функції kernelname (їх може бути кілька). Наступний приклад запускає ядро з ім'ям mуКеrnel паралельно на n нитках, використовуючи одномірний масив з двовимірних (16х16) блоків ниток, і передає на вхід ядру два параметри: а та n. При цьому кожному блоку додатково виділяється 512 байт пам'яті. Запуск проводиться в потоці myStream:
mуКеrnеl<<<dimЗ(n/25б), dimЗ(l6,l6), 512, myStream>>> (а, n);
Додані функції
Крім тoгo, CUDA надає також спеціальний набір функцій зниженою точності, які забезпечують ще більшу швидкодію. Таким aнaлогом для функції обчислення синуса також є функція sinf. Основи CUDA host API CUDA надає в розпорядження програміста ряд функцій, які можуть бути використані тільки CPU (так званий CUDA host АРI). Ці функції відповідають за: - управління GPU; - роботу з контекстом; - роботу з пам'яттю; - роботу з модулями; - управління виконанням коду; - роботу з текстурами; - взаємодія з OpenGL і Direct3D. CUDA host АРI виступає в двох формах: - низькорівневий CUDA driver API; - високорівнева CUDA runtime АРI (реалізований через CUDA driver API). Ці АРI є взаємовиключними. У своїй програмі ви можете працювати тільки з одним з них.
CUDA driver API Низькорівневий API дає більше можливостей програмісту, але й вимагає більшого обсягу коду. Даний АРI реалізований в динамічній бібліотеці nvcuda, і всі імена в ньому починаються із префікса cu. Слід мати на увазі, що у кожної функції CUDA runtime АРI є прямий аналог в CUDA driver АРI, тобто перехід з CUDA runtime АРI на CUDA driver АРI не дуже складний, зворотне в загальному випадку невірно. CUDA driver АРI має зворотну сумісність з більш ранніми версіями. До числа недоліків цього АРI відносяться більший обсяг коду і необхідність явних налаштувань, вимога явної ініціалізації.
CUDA runtime API Це високорівневий API, до того ж CUDA runtime АРI не вимагає явної ініціалізації - вона відбувається автоматично при першому виклику будь-якої його функції. Даний АРI реалізований в динамічній бібліотеці cudart, все імена починаються із префікса cuda. Одним з плюсів даного АРI є можливість використання додаткових бібліотек (CUFFT, CUBLAS, CUDPP та інших). Ми будемо використовувати саме цей АРI як більш простий. При необхідності код, заснований на ньому, може бути переписаний на CUDA driver АРI, оскільки останній містить повні аналоги всіх функцій CUDA runtime АРI. Розглянемо основні елементи CUDA runtime API на прикладі наступної програми, яка виконує операцію складання матриць:
// Ядро __global__ void vectorAdd(float a*, float *b, float *c) { int index = blockIdx.x * blockDim.x + threadIdx.x; с[index] = а[index] + b[index]; } int main() { const unsigned int blockSize = 512; const unsigned int numBlocks = 3; const unsigned int numItems = numBlocks * blockSize; cudaSetDevice(0); float *a = new float[numItems]; float *b = new float[numItems]; float *c = new float[numItems]; // Инициализация массивов //... float *aDev, *bDev, *cDev; cudaMalloc((void**)&aDev, numItems * sizeof(float)); cudaMalloc((void**)&bDev, numItems * sizeof(float)); cudaMalloc((void**)&cDev, numItems * sizeof(float)); cudaMemcpy(aDev, а, numItems * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(bDev, b, numItems * sizeof(float), cudaMemcpyHostToDevice); vectorAdd<<<nurnBlocks,blockSize>>> (aDev, bDev, cDev); cudaMemcpy((void*)с, cDev, numItems * sizeof(float), cudaMemcpyDeviceToHost); delete[] a; delete[] b; delete[] c; cudaFree(aDev); cudaFree(bDev); cudaFree(cDev); } Кожна функція CUDA runtime АРI (крім запуску ядра) повертає значення типу cudaError_t. При успішному виконанні функції повертається значення cudaSuccess, в іншому випадку повертається код помилки. Отримати опис помилки у вигляді рядка по її коду можна за допомогою функції cudaGetErrorString:
char* cudaGetErrorString(cudaError_t code);
Також можна отримати код останньої помилки за допомогою функції cudaGetLastError:
cudaError_t cudaGetLastError(); Важливим моментом роботи з CUDA, на який слід відразу ж звернути увагу, є те, що багато функцій АРІ - асинхронні, тобто управління повертається ще до реального завершення необхідної операції.
- запуск ядра; - функції копіювання пам'яті, імена яких закінчуються на Async; - функції копіювання пам'яті device ↔ device - функції ініціалізації пам'яті.
6.2 Завдання 6.2.1 Розробити паралельну програму інтегрування методом Монте-Карло, використовуючи технологію CUDA Дослідити своє рішення аналогічно п. 5.2.3 Завдання 2. 6.2.2 Отримати індивідуальне завдання у викладача.
6.3 Зміст звіту
6.3.1 Мета лабораторної роботи. 6.3.2 Тексти програм. 6.3.3 Результати розрахунків. 6.3.4 Відповіді на контрольні питання.
6.4 Контрольні питання
6.4.1 В чому полягає сутність технології CUDA? 6.4.2 Види пам'яті в CUDA. 6.4.3 Поясніть використання вбудованих змінних threadldx, blockldx,. gridDim і blockDim. 6.4.4 Поясніть синтаксис директиви виклику ядра. 6.4.5 Які Ви знаєте додані типи, змінні та функції? 6.4.6 Поясніть різницю між CUDA driver API та CUDA runtime API. 6.4.7 Як отримати опис останньої помилки у вигляді рядка?
ЛІТЕРАТУРА 1. Немнюгин С.А., Стесик О.Л. Параллельное программирование для многопроцессорных вычислительных систем. – СПб.: БХВ – Санкт-Петербург, 2002. – 400 с. 2. Эндрюс Г.Р. Основы многопоточного, параллельного и распределенного программирования.: Пер. с англ. – М.: Издательский дом "Вильямс", 2003. – 512 с. 3. Шеховцов В.А. Операційні системи. – К.: Видавнича група BHV, 2005. - 576 C. 4. Гергель В.П., Стронгин Р.Г. Основы параллельных вычислений для многопроцессорных вычислительных систем. Учебное пособие – Нижний Новгород; Изд-во ННГУ им. Н.И. Лобачевского, 2000. –176 с. 5. Parallel.ru – информационно-аналитический центр по параллельным вычислениям – http://www.parallel.ru. 6. Воеводин В.В., Воеводин Вл.В. Параллельные вычисления. – СПб.: БХВ-Петербург, 2002. – 608 с. 7. Корнеев В.Д. Параллельное программирование в MPI. – 2-е изд., испр. – Новосибирск: Изд-во ИВМиМГ СО РАН, 2002. – 215 с. 8. Бройнль Т. Паралельне програмування: Початковий курс: Навч. посібник / Вступ. Слово А.Ройтера; Пер. з нім. В.А. Святного. – К.: Вища школа, 1997. – 358 с. 9. Букатов А.А., Дацюк В.Н., Жегуло А.И. Программирование многопроцессорных вычислительных систем. – Ростов-на-Дону: Изд-во ООО "ЦВВР", 2003. – 208 с. 10. Ортега Дж. Введение в параллельные и векторные методы решения линейных систем. – М.: Мир, 1991. – 367 с. 11. Корнеев В.В. Параллельные вычислительные системы. – М.: Нолидж, 1999. – 320 с. 12.. Параллельные вычисления на GPU. Архитектура и программная модель CUDA[Текст]: Учеб. пособие / А.В. Боресков и др. Предисл.: В.А. Садовничий. – М.: Издательство Московского университета, 2012. – 336 с.
Додаток А Текст програми Thread War // Проста комп'ютерна гра Thread War // Використовуйте клавіші "уліво" і "вправо", щоб переміщати пушку // клавіша "пробіл" робить постріл, // Якщо 30 ворогів підуть із екрана не знищеними, ви програли // Очки даються за кожного вбитого супротивника
#include <windows.h> #include <process.h> #include <stdlib.h> #include <time.h> #include <stdio.h>
// Об'єкти синхронізації HANDLE screenlock; // зміною екрана займається тільки один потік HANDLE bulletsem; // можна вистрілити тільки три рази підряд HANDLE startevt; // гра починається з натисканням клавіші "уліво" або вправо" HANDLE conin, conout; // дескриптори консолі HANDLE mainthread; // основний потік main CRITICAL_SECTION gameover;
CONSOLE_SCREEN_BUFFER_INFO info; // інформація про консоль // кількість влучень і промахів long hit=0; long miss=0; long delayfactor=7; // фактор затримки для ворогів
// Створення випадкового числа від n0 до n1 int random(int n0, int n1) { if (n0==0 && n1==1) return rand()%2; // спеціальний випадок return rand()%(n1-n0)+n0; } // Очищення екрана консолі void cls() { COORD org={0,0}; DWORD res; FillConsoleOutputCharacter(conout, ' ', info.dwSize.X*info.dwSize.Y, org, &res); }
// вивести на екран символ в позицію х и y void writeat(int x, int y, char c) { // Блокувати вивід на екран за допомогою м’ютекса WaitForSingleObject(screenlock, INFINITE); COORD pos={x,y}; DWORD res; WriteConsoleOutputCharacter(conout, &c, 1, pos, &res); ReleaseMutex(screenlock); }
// Одержати натискання на клавішу (лічильник повторень в ct) int getakey(int &ct) { INPUT_RECORD input; DWORD res; while (1) { ReadConsoleInput(conin,&input, 1, &res);
// ігнорувати інші події if (input.EventType!=KEY_EVENT) continue;
// ігнорувати події відпускання клавіш // нас цікавлять тільки натискання if (!input.Event.KeyEvent.bKeyDown) continue; ct=input.Event.KeyEvent.wRepeatCount; return input.Event.KeyEvent.wVirtualKeyCode; } }
// Обробка комбінацій ^C, ^Break, і т.і. BOOL WINAPI ctrl(DWORD type) { exit(0); return TRUE; // не досяжна ділянка коду }
// Визначити символ в заданій позиції екрана int getat(int x, int y) { char c; DWORD res; COORD org={x,y};
// Блокувати доступ до консолі доти, поки процедура не буде виконана WaitForSingleObject(screenlock,INFINITE); ReadConsoleOutputCharacter(conout, &c, 1, org, &res); ReleaseMutex(screenlock); // unlock return c; }
// Відобразити очки в заголовку вікна й перевірити умову завершення гри void score(void) { char s[128]; sprintf(s, "Thread War! Hit: %d Miss: %d", hit, miss); SetConsoleTitle(s); if (miss>=30) { EnterCriticalSection(&gameover); SuspendThread(mainthread); // призупинити головний потік MessageBox(NULL, "Game Over!", "Thread War", MB_OK|MB_SETFOREGROUND); exit(0); // не виходить із критичної секції } if ((hit+miss)%20==0) InterlockedDecrement(&delayfactor); // повинен бути ilock }
char badchar[]="-\\|/"; // це потік супротивника. void badguy(void *_y) { int y=(int) _y; // випадкова координата y int dir; int x; // непарні y з'являються ліворуч, парні y з'являються праворуч x=y%2?0:info.dwSize.X; // установити напрямок залежно від початкової позиції dir=x?-1:1; //поки супротивник перебуває в межах екрана while ((dir==1&&x!=info.dwSize.X)||(dir==-1&&x!=0)) { int dly; BOOL hitme=FALSE; // перевірка на влучення (куля?) if (getat(x,y)=='*') hitme=TRUE;
// вивід символу на екран writeat(x,y,badchar[x%4]);
// ще одна перевірка на влучення if (getat(x,y)=='*') hitme=TRUE; // перевірка на влучення через невеликі // проміжки часу if (delayfactor<3) dly=3; else dly=delayfactor+3; for (int i=0; i<dly; i++) { Sleep(40); if (getat(x,y)=='*') { hitme=TRUE; break; } } writeat(x,y,' '); // ще одна перевірка на влучення if (getat(x,y)=='*') hitme=TRUE; if (hitme) { // у супротивника влучили! MessageBeep(-1); InterlockedIncrement(&hit); score(); _endthread(); } x+=dir; }
//супротивник утік! InterlockedIncrement(&miss); score(); }
// цей потік займається створенням потоків супротивників void badguys(void *) { // чекаємо сигналу до початку гри протягом 15 секунд WaitForSingleObject(startevt, 15000); // створюємо випадкового ворога // кожні 5 секунд з'являється шанс створити //супротивника з координатами від 1 до 10 while (1) { if (random(0,100)<(hit+miss)/25+20) // згодом імовірність збільшується _beginthread(badguy,0, (void *)(random(1,10))); Sleep(1000); // щосекунди } }
// Це потік кулі // кожна куля - це окремий потік void bullet(void *_xy_) { COORD xy=*(COORD *)_xy_; if (getat(xy.X, xy.Y)== '*') return; // тут уже є куля // треба почекати // перевірити семафор // якщо семафор дорівнює 0, пострілу не відбувається if (WaitForSingleObject(bulletsem,0)==WAIT_TIMEOUT) return; while (-іxy.Y) { writeat(xy.X, xy.Y, '*'); // відобразити кулю Sleep(100); writeat(xy.X, xy.Y,' '); // стерти кулю } // постріл зроблений - додати 1 до семафора ReleaseSemaphore(bulletsem,1,NULL); }
// Основна програма void main() { HANDLE me; // Настроювання глобальних змінних conin=GetStdHandle(STD_INPUT_HANDLE); conout=GetStdHandle(STD_OUTPUT_HANDLE); SetConsoleCtrlHandler(ctrl,TRUE); SetConsoleMode(conin,ENABLE_WINDOW_INPUT); me=GetCurrentThread(); // не є реальним дескриптором
// змінити псевдодескриптор на реальний дескриптор поточного потоку DuplicateHandle(GetCurrentProcess(), me, GetCurrentProcess(), &mainthread, 0, FALSE, DUPLICATE_SAME_ACCESS);
startevt=CreateEvent(NULL,TRUE,FALSE,NULL); screenlock=CreateMutex(NULL,FALSE,NULL); InitializeCriticalSection(&gameover); bulletsem=CreateSemaphore(NULL,3,3,NULL); GetConsoleScreenBufferInfo(conout,&info);
// Ініціалізувати відображення інформації про очки score(); // Настроїти генератор псевдовипадкових чисел srand((unsigned)time(NULL)); cls(); // насправді не потрібно // установка початкової позиції пушки int y=info.dwSize.Y-1; int x=info.dwSize.X/2; //запустити потік badguys; нічого не робити доти, // поки не відбудеться подія або минуть 15 секунд _beginthread(badguys,0,NULL); // основний цикл гри while (1) { int c,ct; writeat(x,y,'|'); // намалювати пушку c=getakey(ct); // одержати символ switch (c) { case VK_SPACE: // вогонь! { static COORD xy; xy.X=x; xy.Y=y;
_beginthread(bullet, 0, (void *) &xy); Sleep(100); // дати кулі час полетіти на деяку відстань break; } case VK_LEFT: // команда "уліво!" SetEvent(startevt); // потік badguys працює writeat(x,y,' '); // зтерти з екрана пушку while (ct--) // переміститися if (x) x--; break; case VK_RIGHT: // команда "вправо!"; логіка та ж SetEvent(startevt); writeat(x,y,' '); while (ct-і) if (x!=info.dwSize.X-1) x++; break; } } }   Конфликты в семейной жизни. Как это изменить? Редкий брак и взаимоотношения существуют без конфликтов и напряженности. Через это проходят все...  Что способствует осуществлению желаний? Стопроцентная, непоколебимая уверенность в своем...  ЧТО ПРОИСХОДИТ, КОГДА МЫ ССОРИМСЯ Не понимая различий, существующих между мужчинами и женщинами, очень легко довести дело до ссоры...  ЧТО ТАКОЕ УВЕРЕННОЕ ПОВЕДЕНИЕ В МЕЖЛИЧНОСТНЫХ ОТНОШЕНИЯХ? Исторически существует три основных модели различий, существующих между... Не нашли то, что искали? Воспользуйтесь поиском гугл на сайте:
|