OpenCL: добавление поддержки видеокарт AMD в CMakeLists.txt

 Вот этот пример мы модифицируем так, чтобы он мог компилироваться для AMD видеокарт тоже. CMakeLists.txt получился вот такой:

1. #
2. cmake_minimum_required (VERSION 3.8)
4. project ("MemakePrj")
6. # Add source to this project's executable.
7. add_executable (MemakePrj "main.cpp" "Memake/Memake.cpp" "Memake/Vector2d.cpp")
9. # SDL2 headers
10. target_include_directories(MemakePrj PRIVATE "SDL2-2.0.14/include")
12. # gpu vendor: amd, nvidia
13. set(gpu_vendor "amd")
14. #
15. if("${gpu_vendor}" STREQUAL "amd")
16. message("gpu_vendor is 'amd'")
17. # set define
18. target_compile_definitions(MemakePrj PUBLIC GPU_VENDOR_IS_AMD)
19. # OpenCL library
20. set(opencl_lib_folder "$ENV{OCL_ROOT}/lib")
21. # OpenCL headers
22. target_include_directories(MemakePrj PRIVATE "$ENV{OCL_ROOT}/include")
23. elseif("${gpu_vendor}" STREQUAL "nvidia")
24. message("gpu_vendor is 'nvidia'")
25. # set define
26. target_compile_definitions(MemakePrj PUBLIC GPU_VENDOR_IS_NVIDIA)
27. # OpencCL headers
28. target_include_directories(MemakePrj PRIVATE "$ENV{CUDA_PATH}/include")
29. # OpenCL library
30. set(opencl_lib_folder "$ENV{CUDA_PATH}/lib")
31. else()
32. message("gpu_vendor is 'UNKNOWN'")
33. endif()
34. 
    
35. # add SDL_MAIN_HANDLED definition to avoid
36. # "LNK2019 unresolved external symbol SDL_main referenced in function main_getcmdline"
37. add_definitions( -DSDL_MAIN_HANDLED )
39. # SDL library folder
40. set(SDL2_lib_folder "${PROJECT_SOURCE_DIR}/SDL2-2.0.14/lib")
42. message(${CMAKE_BUILD_TYPE})
44. # check 32 or 64 bits
45. if(CMAKE_SIZEOF_VOID_P EQUAL 8)
46. # 64 bits
47. set(SDL2_lib_folder "${SDL2_lib_folder}/x64")
48. if("${gpu_vendor}" STREQUAL "amd")
49. set(opencl_lib_folder "${opencl_lib_folder}")
50. elseif("${gpu_vendor}" STREQUAL "nvidia")
51. set(opencl_lib_folder "${opencl_lib_folder}/x64")
52. else()
53. message("gpu_vendor is 'UNKNOWN'")
54. endif()
55. elseif(CMAKE_SIZEOF_VOID_P EQUAL 4)
56. # 32 bits
57. set(SDL2_lib_folder "${SDL2_lib_folder}/x86")
58. if("${gpu_vendor}" STREQUAL "amd")
59. message("do not want to support 32-bit version for amd")
60. elseif("${gpu_vendor}" STREQUAL "nvidia")
61. set(opencl_lib_folder "${opencl_lib_folder}/Win32")
62. else()
63. message("gpu_vendor is 'UNKNOWN'")
64. endif()
65. endif()
67. # link SDL2 static lib
68. target_link_libraries(MemakePrj ${SDL2_lib_folder}/SDL2.lib)
69. target_link_libraries(MemakePrj ${SDL2_lib_folder}/SDL2main.lib)
71. # link OpenCL library
72. target_link_libraries(MemakePrj ${opencl_lib_folder}/OpenCL.lib)
74. # copy dynamic lib to folder with executable file
75. file(COPY ${SDL2_lib_folder}/SDL2.dll DESTINATION ${PROJECT_BINARY_DIR})

 Получилось гораздо более громоздко и неудобочитаемо, чем было. В строке 12 мы указываем, какой у нас производитель видеокарты (пока что их два: AMD, Nvidia) IF внутри IF - это не то, что способствует облегчению чтения подобного кода, но это, тем не менее, позволяет собрать проект для двух и более производителей видеокарт, а это главное. Позже подумаю, как правильно разбить CMakeLists.txt на части.

В коде же это выглядит так:

Настройка OpenCL для видеокарт AMD

Прежде, чем начать, отмечу, что компания AMD на данный момент не поддерживает свой OpenCL SDK, к сожалению. Причины мне непонятны, хотя они предлагают вместо него какой-то свой собственный аналог CUDA у NVidia. Последним, к слову, наличие CUDA никак не помешало поддерживать тот же OpenCL. Такое отношение к OpenCL не может радовать, конечно же. Тот набор инструментов, что используется ниже - он поддерживается энтузиастами (за что им огромное спасибо). 

Скачать SDK можно, например, тут. После установки, установим переменную окружения OCL_ROOT и проверим, что она правильно установлена: 

Далее все делаем так же, как мы делали раньше и в тех же местах.

Тут настройка среды закончена и для запуска примера осталось только чуть подредактировать код, вместо заголовка cl.hpp мы возьмем cl2.hpp и добавим определение CL_HPP_ENABLE_PROGRAM_CONSTRUCTION_FROM_ARRAY_COMPATIBILITY перед заголовочным файлом для совместимости со старым кодом:

После этого можно запустить пример и убедиться, что он работает:

Заодно проверим как работает перемножение матриц в наивной реализации для CPU и GPU:

Запускалалось на Ryzen 7 7730U with Vega 8 iGPU

Заодно сравним производительность ноутбучного процессора и встроенной видеокарты с результатами настольного ПК:

Более реалистичный 2д коллайдер / More realistic 2d collider

 В предыдущих двух постах (раз, два) я реализовал 2д-коллайдер, который применяет правило "угол падения равен углу отражения", которое визуально достаточно неплохо работает (особенно когда много частиц), но все же слишком упрощенно показывает отражения после столкновений. Например, частицы всегда имеют одну и ту же скорость и меняется только направление, а хотелось бы, чтобы частицы все же обменивались импульсами более реалистично. Поэтому я решил усложнить задачу и сделать коллайдер, который реализует более сложный закон, а именно закон абсолютно упругого столкновения и закон сохранения импульса. Все формулы и пояснения можно найти здесь. Вот ключевые моменты:

И вот эти формулы:

Собственно, это вся теория, которая нам нужна. Кто-то может заметить, что все эти иллюстрации и формулы работают для одномерного пространства, но нам же нужен двумерный случай! В википедии специально для этого указано следующее:

В случае столкновения двух тел в трёхмерном пространстве векторы импульсов тел до и после столкновения лежат в одной плоскости. Вектор скорости каждого тела может быть разложен на две компоненты: одна по общей нормали поверхности сталкивающихся тел в точке контакта, а другая параллельная поверхности столкновения. Поскольку сила удара действует только по линии столкновения, компоненты скорости, векторы которых проходят по касательной к точке столкновения, не изменяются. Скорости, направленные вдоль линии столкновения, могут быть вычислены с помощью тех же уравнений, что и столкновения в одном измерении. Окончательные скорости могут быть вычислены из двух новых компонентов скоростей и будут зависеть от точки столкновения.

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

В коде это выглядит вот так:

1. void pulseColl(const Ball& b2)
2. {
3. // direction to other ball
4. Point2f toB2 = b2.pos - pos;
5. Point2f toB2Unit = toB2.unitVector();
6. // calculate dot product
7. float dotProdB = f.dotProduct(toB2Unit);
8. // calculate speed v for both balls
9. // that is projection to hit axis
10. float v1 = dotProdB;
11. float dotProdB2 = b2.f.dotProduct(toB2Unit);
12. float v2 = dotProdB2;
13. // if move projection to hit axis
14. // for both balls are directed to move away
15. // one from other then do nothing
16. if (v1 <= 0 && v2 >= 0)
17. {
18. return;
19. }
20. // mass is equal to square of 2d ball
21. float m1 = r * r;
22. float m2 = b2.r * b2.r;
23. // speed of this ball after collision
24. float v1new = (2 * m2 * v2 + v1 * (m1 - m2)) / (m1 + m2);
25. // "to" move component
26. Point2f tmpTo = toB2Unit * v1new;
27. // "tangent" move component
28. Point2f tanUnit = Mat2x2f().rot(k_PI / 2) * toB2Unit;
29. float dotProdTan1 = f.dotProduct(tanUnit);
30. Point2f tmpTan = tanUnit * dotProdTan1;
31. // new move vector
32. Point2f newF = tmpTo + tmpTan;
33. f = newF;
34. }

Сначала мы вычислем единичный вектор вдоль оси столкновения (строки 4-5), после чего вычисляем проекцию на этот вектор нашего вектора движения (строка 7). В строке 11 мы вычисляем такую же проекцию для второго круга. В строке 16 мы проверяем, что оба обьекта двигаются не друг от друга (а если двигаются друг от друга, то ничего не делаем), Такая проверка нужна, т.к. без нее, как экспериментально выяснилось, обьекты могут "сцепляться" и кружиться в такой сцепке друг вокруг друга. Далее, в строке 24 мы производим собственно вычисления нового значения скорости вдоль оси столкновения, после чего в получаем первую компоненту (осевую или радиальную) нового значения скорости. В строках 29-30 мы получаем вторую компоненту скорости (тангенциальную) и в строке 32 вычисляем новый вектор движения путем сложения осевой и тангенциальной компонент. "Осевая" компонента потому, что вдоль оси столкновения. 

Код ядра выглядит так:

1. Ball pulseColl(Ball b1, const Ball b2)
2. {
3. // direction to other ball
4. float2 p1 = (float2)(b1.pos.x, b1.pos.y);
5. float2 p2 = (float2)(b2.pos.x, b2.pos.y);
6. float2 toB2 = p2 - p1;
7. float2 toB2Unit = getUnitVector(toB2);
8. // calculate dot product
9. float2 f1 = (float2)(b1.f.x, b1.f.y);
10. float dotProdB1 = getDotProduct(f1, toB2Unit);
11. // calculate speed v for both balls
12. // that is projection to hit axis
13. float v1 = dotProdB1;
14. float2 f2 = (float2)(b2.f.x, b2.f.y);
15. float dotProdB2 = getDotProduct(f2, toB2Unit);
16. float v2 = dotProdB2;
17. // if move projection to hit axis
18. // for both balls are directed to move away
19. // one from other then do nothing
20. if (v1 <= 0 && v2 >= 0)
21. {
22. return b1;
23. }
24. // mass is equal to square of 2d ball
25. float m1 = b1.r * b1.r;
26. float m2 = b2.r * b2.r;
27. // speed of this ball after collision
28. float v1new = (2 * m2 * v2 + v1 * (m1 - m2)) / (m1 + m2);
29. // "to" move component
30. float2 tmpTo = toB2Unit * v1new;
31. // "tangent" move component
32. float2 tanUnit = rotVect(toB2Unit, M_PI_2_F);
33. float dotProdTan1 = getDotProduct(f1, tanUnit);
34. float2 tmpTan = tanUnit * dotProdTan1;
35. // new move vector
36. float2 newF = tmpTo + tmpTan;
37. b1.f.x = newF[0];
38. b1.f.y = newF[1];
39. return b1;
40. }
41.  

Тут все более-менее то же самое. 

Исходный код находится здесь.

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

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