OpenCL - Векторизация vs In-thread for loop

Question

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

__kernel something(whatever) 
{ 
    unsigned int glIDx = get_global_id(0); 

    for(condition_from_whatever) 
    { 

    }//alternatively, do while 

} 

где идентификатор (0) известен заранее, а не:

__kernel something(whatever) 
{ 
    unsigned int glIDx = get_global_id(0); 
    unsigned int glIDy = get_global_id(1); // max "unroll dimension" 

    if(glIDy_meets_condition) 
     do_something(); 
    else 
     dont_do_anything(); 

} 

, которые обязательно выполнить для полного ВОЗМОЖНО RANGE из glIDy, не имея возможности заранее прекратить, так как в этой дискуссии:

Killing OpenCL Kernels

Я не могу показаться, чтобы найти какой-либо конкретной информации о затраты на динамические измерения forloops/do-while в ядрах, хотя я вижу их везде в ядрах в Nvidia и SDK от AMD. Я помню, что читал что-то о том, насколько более апериодическим является внутренняя ветвь состояния, тем хуже производительность.

Актуальна:

Есть ли более эффективный способ борьбы с этим на архитектуре GPU, чем первая схема я предложил?

Я также открыт для общей информации об этой теме.

Спасибо.

Answer 1

Я не думаю, что есть общий ответ, который можно дать на этот вопрос. Это действительно зависит от вашей проблемы.

Однако вот некоторые соображения по этой теме:

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

Таким образом, если в пределах деформации у вас есть расхождение из-за if if (или цикла for с другим номером итераций), выполнение будет сериализовано.Иными словами, потоки в этом варпе, следуя первому пути, будут выполнять свои задания, остальные будут простаивать. Как только их работа будет завершена, эти потоки будут работать, пока остальные начнут работать.

Другая проблема возникает с этими утверждениями, если вам нужно синхронизировать свои потоки с барьером. У вас будет неопределенное поведение, если не все потоки попадут в барьер.

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

Зная, что деформация состоит из 32 потоков и волнового фронта 64 (возможно, не на старых графических процессорах AMD - не уверен), вы можете сделать размер ваших хорошо организованных рабочих групп равными или кратными этим числам. Обратите внимание, что это довольно упрощено, поскольку необходимо учитывать некоторые другие проблемы. См. Например, this question и ответ, полученный Chanakya.sun (возможно, больше копать в этой теме было бы неплохо).

В случае, если ваша проблема не может быть организована так, как описано выше, я бы предложил рассмотреть возможность использования OpenCL для процессоров, которые неплохо справляются с ветвлением. Если я хорошо помню, обычно у вас будет один рабочий элемент на рабочую группу. В этом случае лучше проверить документацию от Intel и AMD для CPU. Мне также очень нравится глава 6 из Heterogeneous Computing with OpenCL, в которой объясняются различия между использованием OCL с графическими процессорами и процессорами при программировании.

Мне нравится this article тоже. Это в основном дискуссия об увеличении производительности для простого сокращения на GPU (не ваша проблема), но в последней части статьи рассматривается также производительность процессоров.

Последнее, что касается ваших комментариев к ответу, предоставленному @Oak о поддержке поддержки очереди потоков внутри устройства, которая на самом деле называется динамическим параллелизмом. Эта функция, очевидно, решит вашу проблему, но даже с использованием CUDA вам понадобится устройство с возможностью 3.5 или выше. Поэтому даже графические процессоры NVIDIA с архитектурой Kepler GK104 не поддерживают его (возможность 3.0). Для OCL динамический параллелизм является частью стандартной версии 2.0. (насколько я знаю, пока не реализовано).

Answer 2

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

Два предостережения:

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