Подякувати Студентському архіву довільною сумою
Admin
26.02.2023 12:38
Дякуємо, що користуєтесь нашим архівом!
Реалізація алгоритмів пошуку на графічному процесорі
Інформація про навчальний заклад
ВУЗ:
Національний університет Львівська політехніка
Інститут:
Не вказано
Факультет:
Не вказано
Кафедра:
Не вказано
Інформація про роботу
Рік:
2024
Тип роботи:
Розрахункова робота
Предмет:
СП
Група:
ІУСм-12
Частина тексту файла (без зображень, графіків і формул):
МІНІСТЕРСТВО ОСВІТИ І НАУКИ УКРАЇНИ
НАЦІОНАЛЬНИЙ УНІВЕРСИТЕТ «ЛЬВІВСЬКА ПОЛІТЕХНІКА»
Розрахункова робота
з дисципліни
«Проблемно-орієнтовані обчислювальні
комплекси та системи»
на тему
«Реалізація алгоритмів пошуку на графічному процесорі»
ЗМІСТ
І. Бінарний пошук 3
1.1. Бінарний пошук у CPU 3
1.2. Бінарний пошук у GPU 4
ІІ. Реалізація алгоритму SSSP на GPU 11
ІІІ. Реалізація алгоритму BFS на GPU 18
СПИСОК ВИКОРИСТАНОЇ ЛІТЕРАТУРИ 23
І. Бінарний пошук
Бінарний пошук у CPU
Класична реалізація бінарного пошуку у CPU досить проста в програмуванні. Інші модифікації і поліпшення можуть виявитися корисними лише за певних умов. Буде розглянуто загальний випадок. По-друге: максимум за величину порядку ln2 N порівнянь буде знайдено шуканий ключ (або установлено, що його немає в таблиці). Не багато інших алгоритмів зможуть потягатися з такою швидкодією. І так маємо в самому гіршому випадку ln2 N порівнянь. Правда алгоритм застосуємо лише для відсортованого списку. Реалізація на CPU має наступний вигляд:
__host__ void searchCPU ( float * data, int * result )
{
int x = 16777200; // шуканий елемент
int i = 0; // ліва границя масиву
int j = 16 * 1024 * 1024; // права границя масиву
int k = 0;
while ( i <= j )
{
k = i + (j - i) / 2 ;
if ( data[k] < x )
{ i = k + 1; }
else
if ( data[k] > x )
{ j = k - 1; }
else
{
result[0] = data [k];
result[1] = k;
break;
}
}
}
Як можна побачити, дійсно нічого складного. Даний алгоритм значно перевершує по швидкодії простий лінійний пошук (у багато сотень разів), реалізований навіть на GPU, адже для лінійного пошуку маємо N порівнянь в гіршому випадку. Але, як вже говорилося вище, він можливий лише для відсортованого списку (про сортування трохи нижче). До речі кажучи, цей алгоритм набагато стійкіший, оскільки для нього практично не має значення в якій частині списку перебуває елемент, тоді як для лінійного – від цього безпосередньо залежить час пошуку. Але це в реалізації на CPU! На GPU для лінійного пошуку маємо протилежний випадок тобто досить стійкий алгоритм, тому безліч потоків може майже одночасно охопити різні ділянки списку.
Бінарний пошук у GPU
Графік порівняння продуктивності для CPU і GPU, зображений на рисунку 1.
/
Рисунок 1. Графік порівняння продуктивності для CPU і GPU
Як бачимо з графіка, швидкодія GPU починає перевершувати свій аналог на CPU лише на певній позначці, тобто при досить великих обсягах динних. Спробуємо реалізувати бінарний пошук для GPU і підтвердити або спростувати дані у вищевказаному графіку. Але яким чином розділити обчислення між потоками? Перша думка була найпростіша. Розділити весь список на рівні ділянки, в кожній ділянці виконувати бінарний пошук. Потрібно зауважити, що чим менше потоків, тим швидше виконується програма, тобто її швидкодія найбільша при одному потоці.
Алгоритм такий: ділимо весь список на кількість потоків, Кожен потік дивиться один елемент і порівнює його з шуканим. Якщо це шуканий елемент чудово. Якщо поточний елемент менший, запам'ятовуємо його як нижню межу. Якщо більший запам'ятовуємо як верхню межу. Кожен з потоків знайде свій елемент більший або менший шуканого, а нам потрібні тільки граничні елементи, між якими лежить шуканий, отже, потрібно серед менших елементів залишити максимум, а серед більших мінімум. Далі повторюємо процедуру, але вже не для всього списку а для нових його меж. Повторюємо доти поки або не знайдемо елемент, або кількість елементів між нижньою і верхньою межею стане меншою кількості потоків. У такому випадку для решти елементів застосовуємо будь-який інший алгоритм, наприклад однопотоковий бінарний пошук. Розглянемо приклад.
Нехай маємо 2 048 елементів і 32 потоки. Для зручності індекси елементів у масиві і їх значення будемо вважати однаковими. Шукаємо, скажімо, елемент під номером 65 рівний 65. І так перший крок. Для кожного потоку дивимося кожен 32 елемент. тобто перший потік дивиться перший елемент, другий елемент номер 32, третій елемент номер 64 і т.д. Кожен потік знаходить або більший або менший елемент. Серед менших елементів максимальним є 64, серед більших мінімальним є 96. І так маємо нові кордони списку 64..96. Кількість елементів між кордонами рівне 96-64 = 32. Повторюємо ітерацію. 32 елемента ділимо на кількість потоків, кожен потік дивиться свій елемент. Отримує що кожен потік дивиться по одному елементу і знаходить шуканий.
Разом маємо лише два кроки! Теоретично непогано. Дійсно якщо оцінити алгоритм теоретично, то маємо швидкість O = (ln кількість потоків N + ln2 кількість процесорів) в гіршому випадку. Теоретично дуже непогано. Реалізація приведена нижче.
__global__ void searchGPU ( float * data, /*float * search,*/ int * result )
{
int x = 16 * 1024 * 1024; // кількість елементів
int mx = 0; // нижня межа
int mn = x-1; // верхня межа
int step; // крок
x = 16777200; // шуканий елемент
int idx; // даний потік
while ( mn - mx > 32 ) // поки кількість елементів між границями
//більше кількості потоків (в даному випадку потоків 32)
{
step = (mn - mx) / 32;
nt idx = mx + (blockIdx.x * blockDim.x + threadIdx.x) * step;
if ( data[idx] == x ) // найдений шуканий елемент
{
result[0] = data [idx];
result[1] = idx;
break;
}else
if ( data[idx] > x )
{
if(data[idx] < mn)
mn = data[idx]; // мінімальний серед більших
}else
if(data[idx] < x)
{
if(data[idx] > mx)
mx = data[idx]; // максимальний серед меньших
}
}
}
Теоретичний аналіз говорить, що алгоритм повинен мати дуже хорошу швидкість, проте, на жаль, все ще набагато меншу швидкості бінарного пошуку ніж для CPU.
У чому ж справа? Причин багато. Наприклад, копіювання даних на відеокарту і назад. Ресурси на запуск потоків. Затримка при довільному читанні елементів з пам'яті. Всі ці фактори складно враховувати при теоретичному аналізі алгоритму. Однак цілком можливо, що реалізація даного алгоритму перевершить бінарний пошук для CPU в задачах зовнішнього пошуку. Коли надходить дуже багато даних із зовнішнього накопичувача, а не з оперативної пам'яті. Але тут потрібно враховувати багато інших факторів, наприклад швидкість читання даних з пристрою, організацію зберігання самих даних на пристрої (та й дані ще повинні бути відсортовані), і орієнтувати алгоритм вже під ці умови.
Отже, алгоритм бінарного пошуку не дуже хороший в реалізації на GPU. Але це у випадку, якщо ми шукаємо один елемент в масиві. А якщо нам потрібно знайти масив елементів в масиві? Таке завдання ідеально підходить для GPU. Тоді ми організуємо алгоритм таким чином, що б кожен потік шукав свій елемент, що буде відбуватися майже одночасно. Теоретично маємо все ту ж швидкість – O = (ln2N). Тоді як для CPU реалізація представляє з себе послідовний двійковий пошук кожного елемента, що складе O = (ln2N * M), де M кількість шуканих елементів. Приклад реалізації приведено нижче.
__global__ void searchGPU ( float * data, float * search)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
//data масив даних
//search шуканий масив
int i = 0;
int j = 16 * 1024 * 1024;
int k = 0;
while ( i <= j )
{
k = i + (j - i) / 2 ;
if ( data[k] < search[idx] )
{
i = k + 1;
}
else
if ( data[k] > search[idx] )
{
j = k - 1;
}
else
{
break; //елемент знайдено
}
}
}
Дана програма на GPU працює в тисячі разів швидше, ніж на CPU, що підтверджує вищесказане про ефективність алгоритму на GPU у разі пошуку масиву елементів.
Можливі й інші випадки, наприклад, коли дані динамічно змінюються, або просто пошук здійснюється відносно рідко. Чудовим прикладом може служити звичайний текст. Можливість пошуку в текстовому редакторі невід'ємна його функція. В даному випадку, має сенс організувати пошук за допомогою паралельної технології, однопотоковий алгоритм на CPU змушений буде переглядати весь текст послідовно (хай і не по кожній букві, а з певним кроком), що до того ж, робить алгоритм нестійким, тому час пошуку прямо залежить від положення шуканого рядка в тексті. У разі ж алгоритму на GPU маємо більш незалежний пошук, тому безліч потоків може починати перегляд тексту з різних його ділянок.
У найкращому і середньому випадку алгоритм на CPU безумовно виграє за часом пошуку, проте в гіршому випадку, коли шуканий елемент буде в кінці списку, він програє GPU. Зрозуміло, якщо мова йде про достатньо велику кількість даних. Який з алгоритмів краще використовувати однозначно відповісти не можна. У разі CPU ми матимемо пошук, який може значно розтягуватися за часом, а у випадку з GPU матимемо достатньо постійний час пошуку.
І так, підіб'ємо підсумки. Загалом можна сказати, що GPU особливо пристосовані вирішувати проблеми, які можуть бути адаптовані до паралельних обчислень, коли одна і та ж програма (алгоритм) виконується над багатьма елементами даних паралельно, вимагають інтенсивних розрахунків: від простих арифметичних операцій до операцій над пам'яттю. Виконання однієї і тієї ж програма для багатьох елементів даних одночасно знижує вимоги до контролю потоків.
Паралельне обчислення розподіляє елементи даних між одночасно виконуваними потоками. Багато програм, які обробляють великі обсяги даних можуть використовувати модель паралельних обчислень для підвищення продуктивності. Багато алгоритмів не з області обробки зображень можуть бути поліпшені за допомогою паралельних обчислень, починаючи з обробки сигналів або фізичного моделювання систем і закінчуючи обчислювальною економікою чи біологією.
Однак прискорення алгоритмів за допомогою технології nVidia CUDA далеко не панацея і має ряд недоліків і обмежень серед яких особливості архітектури і особливості програмування. Теоретичний аналіз алгоритмів не завжди можна застосувати у разі технології nVidia CUDA, тому не враховує багатьох її особливостей. CUDA ефективна в таких алгоритмах, де CPU змушений повторювати ітерації, а CUDA може зробити те ж саме в різних потоках. Іншими словами, якщо при перенесенні алгоритму з CPU на GPU ми позбавляємося від циклу і переносимо його на потоки.
Також потрібно відзначити загальне висловлювання, про те, що CUDA ефективна в таких завданнях, де на вході маємо великий масив даних і на виході так само маємо великий масив даних.
ІІ. Реалізація алгоритму SSSP на GPU
У даному розділі буде розглянуто можливість ефективного розпаралелення алгоритму SSSP - пошуку найкоротшого шляху в графі з використанням графічних прискорювачів. В якості графічного прискорювача буде розглянута карта GTX Titan архітектури Kepler.
Останнім часом все більшу роль відіграють графічні прискорювачі (GPU) в не графічних обчисленнях. Потреба їх використання обумовлена їх відносно високою продуктивністю і більш низькою вартістю. Як відомо, на GPU добре вирішуються завдання на структурних сітках, де паралелізм так чи інакше легко виділяється. Але є завдання, які вимагають великих потужностей і використовують не структурні сітки.
Прикладом такої задачі є Single Shortest Source Path problem (SSSR) - завдання пошуку найкоротших шляхів від заданої вершини до всіх інших у зваженому графі. Для вирішення даної задачі на CPU існує, принаймні, два відомих алгоритми: алгоритм Дейсктри і алгоритм Форда-Беллмана.
Коротко розглянемо структуру зберігання неорієнтованого зваженого графа, так як надалі вона буде згадуватися і перетворюватися. Граф задається в стислому CSR форматі. Даний формат широко поширений для зберігання розріджених матриць і графів. Для графа з N вершинами і M ребрами необхідно три масиви: xadj, adjncy, weights. Масив xadj розміру N + 1, інші два - 2 * M, оскільки в неорієнтованому графі для будь-якої пари вершин необхідно зберігати пряму і зворотню дуги.
Принцип зберігання графа наступний. Весь список сусідів вершини I знаходиться в масиві adjncy з індексу xadj [I] до xadj [I + 1], не включаючи його. За аналогічним індексам зберігаються ваги кожного ребра з вершини I. Для ілюстрації на рисунку 2 зліва показаний граф з 5 вершин, записаний за допомогою матриці суміжності, а на справа - в CSR форматі.
/
Рисунок 2. Представлення графа
Реалізація алгоритму на GPU
Для того щоб збільшити обчислювальне навантаження на один потоковий мультипроцесор (SMX), необхідно перетворити вхідні дані. Усі перетворення можна розділити на два етапи:
Розширення формату CSR в координатний формат (COO).
Сортування COO формату.
На першому етапі необхідно розширити формат CSR наступним чином: введемо ще один масив startV, в якому будемо зберігати початок дуг. Тоді в масиві adjncy будуть зберігатися їх кінці. Таким чином, замість зберігання сусідів, будемо зберігати дуги. Приклад даного перетворення на графі, описаному вище приведено на рисунку 3.
/
Рисунок 3. Перетворення на графі
На другому етапі необхідно відсортувати отримані ребра так, щоб кожна пара (U, V) зустрічалася рівно один раз. Таким чином, при обробці ребра (U, V) можна відразу обробити ребро (V, U), не зчитуючи повторно дані про це ребро з глобальної пам'яті GPU.
За основу для реалізації на GPU узятий алгоритм Форда-Беллмана. Даний алгоритм добре підходить для реалізації на GPU в силу того, що ребра можна переглядати незалежно один від одного і дані ребер і їх ваг розташовані поспіль, що покращує пропускну здатність пам'яті GPU:
int k = blockIdx.x * blockDim.x + threadIdx.x;
if(k < maxV)
{
double w = weights[k];
unsigned en = endV[k];
unsigned st = startV[k];
if(dist[st] > dist[en] + w) // (*)
{
dist[st] = dist[en] + w;
modif[0] = iter;
}
else if(dist[en] > dist[st] + w) // (**)
{
dist[en] = dist[st] + w;
modif[0] = iter;
}
}
У заданому ядрі, кожна нитка обробляє два ребра (пряме і зворотне), намагаючись поліпшити відстань по одному з них. Ясно, що обидві умови if блоку не можуть виконатися одночасно. На відміну від алгоритму Форда-Беллмана, де кожне ребро проглядається послідовно, в реалізованому на GPU алгоритмі може виникнути ситуація «гонки» потоків - коли два або більше потоку заходять оновити одну і ту ж комірку dist [I]. Покажемо, що в даному випадку алгоритм залишиться коректним.
Припустимо, що є дві нитки K1 і K2, які хочуть оновити комірку dist [I]. Це означає, що виконалась умова (*) або (**). Можливі два випадки. Перший - одна з двох ниток записала найменше значення. Тоді на наступній ітерації для цих двох ниток (потоків) умова буде помилково, а значення в комірці dist [I] виявиться мінімальним. Другий - одна з двох ниток записати не мінімальне значення. Тоді на наступній ітерації для однієї з ниток умова буде істинно, а для іншого - помилково. Тим самим, результат буде однаковий в обох випадках, але досягнутий за різну кількість ітерацій.
Згідно оптимізованому варіанту алгоритми Форда-Беллмана, якщо на якій-небудь ітерації не було змін в масиві dist, то далі ітегрувати не має сенсу. На GPU для цих цілей була введена змінна modif, в яку нитки записували номер поточної ітерації.
Одна ітерація - один запуск ядра. У базовому варіанті, на CPU в циклі запускаємо ядро, далі зчитуємо змінну modif і, якщо вона не змінилася з попереднього ітерації, то в масиві dist відповідь завдання - найкоротший шлях із заданої вершини до всіх інших.
Далі розглянемо деякі оптимізації, які дозволяють істотно поліпшити продуктивність алгоритму. Знання кінцевої архітектури може бути корисним при виконанні оптимізацій.
Сучасні CPU мають дворівневий кеш. Розмір кеша першого рівня дорівнює 64Кб і міститься на всіх обчислювальних ядрах процесора. Розмір кеша другого рівня варіюється від 1 до 2МБ. Кеш третього рівня є загальним для всього CPU і має розмір близько 12-15МБ.
Сучасні GPU мають дворівневий кеш. Розмір кеша першого рівня дорівнює 64Кб. Використовується для розділення пам'яті і витіснення регістрів. Для пам'яті, що розглядається доступно не більше 48Кб. Максимальний розмір кеша другого рівня становить 1,5МБ і є загальним для всього GPU. Використовується для кешування даних, що завантажуються з глобальної пам'яті GPU.
У найсучаснішому чипі GPU GK110 є 15 обчислювальних блоків. Виходить, що на один блок припадає приблизно 48Кб кеша першого рівня і 102КБ кеша другого рівня. У порівнянні з CPU, це дуже мало, тому операції читання з глобальної пам'яті графічного процесора дорожче, ніж з оперативної пам'яті центрального процесора. Також в архітектурі Kepler з'явилася можливість прямого доступу в текстурний read-only кеш. Для цього необхідно додати в ядрі перед відповідним формальним параметром const__restrict.
У даній задачі потрібно постійно оновлювати і зчитувати масив відстаней dist. Даний масив займає досить мало місця в глобальній пам'яті GPU у порівнянні з інформацією про дуги і їх ваги. Наприклад, для графа з кількістю вершин 220 (приблизно 1 млн) масив dist займатиме 8МБ. Незважаючи на це, доступ до цього масиву здійснюється випадково, що погано для GPU, так як породжуються додаткові завантаження з глобальної пам'яті на кожен обчислювальний варп. Для того щоб мінімізувати кількість завантажень на один варп, необхідно зберегти дані в L2 кеші, прочитані, але не використані дані іншими варпами.
Так як текстурний кеш є read-only, то для того щоб ним скористатися довелося ввести дві різні посилання на один і той же масив відстаней dist. Відповідний код:
__global__ void relax_ker (const double * __restrict dist, double *dist1, … …)
{
int k = blockIdx.x * blockDim.x + threadIdx.x + minV;
if(k < maxV)
{
double w = weights[k];
unsigned en = endV[k];
unsigned st = startV[k];
if(dist[st] > dist[en] + w)
{
dist1[st] = dist[en] + w;
modif[0] = iter;
}
else if(dist[en] > dist[st] + w)
{
dist1[en] = dist[st] + w;
modif[0] = iter;
}
}
}
У підсумку вийшло так, що всередині ядра всі операції читання ведуться з одним масивом, а всі операції запису з іншим. Але обидва посилання dist і dist1 вказують на одну і ту ж ділянку пам'яті GPU.
Для кращої роботи описаної вище оптимізації необхідно, щоб якомога довше завантажені дані знаходилися в L2 кеші. Доступ до масиву dist здійснюється за заздалегідь відомими індексами, що зберігаються в масивах endV і startV. Для локалізації звернень поділимо масив dist на сегменти певної довжини, наприклад, P елементів. Так як усього в графі N вершин, то отримаємо (N / P + 1) різних сегментів. Далі, будемо перевпорядковувати ребра в дані сегменти таким чином. У першу групу віднесемо ребра, кінці яких потрапляють в нульовій сегмент, а початки - спочатку в нульовий, потім у перший і т.д. У другу групу віднесемо ребра, кінці яких потрапляють в перший сегмент, а початки - спочатку в нульовий, потім у перший і т.д.
Після даної перестановки ребер, значення елементів масиву dist, відповідні, наприклад, першій групі будуть перебувати в кеші максимально довго, так як нитки в першій групі будуть запитувати дані з нульового сегмента для кінцевих вершин і нульового, першого і т.д. для початкових вершин. Причому ребра розташовані так, що нитки будуть запитувати дані не більше ніж з трьох різних сегментів.
ІІІ. Реалізація алгоритму BFS на GPU
У даному розділі буде розглянуто можливість ефективно розпаралелити алгоритм BFS - пошук в ширину в графі з використанням графічних прискорювачів.
Останнім часом все більшу роль відіграють графічні прискорювачі (GPU) В не графічних обчисленнях. Потреба їх використання обумовлена їх відносно високою продуктивністю і більш низькою вартістю. Як відомо, на GPU добре вирішуються завдання на структурних сітках, де паралелізм так чи інакше легко виділяється. Але є завдання, які вимагають великих потужностей і використовують не структурні сітки.
Прикладом завдання на неструктурних сітках є завдання Breadth First Search (BFS) - пошуку в ширину в неорієнтованому графі. Дане завдання є основною в ряді алгоритмів на графах. Також вона трохи простіша, ніж пошук найкоротшого шляху. На даний момент алгоритм BFS використовується як основний тест для рейтингу Graph500. Далі розглянемо, як можна використовувати ідеї рішення задачі SSSP в задачі BFS. Про архітектуру GPU компанії Nvidia і про згадані алгоритмах вже багато написано, тому в цій статті я не стану додатково писати про це. Так само, сподіваюся, що поняття warp, cuda блок, SMX, та інші базові речі, пов'язані з CUDA читачеві знайомі.
Так само як і в задачі SSSR будемо використовувати ті ж самі перетворення, для того, щоб збільшити навантаження на один SMX і знизити кількість однакових даних, що знаходяться в глобальній пам'яті GPU. Основна відмінність - для алгоритму BFS не потрібні ваги в графі. Так само варто відзначити, що нам необхідно зберігати не найкоротші відстані, а номер рівня, в якому знаходиться дана вершина (див. рис. 4).
/
Рисунок 4. Рівні графа
Після проведення тестових запусків з'ясувалося, що кількість рівнів в RMAT графах із середнім ступенем зв'язності 32 не перевищує 10. Тому для зберігання даних значень буде достатньо unsigned char. Тим самим масив рівнів буде займати в 8 разів менше місця, ніж масив відстаней, що дуже важливо, так як розмір кеша в архітектурі Kepler всього 1,5МБ.
На CPU був реалізований нативний варіант алгоритму обходу в ширину, а саме створення черги ще не переглянутих вершин. Код CPU реалізації виглядає наступним чином:
queue<vertex_id_t> q;
bool *used = new bool[G->n];
for (unsigned i = 0; i < G->n; ++i)
used[i] = false;
used[root] = true;
q.push(root);
dist[root] = 0;
while (!q.empty())
{
vertex_id_t nextV = q.front();
q.pop();
for (unsigned k = G->rowsIndices[nextV]; k < G->rowsIndices[nextV + 1]; ++k)
{
if (used[G->endV[k]] == false)
{
used[G->endV[k]] = true;
q.push(G->endV[k]);
dist[G->endV[k]] = dist[nextV] + 1;
}
}
}
Даний код досить простий і, цілком можливо, не оптимальний. Він був використаний для перевірки правильності роботи алгоритму на GPU. Мети написати оптимальний алгоритм на CPU не було, тому продуктивність на CPU буде отримана даним алгоритмом.
За основу реалізації на GPU був узятий все той же алгоритм Форда-Беллмана і ядро в розглянутому завданню SSSP. Основне обчислювальне ядро для BFS буде виглядати наступним чином:
if(k < maxV)
{
unsigned en = endV[k];
unsigned st = startV[k];
if(levels[st] == iter)
{
if(levels[en] > iter)
{
levels_NR[en] = iter + 1;
modif[0] = iter;
}
}
else if(levels[en] == iter)
{
if(levels[st] > iter)
{
levels_NR[st] = iter + 1;
modif[0] = iter;
}
}
}
Ідея алгоритмe в наступному. Нехай масив levels спочатку заповнений максимальним для обраного типу значенням (255 для unsigned char). Будемо передавати в ядро номер поточної ітерації - iter. Далі необхідно переглянути всі ребра і перевірити, чи є початкова або кінцева вершина поточним батьком, тобто чи належить вона рівню iter. Якщо так, то необхідно помітити протилежну вершину дуги значенням на одиницю більше, щоб «включити» цю вершину в список батьків на наступній ітерації. Також як і в SSSP, залишається змінна modif, що показує необхідність продовження розмітки графа.
Даний код вже містить застосування оптимізації в задачі SSSR - використання const __restrict для масиву levels і використання іншого посилання levels_NR, що вказує на ту саму ділянку пам'яті, необхідну для запису. Друга оптимізація у вигляді перестановки для кращої локалізації даних в кеші теж була використана. Для алгоритму BFS довжина оптимальної кеш лінії дорівнює 1024КБ або приблизно 1млн елементів масиву levels (1024 * 1024) незалежно від розміру графа.
СПИСОК ВИКОРИСТАНОЇ ЛІТЕРАТУРИ
Кнут Д.Э. Искусство программирования, том 3. Сортировка и поиск, 2-е изд. М.: Вильямс , 2007. С. 824. ISBN 0-201-89685-0
Parallel search on video cards, article, oracle server technologies special projects.
Офіційний сайт компанії NVidia <http://www.nvidia.ru/page/home.html>
Заметки на полях. Что быстрее, CPU или GPU? <http://www.3dnews.ru/video/what_is_faster_gpu_or_cpu>
Что такое вычисления с gpu-ускорением? <http://www.nvidia.com.ua/object/gpu-computing-ru.html>
Параллельные вычисления с cuda. <http://www.nvidia.com.ua/object/cuda-parallel-computing-ru.html>
CPU vs GPU. Доступний з: <http://www.fssr.ru/hz.php?name=News&file=article&sid=9504>
08.12.2015 09:12-
Ділись своїми роботами та отримуй миттєві бонуси!
0%
Оголошення від адміністратора