Доступ/синхронизация с локальной памятью

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

У меня есть проблема: моя программа работает правильно до тех пор, пока размер проблемы не будет превышать 32.

__kernel void assort(
__global float *array, 
__local float *currentOutput, 
__local float *stimulations, 
__local int *noOfValuesAdded, 
__local float *addedValue, 
__local float *positionToInsert, 
__local int *activatedIdx, 
__local float *range, 
int size, 
__global float *stimulationsOut 
) 
{ 
int id = get_local_id(0); 
if (id == 0) {...} 

barrier(CLK_LOCAL_MEM_FENCE); 

for (int i = 2; i < size; i++) 
{ 
    int maxIdx; 
    if (id == 0) 
    { 
    addedValue[0] = array[i]; 
    {...} 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 


    if (id < noOfValuesAdded[0]){...} 
    else 
     barrier(CLK_LOCAL_MEM_FENCE); 
    barrier(CLK_LOCAL_MEM_FENCE); 
    if (activatedIdx[0] == -2) {...} 
    else {...} 

    barrier(CLK_LOCAL_MEM_FENCE); 
    if (positionToInsert[0] != -1.0f) {...} 

    barrier(CLK_LOCAL_MEM_FENCE); 
    stimulationsOut[id] = addedValue[0]; 
    return; 
    }

После некоторого следствия attemp я понял (путем проверки stimulationsOut), что addedValue [0] имеет diffrent значение из 33-й instanction ядра, а затем другое значение из 65-й (так что его, как [123 123 123 ... 123 (33-й элемент) 66 ... 66 66 66 66 66 .. (65-й элемент) 127 ... .. 127 ...])

__gl obal float * array READ_ONLY, и я не меняю addValue [0] рядом с первым, если в цикле for. Что может решить эту проблему?

Мои GPU функции: [https://devtalk.nvidia.com/default/topic/521502/gt650m-a-kepler-part-/]

После закомментировав из двух, если наша проблема тело не reccuring:

  /*if (activatedIdx[0] == -2) 
     { 
      if (noOfValuesAdded[0] == 2) 
      { 
       positionToInsert[0] = 0.99f; 
      } 
      else if (id != 0 && id != maxIdx 
        && stimulations[id] >= stimulations[(id - 1)] 
        && stimulations[id] >= stimulations[(id + 1)]) 
      { 
       if ((1.0f - (fabs((currentOutput[(id - 1)] - currentOutput[id]))/range[0])) < stimulations[(id - 1)]) 
        positionToInsert[0] = (float)id - 0.01f; 
        else 
       positionToInsert[0] = (float)id + 0.99f; 
      } 
     }*/

if (positionToInsert[0] != -1.0f) 
    { 
     float temp = 0.0f; 
     /*if ((float)id>positionToInsert[0]) 
     { 
      temp = currentOutput[id]; 
      barrier(CLK_LOCAL_MEM_FENCE); 
      currentOutput[id + 1] = temp; 
     } 
     else 
     { 
      barrier(CLK_LOCAL_MEM_FENCE); 
     }*/ 
     barrier(CLK_LOCAL_MEM_FENCE); 

     if (id == round(positionToInsert[0])) 
     { 
      currentOutput[id] = addedValue[0]; 
      noOfValuesAdded[0] = noOfValuesAdded[0] + 1; 
     } 
    }

Update: После устранения барьеров, алгоритм работы правильно, пока размер не превысит 768 (это странно 2 раза больше ядер на моем gpu). Я ожидал, что он будет работать до 1024 элементов, это максимальный размер рабочей группы. Я что-то упускаю?

источник

2016-12-07 Karol K

Я отправляю усеченную версию кода, уже измененную для расследования, до , чтобы облегчить ответы, поскольку ее не очень легко прочитать. –

Все рабочие элементы в деформации выполняют одну и ту же инструкцию в режиме блокировки. Размер Warp на Nvidia - 32 рабочих элемента. Если ядро работает правильно до 32 рабочих элементов, это говорит о том, что что-то не так с барьерами.

Docs для barrier сказать:

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

Я вижу, что это проблема в вашем ядре. Например здесь:

if ((float)id>positionToInsert[0]) 
{ 
    temp = currentOutput[id]; 
    barrier(CLK_LOCAL_MEM_FENCE); // <---- some work items may meet here 
    currentOutput[id + 1] = temp; 
} 
else 
{ 
    barrier(CLK_LOCAL_MEM_FENCE); // <---- other work items may meet here 
}

Вы могли бы исправить это:

if ((float)id>positionToInsert[0]) 
    temp = currentOutput[id]; 
barrier(CLK_LOCAL_MEM_FENCE); // <---- here all work items meet at the same barrier 
if ((float)id>positionToInsert[0]) 
    currentOutput[id + 1] = temp;

источник

2016-12-08 07:20:14 doqtor

Спасибо. Когда проект был запущен, он был написан в Java-инфраструктуре Aparapi и утверждает, что только количество вызовов localBarrier должно быть одинаковым для каждого ядра (https://github.com/aparapi/aparapi/blob/master/doc/UsingLocalMemory .md) –

После устранения барьеров, алгоритм работает правильно, пока размер не превышает 768 (который Жутко 2 раза количества ядер на моем ГПА). Я ожидал, что он будет работать до 1024 элементов, это максимальный размер рабочей группы. Я что-то упускаю?

источник

2016-12-08 10:08:13

CL_KERNEL_WORK_GROUP_SIZE возвращает 1024 –

Можете ли вы запустить с nvprof -print-gpu-trace и посмотреть размеры блока и сетки? –

Доступ/синхронизация с локальной памятью

ответ

Смежные вопросы