OpenCL Kernel Troubles

Question

Привет, я создал два ядра, чтобы выполнить простую соответствующую программу deshredder, которая будет запускаться с OpenCL и синхронизирована. Эти два ядра делают то, что они должны делать, но один работает намного медленнее, чем другой, по причине, которую я не могу расшифровать:/Единственное реальное различие заключается в том, как я храню отправленные данные и как происходит сопоставление.

__kernel void Horizontal_Match_Orig( 
__global int* allShreds, 
__global int* matchOut, 
const unsigned int shredCount, 
const unsigned int pixelCount) 

{ 
    int match = 0; 
    int GlobalID = get_global_id(0); 
    int currShred = GlobalID/pixelCount; 
    int thisPixel = GlobalID - (currShred * pixelCount); 
    int matchPixel = allShreds[GlobalID];//currShred*pixelCount+thisPixel]; 
    for (int i = 0; i < shredCount; i++) 
    { 

     match = 0; 
     if (matchPixel == allShreds[(i * pixelCount) + thisPixel]) 
     { 
      if (matchPixel == 0) 
      { 
       match = match + 150; 
      } 
      else match = match + 1; 
     } 
     else match = match - 50; 
     atomic_add(&matchOut[(currShred * shredCount) + i], match); 
    } 
} 

Это ядро ​​получает вскрывать края по горизонтали, так что пиксели одного клочка занимают положение 0 до п в allShreds массива, а затем пиксели следующего клочка сохраняются из поз п + 1 т (где n = количество пикселей, а m - количество добавленных пикселей). Каждый поток ГПУ получает один пиксель, чтобы работать с и соответствует его против соответствующего пикселя всех остальных клочков (включая себя)

__kernel void Vertical(
    __global int* allShreds, 
    __global int* matchOut, 
    const int numShreds, 
    const int pixelsPerEdge) 
{ 
    int GlobalID = get_global_id(0); 
    int myMatch = allShreds[GlobalID]; 
    int myShred = GlobalID % numShreds; 
    int thisRow = GlobalID/numShreds; 
    for (int matchShred = 0; matchShred < numShreds; matchShred++) 
    { 
     int match = 0; 
     int matchPixel = allShreds[(thisRow * numShreds) + matchShred]; 
     if (myMatch == matchPixel) 
     { 
      if (myMatch == 0) 
       match = 150; 
      else 
       match = 1; 
     } 
     else match = -50; 
      atomic_add(&matchOut[(myShred * numShreds) + matchShred], match); 
    } 
} 

Это ядро ​​получает вскрывать края по вертикали, так что первые пиксели всех клочки сохраняются в позиции от 0 до n, тогда 2-й пиксел всех фрагментов сохраняется в pos n + 1 ot m (где n = количество фрагментов, а m = количество клочков, добавленных в n). Процесс похож на предыдущий, где каждый поток получает пиксель и сопоставляет его с соответствующим пикселем каждого из других фрагментов.

Оба дают те же результаты, что и результаты, проверенные на чисто последовательную программу. Теоретически они должны работать примерно столько же времени, что и вероятность того, что вертикальный работает быстрее, так как атомная добавка не должна влиять на нее так же ... Однако она работает намного медленнее ... Любые идеи?

Это код, я использую, чтобы запустить его (я использую C# обертку для него):

theContext.EnqueueNDRangeKernel(1, null, new int[] { minRows * shredcount }, null, out clEvent); 

с общей глобальной рабочей нагрузкой, равной общее количество пикселей (#Shreds X #Pixels в каждом один).

Любая помощь будет принята с благодарностью

Answer 1

Два ядра делать то, что они должны делать, но один работает гораздо медленнее, чем другой по причине я не могу расшифровать:/Единственное реальное отличие состоит в том, как Я храню отправленные данные и как происходит сопоставление.

И все это имеет значение. Это классическая проблема слияния. Вы не указали свою модель GPU или поставщика в своем вопросе, поэтому мне придется оставаться неопределенным, поскольку фактические цифры и поведение полностью зависят от оборудования, но общая идея достаточно переносима.

Рабочие элементы в запросах памяти памяти GPU (чтение и запись) вместе (с помощью «warp»/«wavefront»/«sub-group») к движку памяти. Этот движок обслуживает память в транзакциях (куски размером от двух до четырех от 128 до 128 байтов). Давайте предположим размер 128 для следующего примера.

Введите коалесцирующем доступа к памяти: если 32 пунктов работы перекоса чтения 4 байта (int или float), которые последовательны в памяти, двигатель памяти будет выпуска одной транзакции обслуживать все 32 запросов. Но для каждого чтения, которое составляет более 128 байт, кроме другого, должна быть выпущена другая транзакция. В худшем случае это 32 транзакции по 128 байт, что намного дороже.

---

горизонтальное ядро ​​делает следующий доступ:

allShreds[(i * pixelCount) + thisPixel] 

(i * pixelCount) постоянен по рабочим элементам, только thisPixel меняется. Учитывая ваш код и при условии, что рабочий элемент 0 имеет thisPixel = 0, тогда рабочий элемент 1 имеет thisPixel = 1 и так далее. Это означает, что ваши рабочие элементы запрашивают смежные чтения, поэтому вы получаете полностью объединенный доступ. Аналогично для звонка atomic_add.

С другой стороны, ваше вертикальное ядро ​​выполняет следующие доступы:

allShreds[(thisRow * numShreds) + matchShred] 
// ... 
matchOut[(myShred * numShreds) + matchShred] 

matchShred и numShreds постоянны между потоками, только thisRow и myShred изменяются. Это означает, что вы запрашиваете чтение, которое находится в numShreds друг от друга. Это не последовательный доступ и поэтому не объединен.