2d-коллайдер и его реализация на OpenCL

Улучшаем предыдущий пример. Теперь реализуем этот 2д-коллайдер на OpenCL. Учитывая, что нам надо будет параллельно вычислять и менять текущую позицию и вектор движения одновременно для всех обьектов, то использовать один массив обьектов будет проблематично, т.к. нужно будет использовать различные примитивы синхронизации, что безусловно пагубно скажется на производительности. Поэтому обход/модификация всех обьектов будет производиться следующим образом: один массив обьектов будет только для чтения, а второй только для записи и каждый тред на GPU будет записывать изменения в свою ячейку, не пересекаясь с другими.

1. __kernel void collideAndUpdate(
2. __global Ball* b1,
3. __global Ball* b2,
4. const int ballCnt,
5. const int scrW,
6. const int scrH,
7. const double frameTimeMs)
8. {
9. // Get work-item identifiers.
10. int i = get_global_id(0);
11. Ball tmp = b1[i];
12. for (int j = 0; j < ballCnt; j++)
13. {
14. // check collision between one ball to others,
15. // but don't check collision to itself
16. if (j != i)
17. {
18. Ball tmp2 = b1[j];
19. tmp = checkCollision(tmp, tmp2);
20. }
21. }
22.  
23. // check borders
24. if (tmp.pos.x <= 0 || tmp.pos.x >= scrW)
25. {
26. tmp.f.x *= -1;
27. }
28. if (tmp.pos.y <= 0 || tmp.pos.y >= scrH)
29. {
30. tmp.f.y *= -1;
31. }
32.  
33. // update positions
34. tmp.pos.x += tmp.f.x * frameTimeMs;
35. tmp.pos.y += tmp.f.y * frameTimeMs;
36.  
37. b2[i] = tmp;
38. }

Поскольку не получилось подключить C++ заголовочный файл в код ядра, то пришлось по сути переписывать все заново для OpenCL:

1. Ball checkCollision(Ball b1, const Ball b2)
2. {
3. float2 p1 = (float2)(b1.pos.x, b1.pos.y);
4. float2 p2 = (float2)(b2.pos.x, b2.pos.y);
5. const float dist = getDistanceBetween(p1, p2);
6. if (dist < b1.r + b2.r)
7. {
8. // direction to other ball
9. float2 to2 = p2 - p1;
10. float2 f = (float2)(b1.f.x, b1.f.y);
11. // calculate dot product
12. float dotProd = dot(f, to2);
13. // if dot product is negative then force directed away from B ball
14. // and we do nothing
15. if (dotProd > 0)
16. {
17. // angle between normal and force (moving) vectors
18. float angle = getAngleTo(f, to2);
19. // the angle of incidence is equal to the angle of reflection
20. f = rotVect(f, angle * 2);
21. f = f * (-1);
22. b1.f.x = f.x;
23. b1.f.y = f.y;
24. }
25. }
26. return b1;
27. }

И структуру Ball и все векторные операции тоже пришлось переписывать. Потому что OpenCL компилятор и С++ компилятор отличаются очень сильно. Фактически настолько, что и там и там можно использовать только какие-то простые структуры и константы с дефайнами. Вот как выглядят векторные операции:

1. float getDistanceBetween(float2 p1, float2 p2)
2. {
3. float2 p1p2 = p2 - p1;
4. float dist = native_sqrt(p1p2[0] * p1p2[0] + p1p2[1] * p1p2[1]);
5. return dist;
6. }
7.  
8. float getLen(float2 p)
9. {
10. float len = getDistanceBetween((float2)(0, 0), p);
11. return len;
12. }
13.  
14. float getCrossProd(float2 p1, float2 p2)
15. {
16. float crossProd = p1[0] * p2[1] - p1[1] * p2[0];
17. return crossProd;
18. }
19.  
20. float getAngleTo(float2 p1, float2 p2)
21. {
22. return asin(getCrossProd(p1, p2) / (getLen(p1) * getLen(p2)));
23. }
24.  
25. float2 rotVect(float2 v, float angle)
26. {
27. // first of all create rotation matrix
28. float c = cos(angle);
29. float s = sin(angle);
30. float2 mr0 = { c, -s }; // first row
31. float2 mr1 = { s, c }; // second row
32. // get rotated vector by multiply matrix to vector
33. float2 tmp;
34. tmp[0] = (v[0] * mr0[0] + v[1] * mr0[1]);
35. tmp[1] = (v[0] * mr1[0] + v[1] * mr1[1]);
36. return tmp;
37. }

Я не нашел как работать с матрицами, поэтому в качестве матрицы 2х2 я использовал просто два вектора float2, каждый из которых играет роль строки в матрице (строки 30-31). То есть довольно много заново написанного кода и если кто-то задумал перенести что-то с помощью OpenCL на видеокарту, то пусть имеют ввиду, что такой даже не копипасты, а полной переработки кода с учетом кучи нюансов будет очень много. 

Чтобы увидеть реальную разницу в производительности, понадобилось увеличить количество шариков с 500 до 20000 и уменьшить изх диаметр до 1, чтобы они все поместились. Результат:

Производительность на GPU при 20 тысячах шариков получилась 29.3 кадра в секунду, а на CPU всего лишь 1.6 кадра в секунду. Разница, как говорится, налицо! 

Тестовая платформа: Ryzen 3700X, 16GB RAM, RTX 3060 12GB

Весь код здесь.