Какие проблемы хорошо подходят для вычислений на GPU?

84

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

Я знаю, что это основной вопрос, но большая часть моих поисков оказывается у людей, явно выступающих за то или иное без реального объяснения причин или каких-то смутных практических правил. Ищете более полезный ответ здесь.

фомиты
источник

Ответы:

63

У оборудования GPU есть две сильные стороны: необработанные вычисления (FLOP) и пропускная способность памяти. Самые сложные вычислительные проблемы попадают в одну из этих двух категорий. Например, плотная линейная алгебра (A * B = C или Solve [Ax = y] или Diagonalize [A] и т. Д.) Находится где-то в спектре пропускной способности вычислений / памяти в зависимости от размера системы. Быстрые преобразования Фурье (БПФ) также подходят для этой формы с высокими требованиями к совокупной пропускной способности. Как и другие преобразования, алгоритмы на основе сетки / сетки, Монте-Карло и т. Д. Если вы посмотрите на примеры кода NVIDIA SDK , вы сможете почувствовать проблемы, которые чаще всего решаются.

Я думаю, что более поучительный ответ на вопрос: «Какие проблемы действительно плохи для графических процессоров?» Большинство проблем, которые не попадают в эту категорию, могут быть запущены на GPU, хотя некоторые требуют больше усилий, чем другие.

Проблемы, которые плохо отображаются, как правило, слишком малы или слишком непредсказуемы. Очень мелким проблемам не хватает параллелизма, необходимого для использования всех потоков на графическом процессоре, и / или они могут помещаться в низкоуровневый кэш на процессоре, что существенно повышает производительность процессора. Непредсказуемые проблемы имеют слишком много значимых ветвей, которые могут помешать эффективной потоковой передаче данных из памяти графического процессора в ядра или уменьшить параллелизм, нарушая парадигму SIMD (см. « Различные отклонения »). Примеры таких проблем:

  • Большинство графовых алгоритмов (слишком непредсказуемо, особенно в пространстве памяти)
  • Разреженная линейная алгебра (но это плохо для процессора)
  • Небольшие проблемы с обработкой сигналов (например, FFT менее 1000 точек)
  • Поиск
  • Сортировать
Макс Хатчинсон
источник
3
Тем не менее, GPU решения для этих «непредсказуемых» проблемы являются возможным и, в то время как в настоящее время не представляется возможным , как правило, может получить значение в будущем.
оставил около
6
Я хотел бы специально добавить ответвления в список ограничителей производительности графических процессоров. Вы хотите, чтобы все ваши (сотни) выполняли одну и ту же инструкцию (как в SIMD) для выполнения действительно параллельных вычислений. Например, на картах AMD, если какой-либо из потоков команд встречает ветвь и должен расходиться - все волновой фронт (параллельная группа) расходится. Если другие подразделения с фронта волны не должны расходиться - они должны выполнить второй проход. Это то, что maxhutch подразумевает под предсказуемостью, я думаю.
Фиолетовый Жираф
2
@VioletGiraffe, это не обязательно так. В CUDA (то есть на графических процессорах Nvidia) расхождение ветвей влияет только на текущий перекос, который составляет не более 32 потоков. Различные деформации, хотя и выполняют один и тот же код, не являются синхронными, если они явно не синхронизированы (например, с помощью __synchtreads()).
Педро
1
@Pedro: Да, но в целом ветвление снижает производительность. Для высокопроизводительных кодов (чем код GPU не является?), Это практически необходимо учитывать.
jvriesem
21

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

Основная трудность в создании высокопроизводительного кода GPU состоит в том, что у вас есть тонна ядер, и вы хотите, чтобы все они использовались максимально эффективно. Проблемы, которые имеют нерегулярные схемы доступа к памяти или не имеют высокой арифметической интенсивности, усложняют эту задачу: либо вы тратите много времени на передачу результатов, либо вы тратите много времени на извлечение данных из памяти (что медленно!), И не хватает времени на обработку чисел. Конечно, возможность параллелизма в вашем коде имеет решающее значение для его способности хорошо реализовываться и на GPU.

Reid.Atcheson
источник
Можете ли вы указать, что вы подразумеваете под обычными шаблонами доступа к памяти?
Fomite
1
ответ maxhutch лучше моего Под шаблоном обычного доступа я подразумеваю, что доступ к памяти осуществляется временным и пространственно локальным способом. То есть: вы не делаете огромные скачки вокруг памяти неоднократно. Это также что-то вроде пакетной сделки, которую я заметил. Это также означает, что ваши шаблоны доступа к данным могут быть заранее определены компилятором или программистом, так что ветвление (условные выражения в коде) сведено к минимуму.
Reid.Atcheson
15

Это не предназначено как ответ самостоятельно, а скорее как дополнение к другим ответам maxhutch и Reid.Atcheson .

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

Точнее, ядро ​​должно вписываться в регистр каждого многопроцессорного блока (или вычислительного блока ) графического процессора. Точный размер регистра зависит от графического процессора.

Учитывая, что ядро ​​достаточно маленькое, необработанные данные проблемы должны помещаться в локальную память графического процессора (читай: локальная память (OpenCL) или разделяемая память (CUDA) вычислительного устройства). В противном случае даже высокая пропускная способность памяти графического процессора будет недостаточной, чтобы постоянно обрабатывать элементы обработки .
Обычно эта память составляет около 16 до 32 KiByte больших .

Torbjörn
источник
Разве локальная / разделяемая память каждого процессора не распределяется между всеми десятками (?) Потоков, работающих в одном кластере ядер? В этом случае вам не нужно на самом деле уменьшать рабочий набор данных, чтобы получить максимальную производительность от графического процессора?
Дэн Нили
Локальная / совместно используемая память блока обработки доступна только самому вычислительному блоку и, таким образом, совместно используется только элементами обработки этого вычислительного блока. Глобальная память видеокарты (обычно 1 ГБ) доступна для всех процессоров. Пропускная способность между элементами обработки и локальной / разделяемой памятью очень быстрая (> 1 ТБ / с), но пропускная способность глобальной памяти намного медленнее (~ 100 ГБ / с) и должна быть распределена между всеми вычислительными блоками.
Торбьерн
Я не спрашивал об основной памяти GPU. Я думал, что оперативная память была выделена только на уровне ядра кластера, а не на отдельное ядро. например, для NVIDIA GF100 / 110 GPU; для каждого из 16 кластеров SM не 512 ядер CUDA. Для каждого SM, предназначенного для одновременной работы до 32 потоков, для максимизации производительности графического процессора потребуется поддерживать рабочий набор в диапазоне 1 КБ / поток.
Дэн Нили
@Torbjoern Вы хотите, чтобы все конвейеры выполнения GPU были заняты, графические процессоры достигают этого двумя способами: (1) наиболее распространенный способ - увеличить занятость, или, иначе говоря, увеличить число параллельных потоков (небольшие ядра используют меньше общие ресурсы, чтобы вы могли иметь больше активных потоков); может быть, лучше, это (2) увеличить параллелизм уровня команд в вашем ядре, чтобы вы могли иметь более крупное ядро ​​с относительно низкой загрузкой (небольшое количество активных потоков). Смотрите bit.ly/Q3KdI0
fcruz
11

Вероятно, более техническое дополнение к предыдущим ответам: графические процессоры CUDA (то есть Nvidia) можно описать как набор процессоров, которые работают автономно на 32 потоках каждый. Потоки в каждом процессоре работают в режиме блокировки (например, SIMD с векторами длины 32).

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

Если у вашего кода не распараллеливание красиво / автоматически сотни / тысячи нитей, вы можете быть в состоянии разбить его на отдельные асинхронные задачи , которые действительно Распараллеливать хорошо, и выполнять те с только 32 потоков , выполняющихся в карцер шаге. CUDA предоставляет набор элементарных инструкций, которые позволяют реализовать взаимные исключения, которые, в свою очередь, позволяют процессорам синхронизироваться между собой и обрабатывать список задач в парадигме пула потоков . В этом случае ваш код будет работать так же, как в многоядерной системе, но имейте в виду, что каждое ядро ​​имеет 32 собственных потока.

Вот небольшой пример использования CUDA, как это работает

/* Global index of the next available task, assume this has been set to
   zero before spawning the kernel. */
__device__ int next_task;

/* We will use this value as our mutex variable. Assume it has been set to
   zero before spawning the kernel. */
__device__ int tasks_mutex;

/* Mutex routines using atomic compare-and-set. */
__device__ inline void cuda_mutex_lock ( int *m ) {
    while ( atomicCAS( m , 0 , 1 ) != 0 );
    }
__device__ inline void cuda_mutex_unlock ( int *m ) {
    atomicExch( m , 0 );
    }

__device__ void task_do ( struct task *t ) {

    /* Do whatever needs to be done for the task t using the 32 threads of
       a single warp. */
    }

__global__ void main ( struct task *tasks , int nr_tasks ) {

    __shared__ task_id;

    /* Main task loop... */
    while ( next_task < nr_tasks ) {

        /* The first thread in this block is responsible for picking-up a task. */
        if ( threadIdx.x == 0 ) {

            /* Get a hold of the task mutex. */
            cuda_mutex_lock( &tasks_mutex );

            /* Store the next task in the shared task_id variable so that all
               threads in this warp can see it. */
            task_id = next_task;

            /* Increase the task counter. */
            next_tast += 1;

            /* Make sure those last two writes to local and global memory can
               be seen by everybody. */
            __threadfence();

            /* Unlock the task mutex. */
            cuda_mutex_unlock( &tasks_mutex );

            }

        /* As of here, all threads in this warp are back in sync, so if we
           got a valid task, perform it. */
        if ( task_id < nr_tasks )
            task_do( &tasks[ task_id ] );

        } /* main loop. */

    }

Затем вам нужно вызвать ядро, main<<<N,32>>>(tasks,nr_tasks)чтобы убедиться, что каждый блок содержит только 32 потока и, таким образом, помещается в одну деформацию. В этом примере я также для простоты предположил, что задачи не имеют никаких зависимостей (например, одна задача зависит от результатов другой) или конфликтов (например, работа с одной и той же глобальной памятью). Если это так, то выбор задачи становится немного сложнее, но структура по сути та же самая.

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

Pedro
источник
2
Это технически верно, но для получения высокой пропускной способности памяти необходим высокий параллелизм, и существует ограничение на количество асинхронных вызовов ядра (в настоящее время 16). У тебя также есть тонны недокументированного поведения, связанного с расписанием в текущем выпуске. Я бы посоветовал не полагаться на асинхронные ядра для улучшения производительности на данный момент ...
Макс Хатчинсон
2
Все, что я описываю, может быть сделано за один вызов ядра. Вы можете сделать N блоков по 32 потока каждый, чтобы каждый блок помещался в одну основу. Затем каждый блок получает задачу из глобального списка задач (доступ контролируется с помощью атомарных элементов / мьютексов) и вычисляет ее, используя 32 ступени с блокировкой. Все это происходит за один вызов ядра. Если вам нужен пример кода, дайте мне знать, и я опубликую его.
Педро
4

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

Брайан Борхерс
источник
Я хочу не согласиться. Большинство (или все) более новых графических процессоров имеют встроенную поддержку двойной точности. Почти каждый такой графический процессор сообщает о вычислениях с двойной точностью, работающих примерно на половину скорости одинарной точности, вероятно, из-за простого удвоения требуемого доступа к памяти / пропускной способности.
Годрик Провидец
1
Хотя верно и то, что новейшие и лучшие карты Nvidia Tesla действительно обеспечивают максимальную производительность с двойной точностью, что составляет половину максимальной производительности с одинарной точностью, соотношение составляет 8: 1 для более распространенных карт потребительского уровня архитектуры Fermi.
Брайан Борчерс
@GodricSeer Соотношение 2: 1 с плавающей точкой SP и DP имеет очень мало общего с пропускной способностью и почти все зависит от того, сколько аппаратных единиц существует для выполнения этих операций. Распространено повторно использовать файл регистров для SP и DP, следовательно, модуль с плавающей запятой может выполнить 2 операции SP как операции DP. Существует множество исключений из этого проекта, например, IBM Blue Gene / Q (не имеет логики SP и, следовательно, SP работает с ~ 1,05x DP). Некоторые графические процессоры имеют коэффициенты, отличные от 2, например, 3 и 5.
Джефф
Прошло четыре года с тех пор, как я написал этот ответ, и текущая ситуация с графическими процессорами NVIDIA заключается в том, что для линий GeForce и Quadro соотношение DP / SP теперь составляет 1/32. Графические процессоры NVIDIA Tesla обладают гораздо более высокой производительностью с двойной точностью, но и стоят намного дороже. С другой стороны, AMD не наносит ущерба производительности с двойной точностью на своих графических процессорах Radeon.
Брайан Борчерс
4

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

Это может быть принято как дополнительный ответ на превосходные ответы выше.

labotsirc
источник
4

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

GT.
источник
2

Графические процессоры имеют большие задержки ввода / вывода, поэтому для насыщения памяти необходимо использовать множество потоков. Чтобы держать деформацию занятой, требуется много потоков. Если путь кода равен 10 часам, а задержка ввода-вывода 320 тактов, 32 потока должны приблизиться к насыщению деформации. Если путь к коду составляет 5 часов, то удвойте потоки.

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

Доступ к памяти осуществляется по строке кэша, обычно 32 байта. Загрузка одного байта имеет сравнимую стоимость с 32 байтами. Итак, объедините хранилище, чтобы увеличить локальность использования.

Есть много регистров и локальной оперативной памяти для каждой деформации, что позволяет для совместного использования соседей.

Моделирование близости больших наборов должно хорошо оптимизироваться.

Случайный ввод / вывод и однопоточность - радость убийства ...

user14381
источник
Это действительно захватывающий вопрос; Я спорю со мной относительно того, возможно ли (или стоит ли это усилие) «параллельное изложение» достаточно простой задачи (обнаружение краев на аэрофотоснимках), когда каждая задача занимает ~ 0,06 с, но для выполнения необходимо ~ 1,8 млн. Задач ( в год, для данных за 6 лет: задачи определенно разделимы) ... таким образом, ~ 7,5 дней вычислительного времени на одном ядре. Если каждый вычисление выполнялось быстрее на графическом процессоре, и задание можно распараллелить по 1 на nGPUcores [n small], то действительно ли время выполнения задания может упасть до ~ 1 часа? Кажется маловероятным
GT.
0

Представьте себе проблему, которая может быть решена с помощью грубой силы, например, коммивояжера. Тогда представьте, что у вас есть серверные стойки с 8-ю шипучими видеокартами каждая, и каждая карта имеет 3000 ядер CUDA.

Просто решите ВСЕ возможные маршруты продавца, а затем сортируйте по времени / расстоянию / некоторому показателю. Конечно, вы отбрасываете почти 100% своей работы, но грубая сила иногда является жизнеспособным решением.

Criggie
источник
У меня был доступ к небольшой ферме из 4 таких серверов в течение недели, и за пять дней я сделал больше распределенных блоков.net, чем за предыдущие 10 лет.
Кригги
-1

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

Многие формулы могут быть простыми в написании, но трудными для вычисления, например, в матричной математике вы получаете не один ответ, а много значений.

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

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

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

ГПУ - это трудолюбивый работник. Процессор управляет полным хаосом и не может обрабатывать каждую деталь.

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

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

Когда я ухожу, не забывайте начинать с простого, начинать быстро, быстро, быстро, быстро и никогда не прекращать попытки.

Эндрю Дж. Корби
источник