Каков наилучший способ программного выбора лучшего графического процессора в OpenCL?

Question

На моем ноутбуке у меня две графические карты - Intel Iris и Nvidia GeForce GT 750M. Я пытаюсь сделать простой вектор, используя OpenCL. Я знаю, что карта Nvidia работает намного быстрее и может улучшить работу. В принципе, я могу поставить оператор if в код, который будет искать NVIDIA в атрибуте VENDOR. Но я хотел бы иметь что-то элегантное. Каков наилучший способ выбора лучшего (более быстрого) GPU программно в OpenCL C/C++?

Answer 1

Если это просто векторный добавочный файл, и ваше приложение находится на стороне хоста, то процессор победит. Или даже лучше, интегрированный процессор будет намного быстрее. Общая производительность зависит от algortihms, типов буфера opencl (use_host_ptr, read_write и т. Д.) И вычислить отношение данных. Даже если вы не скопируете, а закрепите массив и получите доступ, время ожидания процессора будет меньше, чем pci-e.

Если вы собираетесь использовать opengl + opencl interop, вам нужно будет знать, является ли ваше вычислительное устройство одним и тем же устройством с устройством вывода рендеринга. (если ваш экран получает свои данные из igpu, то это радужная оболочка, если нет, то это nvidia)

Если вам просто нужно сделать некоторые операции с массивами C++ (хост-сторона) и получить результаты с самым быстрым способом, я предлагаю вам «балансировка нагрузки».

Пример вектор-сложения элементов из 4k на ядре i7-5775C с Iris про и два gt750m (один разогнали на 10%)

Во-первых, дать равное число ndrange бушует ко всем устройствам. В конце каждой фазы расчета проверьте тайминги.

CPU  iGPU  dGPU-1  dGPU-2 oc 
Intel Intel  Nvidia  Nvidia 
1024  1024  1024   1024 
34 ms 5ms   10ms   9ms  

затем рассчитать взвешенным (зависит от последнего диапазона ndrange), но расслабленный (не точно, но близко) приближения расчета ширины полосы и изменения ndrange диапазонов соответственно:

Intel Intel  Nvidia  Nvidia 
512  1536  1024   1024 
16 ms 8ms   10ms   9ms  

затем продолжить вычисления, пока она действительно не станет стабильной ,

Intel Intel  Nvidia  Nvidia 
256  1792  1024   1024 
9ms  10ms   10ms   9ms 

или пока не сможете включить мелкие зерна.

Intel Intel  Nvidia  Nvidia 
320  1728  1024   1024 
10ms  10ms  10ms   9ms 

Intel Intel  Nvidia  Nvidia 
320  1728  960   1088 
10ms  10ms  10ms   10ms 

     ^  ^
     |   | 
     |   PCI-E bandwidth not more than 16 GB/s per device 
     closer to RAM, better bandwidth (20-40 GB/s) and less kernel overhead 

Вместо того, чтобы только последней итерации для балансировки, вы можете получить среднее (или PID) из последних 10 результатов, чтобы устранить спайки, что ввести в заблуждение балансировки. Кроме того, копии буфера могут занимать больше времени, чем вычисление, если вы включите это в балансировку, вы можете отключить ненужные/непригодные для использования устройства.

Если вы создадите библиотеку, вам не придется проверять результаты для каждого вашего нового проекта. Они будут автоматически сбалансированы между устройствами при ускорении матричных умножений, движений жидкости, соединений в sql table и финансовых приближений.

Для решения балансировки:

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

Althogh вектор-надстройка только ядро ​​не реальный мир сценарий, вот реальный тест из моей системы:

__kernel void bench(__global float * a, __global float *b, __global float *c) 
       { 
        int i=get_global_id(0); 
        c[i]=a[i]+b[i]; 
       } 
2560 768 768 
AMD FX(tm)-8150 Eight-Core Processor   Oland Pitcairn 

это после нескольких итераций (FX быстрее даже при дополнительных копиях буфера, не используя хост-указатель). Даже oland gpu ловит питчер, потому что их пропускная способность pci-e одинакова.

Теперь с некоторыми тригонометрическими функциями:

__kernel void bench(__global float * a, __global float *b, __global float *c) 
    { 
     int i=get_global_id(0); 
     c[i]=sin(a[i])+cos(b[i])+sin(cos((float)i)); 
    } 

    1792 1024 1280 

тестирование gddr3-128bit против gddr5-256bit (разогнанный) и кэширование.

__kernel void bench(__global float * a, __global float *b, __global float *c) 
{ 
        int i=get_global_id(0); 

        c[i]=a[i]+b[i]-a[i]-b[i]+a[i]+b[i]-b[i]-a[i]+b[i]; 
        for(int j=0;j<12000;j++) 
         c[i]+=a[i]+b[i]-a[i]-b[i]+a[i]+b[i]-b[i]-a[i]+b[i]; 

} 



256 256 3584 

Высокая вычислительная соотношение данных:

__kernel void bench(__global float * a, __global float *b, __global float *c) 
      { 
       int i=get_global_id(0); 

       c[i]=0.0f; float c0=c[i];float a0=a[i];float b0=b[i]; 
       for(int j=0;j<12000;j++) 
        c0+=sin(a0)+cos(b0*a0)+cos(sin(b0)*19.95f); 
       c[i]=c0; 

      } 

256 2048 1792 

Теперь Oland видеокарте снова достойно и выиграл даже только с 320 ядрами. Поскольку элементы 4k легко обернуты вокруг всех 320 ядер более чем в 10 раз, но pitcairn gpu (1280 ядер) недостаточно заполнена сложенными массивами (волновыми фронтами), что привело к более низкому заполнению исполнительных блоков ---> не могло скрывать задержки. Низкоуровневые устройства для низких нагрузок лучше, я думаю. Возможно, я мог бы использовать это, когда directx-12 выйдет с некоторым loadbalancer, и этот Oland может вычислить физику из 5000 - 10000 частиц из игровых взрывов, в то время как pitcairn может вычислять плотности дыма.

Answer 2

Я разработал трассировщик лучей в реальном времени (не только лучевой заклинатель), который программным образом выбрал два графических процессора и процессор, а также сбалансировал нагрузку на все три в реальном времени. Вот как я это сделал.

Предположим, что имеется три устройства, d1, d2 и d3. Назначьте каждое устройство весом: w1, w2 и w3. Вызвать количество пикселов, которые будут отображаться n. Предположим, что бесплатный параметр называется alpha.

1. Назначьте каждое устройство весом 1/3.
2. alpha = 0.5.
3. Рендер первые n1=w1*n пикселей на d1, следующие n2=w2*n пикселей на d2, а последние n3=w3*n пикселей на d3 и записать время, чтобы сделать для каждого deivce t1, t2 и t3.
4. Рассчитать стоимость vsum = n1/t1 + n2/t2 + n3/t3.
5. Сосчитать весы w_i = alpha*w_i + (1-alpha)*n_i/t_i/vsum.
6. Вернитесь к шагу 3.

Точка значения alpha является обеспечение плавного перехода. Вместо того, чтобы переназначить весь вес, в зависимости от времени, когда он смешивается с некоторым старым весом. Без использования alpha у меня были нестабильности. Можно настроить значение alpha. На практике он может быть установлен примерно на 1%, но не на 0%.

---

Давайте выберем пример.

У меня был GTX 590, который был двойной графикой с двумя GTX580 с тактовой частотой. У меня также был процессор Sandy Bridge 2600K. Графические процессоры были намного быстрее, чем процессор. Предположим, что они были примерно в 10 раз быстрее. Давайте также скажем, что было 900 пикселей.

Перенесите первые 300 пикселей на GPU1, следующие 300 пикселей с GPU2 и последние 300 пикселей с CPU1 и запишите время 10 s, 10 s, and 100 s соответственно. Таким образом, один графический процессор для всего изображения займет 30 с, а только процессор займет 300 с. Оба GPUS вместе взяли бы 15 s.

Рассчитать vsum = 30 + 30 + 3 = 63. Снова пересчитайте вес: w1,w2 = 0.5*(1/3) + 0.5*300/10/63 = 0.4 и w3 = 0.5*(1/3) + 0.5*300/100/63 = 0.2.

Render the next frame: 360 пикселей с GPU1, 360 PIXELS с GPU2 и 180 PIXELS с CPU1, а времена станут немного более сбалансированными, говорят 11 s, 11 s, and 55 s.

После того, как число кадров будет доминировать до тех пор, пока в конечном итоге весы не будут основаны на этом термине. В этом случае веса становятся 47% (427 пикселей), 47%, 6% (46 пикселей) соответственно, и время становится, скажем, 14 s, 14 s, 14 s соответственно. В этом случае процессор только улучшает результат использования только графических процессоров на одну секунду.

Я предположил, что в этом рассчитывается равномерная нагрузка. В реальном трассировщике лучей нагрузка зависит от линии сканирования и пикселя, но алгоритм остается неизменным для определения весов.

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

Этот метод легко распространить на несколько устройств с петлей. Однажды я протестировал свой трассировщик лучей на четырех устройствах. Два 12-ядерных процессора Xeon и два графических процессора. В этом случае процессоры имели гораздо большее влияние, но графические процессоры все еще доминировали.

---

В случае, если кому-то интересно. Я создал контекст для каждого устройства и использовал каждый контекст в отдельном потоке (используя pthreads). Для трех устройств я использовал три потока.

Фактически вы можете использовать его для работы на одном устройстве от разных поставщиков. Например, я использовал оба процессора AMD и Intel одновременно (каждый из которых генерировал около половины кадра) на моем 2600K, чтобы увидеть, какой поставщик был лучше. Когда я впервые это сделал (2012), если я правильно помню, AMD по иронии судьбы поиграла Intel на процессоре Intel.

---

В случае, если кто заинтересован в том, как я придумал формулу для весов я использовал идею из физики (мой фон физика не программирование).

Скорость (v) = расстояние/время. В этом случае расстояние (d) - это количество пикселов для обработки.Затем общее расстояние

d = v1*t1 + v2*t2 + v3*t3 

, и мы хотим, чтобы они каждый финиш в то же время так

d = (v1 + v2 + v3)*t 

затем, чтобы получить вес определяют

v_i*t = w_i*d 

который дает

w_i = v_i*t/d 

и замена (t/d) из (d = (v1 + v2 + v3)*t) дает:

w_i = v_i /(v1 + v2 + v3) 

Легко видеть, это может быть обобщена на любое количество устройств k

w_i = v_i/(v1 + v2 + ...v_k) 

Так vsum в моем алгоритме означает «сумму скоростей». Наконец, поскольку v_i является пиксели со временем это n_i/t_i, который, наконец, дает

w_i = n_i/t_i/(n1/t1 + n2/t2 + ...n_k/t_k) 

, который является вторым членом в моей формуле для расчета весов.

Answer 3

Хорошо: Просто выберите первое совместимое устройство. В большинстве систем есть только один.

Лучше: Вы можете очень грубо оценка производительности устройства путем умножения информации результата CL_DEVICE_MAX_COMPUTE_UNITS устройства по информации результата CL_DEVICE_MAX_CLOCK_FREQUENCY устройства. В зависимости от вашей рабочей нагрузки вы можете включить другие показатели, такие как размер памяти. Вы можете смешать их в зависимости от вашей рабочей нагрузки.

Лучшее: Контрольный показатель с вашим конкретным рабочим процессом на каждом устройстве. Это действительно единственный способ узнать наверняка, поскольку все остальное - всего лишь догадка.

Наконец, пользователь может заботиться о том, какой из своих графических процессоров вы используете, поэтому у вас должен быть способ переопределить ваш программный выбор независимо от выбранного вами метода.