Редагувати 2: Будь ласка, подивіться на це цей поперечний стовп для TLDR.
Редагувати: Враховуючи, що частинки сегментовані на клітини сітки (скажімо 16^3 grid), чи є кращою ідеєю запустити одну робочу групу для кожної комірки сітки та стільки робочих елементів в одній робочій групі, скільки може бути максимальна кількість частинок на комірку сітки?
У цьому випадку я міг завантажити всі частинки зсусідні клітини в локальну пам’ять і перебирають через них обчислюючи деякі властивості. Тоді я міг би записати конкретне значення в кожну частинку поточної комірки сітки.
Чи був би такий підхід вигідним при запуску ядра для всіх частинок і для кожної ітерації над (більшу частину однакових) сусідів?
Крім того, яке є ідеальне співвідношення number of particles/number of grid cells?
Я намагаюся здійснити (та змінити) Частинки CUDA для OpenCL і використовувати його для запиту найближчих сусідів для кожної частки. Я створив такі структури:
- Буфер
Pутримуючи всі частинки "3D положеннями (float3) Буфер
Spзберіганняint2пари ідентифікаторів частинок та їх просторові хеші.Spсортується за хешем. (Хеш - це просто просте лінійне відображення від 3D до 1D - ще немає Z-індексації.)Буфер
Lзберіганняint2пари початкового та кінцевого положень певних просторових хешів у буферіSp. Приклад:L[12] = (int2)(0, 50).L[12].x- індекс (вSp) з перший частинка з просторовим хешем12.L[12].y- індекс (вSp) з останній частинка з просторовим хешем12.
Тепер, коли у мене є всі ці буфери, я хочу пройти всі частинки всередині P і для кожної частинки ітерація через найближчих сусідів. На даний момент у мене є ядро, яке виглядає так (псевдокод):
__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
size_t gid = get_global_id(0);
float3 curr_particle = P[gid];
int processed_value = 0;
for(int x=-1; x<=1; x++)
for(int y=-1; y<=1; y++)
for(int z=-1; z<=1; z++) {
float3 neigh_position = curr_particle + (float3)(x,y,z)*GRID_CELL_SIDE;
// ugly boundary checking
if ( dot(neigh_position<0, (float3)(1)) +
dot(neigh_position>BOUNDARY, (float3)(1)) != 0)
continue;
int neigh_hash = spatial_hash( neigh_position );
int2 particles_range = L[ neigh_hash ];
for(int p=particles_range.x; p<particles_range.y; p++)
processed_value += heavy_computation( P[ Sp[p].y ] );
}
Out[gid] = processed_value;
}
Проблема цього коду полягає в тому, що він повільний. Я підозрюю, що нелінійний доступ до пам'яті графічного процесора (зокрема P[Sp[p].y] у внутрішньому-самому for петля), що спричиняє повільність.
Що я хочу зробити, це використовувати Крива Z-порядку як просторовий хеш. Таким чином я міг мати лише 1 for ітерація циклу через безперервний діапазон пам'яті при запиті сусідів. Єдина проблема полягає в тому, що я не знаю, якими повинні бути значення початку та зупинки Z-індексу.
Святий Грааль, якого я хочу досягти:
__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
size_t gid = get_global_id(0);
float3 curr_particle = P[gid];
int processed_value = 0;
// How to accomplish this??
// `get_neighbors_range()` returns start and end Z-index values
// representing the start and end near neighbors cells range
int2 nearest_neighboring_cells_range = get_neighbors_range(curr_particle);
int first_particle_id = L[ nearest_neighboring_cells_range.x ].x;
int last_particle_id = L[ nearest_neighboring_cells_range.y ].y;
for(int p=first_particle_id; p<=last_particle_id; p++) {
processed_value += heavy_computation( P[ Sp[p].y ] );
}
Out[gid] = processed_value;
}
Відповіді:
-1 для відповіді № 1Вам слід уважно вивчити алгоритми Мортона. Ерісонс в реальному часі виявлення зіткнень пояснює це дуже добре.
Еріксон - Виявлення зіткнень у реальному часі
Ось ще одне приємне пояснення, включаючи деякі тести:
Кодування / декодування Мортона за допомогою бітового чергування: реалізації
Алгоритми Z-Order визначають лише шляхи докоординати, в яких ви можете хешувати від 2 або 3D координат до цілого числа. Хоча алгоритм заглиблюється на кожну ітерацію, ви повинні встановлювати межі самостійно. Зазвичай індекс зупинки позначають сторожовим. Якщо дозволити сторожовій зупинці, ви дізнаєтесь, на якому рівні розміщена частинка. Тож максимальний рівень, який ви хочете визначити, покаже вам кількість клітинок на вимір. Наприклад, при максимальному рівні 6 ви маєте 2 ^ 6 = 64. У вашій системі (3D) буде 64x64x64 комірок. Це також означає, що ви повинні використовувати цілочисельні координати. Якщо ви використовуєте floats, вам доведеться конвертувати like coord.x = 64*float_x і так далі.
Якщо ви знаєте, скільки комірок у вашій системі, ви можете визначити свої обмеження. Ви намагаєтесь використовувати двійковий октрі?
Оскільки частинки перебувають у русі (у цьому прикладі CUDA), вам слід спробувати розпаралелювати кількість частинок замість клітин.
Якщо ви хочете скласти списки найближчих сусідівви повинні зіставити частинки з клітинами. Це робиться за допомогою таблиці, яка сортується згодом по клітинах до частинок. Тим не менш, вам слід перебирати частинки та отримувати доступ до сусідів.
Про ваш код:
Проблема цього коду полягає в тому, що він повільний. Я підозрюю, що нелінійний доступ до пам'яті графічного процесора (зокрема P [Sp [p] .y] у внутрішньому циклі for) спричиняє повільність.
Згадайте Дональда Ннута. Ви повинні виміряти, де знаходиться горлечко пляшки. Ви можете використовувати NVCC Profiler і шукати вузьке місце. Не впевнений, що OpenCL має як профайлер.
// ugly boundary checking if ( dot(neigh_position<0, (float3)(1)) + dot(neigh_position>BOUNDARY, (float3)(1)) != 0) continue;
Я думаю, вам не слід розгалужувати його таким чином, як щодо повернення нуля при дзвінку heavy_computation. Не впевнений, але, можливо, у вас тут є своєрідне передбачення галузі. Спробуйте це якось видалити.
Паралельний запуск по клітинках - хороша ідеялише якщо у вас немає доступу до даних частинок, інакше доведеться використовувати атоми. Якщо ви переходите діапазон часток, ви натомість читаєте доступ до комірок та сусідів, але паралельно створюєте свою суму, і вас не змушують до якоїсь парадигми умовності раси.
Крім того, яке ідеальне співвідношення кількості частинок / кількості клітин сітки?
Дійсно залежить від ваших алгоритмів таупаковки частинок у вашому домені, але у вашому випадку я б визначив розмір комірки, еквівалентний діаметру частинок, і просто використав би кількість отриманих клітинок.
Отже, якщо ви хочете скористатися Z-порядком і досягти свого святого Грааля, спробуйте використовувати цілочисельні координати та хеш їх.
Також намагайтеся використовувати більшу кількість частинок.Близько 65000 частинок, таких як приклади CUDA, слід врахувати, оскільки таким чином розпаралелювання є в основному ефективним; запущені блоки обробки використовуються (менше простоїв потоків).