Почему локальная память в этом алгоритме OpenCL такая медленная? - PullRequest
Новая
       24

Почему локальная память в этом алгоритме OpenCL такая медленная?

0 голосов
/ 16 января 2019



Я пишу код OpenCL. Мое ядро ​​должно создать специальный «аккумуляторный» вывод на основе входного изображения. Я попробовал два понятия, и оба одинаково медленны, хотя второе использует локальную память. Не могли бы вы помочь мне определить, почему версия локальной памяти такая медленная? Целевым графическим процессором для ядер является AMD Radeon Pro 450. 

// version one
__kernel void find_points(__global const unsigned char* input, __global unsigned int* output) {
  const unsigned int x = get_global_id(0);
  const unsigned int y = get_global_id(1);

  int ind;

  for(k = SOME_BEGINNING; k <= SOME_END; k++) {
    // some pretty wild calculation
    // ind is not linear and accesses different areas of the output
    ind = ...
    if(input[y * WIDTH + x] == 255) {
      atomic_inc(&output[ind]);
    }
  }

}

// variant two
__kernel void find_points(__global const unsigned char* input, __global unsigned int* output) {
  const unsigned int x = get_global_id(0);
  const unsigned int y = get_global_id(1);

  __local int buf[7072];
  if(y < 221 && x < 32) {
    buf[y * 32 + x] = 0;
  }
  barrier(CLK_LOCAL_MEM_FENCE);

  int ind;
  int k;

  for(k = SOME_BEGINNING; k <= SOME_END; k++) {
    // some pretty wild calculation
    // ind is not linear and access different areas of the output
    ind = ...
    if(input[y * WIDTH + x] == 255) {
      atomic_inc(&buf[ind]);
    }
  }

  barrier(CLK_LOCAL_MEM_FENCE);
  if(get_local_id(0) == get_local_size(0) - 1)
    for(k = 0; k < 7072; k++)
      output[k] = buf[k];
  }

}

Я ожидаю, что второй вариант быстрее первого, но это не так. Иногда это даже медленнее.

1 Ответ

0 голосов
/ 16 января 2019

Размер локального буфера __local int buf[7072] (28288 байт) слишком велик. Я не знаю, насколько велика общая память для AMD Radeon Pro 450, но, скорее всего, она составляет 32 или 64 КБ на вычислительную единицу.

32768/28288 = 1, 65536/28288 = 2 означает, что только 1 или максимум 2 волновых фронта (64 рабочих элемента) могут работать только одновременно, поэтому занятость вычислительного устройства очень очень низкая, следовательно, низкая производительность.

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

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

...