Странное поведение с использованием локальной памяти в OpenCL

Question

В настоящее время я работаю над проектом, подающим в суд на OpenCL на NVIDIA Tesla C1060 (версия драйвера 195.17). Однако у меня странное поведение, которое я не могу объяснить. Вот код, который озадачивает меня (сокращен для ясности и в целях тестирования):

kernel void TestKernel(global const int* groupOffsets, global       float* result,     
                       local        int* tmpData,             const int    itemcount)
{
   unsigned int groupid    = get_group_id(0);
   unsigned int globalsize = get_global_size(0);
   unsigned int groupcount = get_num_groups(0);

   for(unsigned int id = get_global_id(0); id < itemcount; id += globalsize, groupid += groupcount)
   {
      barrier(CLK_LOCAL_MEM_FENCE);
      if(get_local_id(0) == 0)
         tmpData[0] = groupOffsets[groupid]; 
      barrier(CLK_LOCAL_MEM_FENCE);
      int offset = tmpData[0];
      result[id]   = (float) offset;
   }
}

Этот код должен загрузить смещение для каждой рабочей группы в локальную память, а затем прочитать его обратно и записать в соответствующую запись выходного вектора. Для большинства рабочих элементов это работает, но для каждой рабочей группы рабочие элементы с локальными идентификаторами с 1 по 31 читают неправильное значение. Мой выходной вектор (для рабочей группы =128) выглядит следующим образом:

index       0: 0
index   1- 31: 470400
index  32-127: 0
index     128: 640
index 129-159: 471040
index 160-255: 640
index     256: 1280
index 257-287: 471680
index 288-511: 1280
...

результат, который я ожидал, будет

index   0-127: 0
index 128-255: 640
index 256-511: 1280
...

Странно то, что проблема возникает только тогда, когда я использую меньше рабочих элементов itemcount (поэтому он работает, как и ожидалось, когда globalsize>=itemcount, то есть каждый рабочий элемент обрабатывает только одну запись). Я предполагаю, что это как-то связано с циклом. Кто-нибудь знает, что я делаю не так и как это исправить?

Обновление: я узнал, что, кажется, работает, если я изменяю

if(get_local_id(0) == 0)
     tmpData[0] = groupOffsets[groupid]; 

в

if(get_local_id(0) < 32)
     tmpData[0] = groupOffsets[groupid]; 

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

Answer 1

Во-первых, и это важно, вы должны быть осторожны, что itemcount является кратным локального размера работы, чтобы избежать расхождения при выполнении барьера.

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

Вы можете реализовать это следующим образом:

unsigned int itemcountrounded = get_local_size(0) * ((itemcount + get_local_size(0) - 1) / get_local_size(0));
for(unsigned int id = get_global_id(0); id < itemcountrounded; id += globalsize, groupid += groupcount)
{
    // ...
    if (id < itemcount)
        result[id]   = (float) offset;
}

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