Предотвращение вычисления дважды в opencl

Question

Я довольно новичок в программировании ocl.

У меня есть 2 миллиона полигонов (по четыре строки - для этих примеров - но это вариант), что мне нужно найти пересечение с 1000 эллипсами.

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

Для этого я создал буфер poylgons со всеми точками и буфером эллипсов.

Мой выходной буфер - это 1000 элементов int, равных 0 во всех элементах. и ядро ​​установит 1 (когда найдет пересечение) в правом индексе в соответствии с индексом эллипса.

Я запустил ядро ​​с глобальным - 2 dim, {2million, 1000}.

__kernel void polygonsIntersectsEllipses( __global const Point* pts, 
              __global const Ellipse* ellipses, 
              __global int* output) 
{ 
    int polygonIdx = get_global_id(0); 
    int ellipseIdx = get_global_id(1); 

    if (<<isIntersects>>) { 
     output[ellipseIdx] = 1; 
    } 
} 

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

Я попытался проверить output[ellipseIdx] != 0 перед тестом на пересечение, но производительность не изменилась так сильно.

Я попытался сделать одномерное глобальное - давая 1000 (эллипсы), а в ядре пробегают миллионы полигонов и останавливаются, когда я его найду, но все равно не так сильно меняются.

Я делаю что-то неправильно? могу ли я ускорить эту операцию? любые подсказки?

EDIT

Используя наконечник от @Moises и делать исследования, выделяют я изменить свой код, чтобы запустить 2 миллиона раз, один размер. используя элементы групповой работы. изменил все мои структуры на нативные типы, пропустил работу модуля. И в основном, когда я мог копировать данные из глобальной в частную/локальную память, я это сделал.

Мой местный размер - мое устройство CL_DEVICE_MAX_WORK_GROUP_SIZE, в моем CPU & gpu - 1024, поэтому за один проход я покрываю все свои 1000 эллипсов.

стороне хоста

size_t global = 1999872; // 2 million divided by 1024 - for the test 
size_t local = 1024; 

Мой код выглядеть как это сейчас

__kernel void polygonsIntersectsEllipses( __global const float4* pts, 
              __global const float4* ellipses, 
              int ellipseCount, 
              __local float4* localEllipses, 
              __global int* output) 
{ 
    // Saving the ellipses to local memory 
    int localId = get_local_id(0); 
    if (localId < eCount) 
     localEllipses[localId] = ellipses[localId]; 

    barrier(CLK_LOCAL_MEM_FENCE); 

    // Saving the current polygon into private memory 
    int polygonIdx = get_global_id(0); 
    float2 private_pts[5]; 
    for (int currPtsIdx = 0; currPtsIdx < 4; currPtsIdx++) 
    { 
     private_pts[currPtsIdx] = pts[polygonIdx * 4 + currPtsIdx]; 
    } 

    // saving the last point as first so i will not use modulus for cycling, in the intersection algorithm 
    private_pts[4] = private_pts[0]; 

    // Run over all the ellipse in the local memory including checking if already there is an output 
    for (int ellipseIdx = 0; ellipseIdx < ellipseCount && output[ellipseIdx] == 0; ++ellipseIdx) { 
     if (<<isIntersects Using localEllipses array and private_pts>>) { 
      output[ellipseIdx] = 1; 
     } 
    } 
} 

Результаты деятельности

CPU не так значительно улучшилось - 1,1 быстрее после изменения.

GPU - 6,5 раз быстрее (я очень рад)

Есть ли место я могу улучшить еще больше? Напоминание, когда один из полигонов пересекает эллипс, нам не нужно проверять остальную часть полигонов. Как мне это сделать ? мой трюк с запросом на выходное значение на самом деле не работает - показатели одинаковы с или без него

Answer 1

Я понимаю, что все ваши потоки 2millionx1000, прочитайте его собственные данные о полигонах и эллипсе справа?Итак, для каждого многоугольника каждый поток читает 1000 раз по той же позиции памяти (с данными полигона), не так ли? Чтобы избежать этого связанного с памятью поведения, вы можете создать только 2 миллиона потоков и использовать цикл for из 1000 итераций для повторения числа эллипсов. Или промежуточное решение, имеющее сетку потоков 2millionx64, где каждый поток вычисляет 16 эллипсов для каждого многоугольника. Я не знаю, лучше ли это, чем ваше решение, но они избегают излишних доступов к памяти.

Привет, Моисес

Answer 2

Оптимизации:

• Используйте минимальное количество памяти. __global int* output для хранения булевых слишком много, используйте вместо этого char. Или даже лучше, используйте двоичный массив. (двоичные операции быстрые по сравнению с глобальными)

• Вы не должны читать снова из глобальной памяти из каждой нити output[ellipseIdx] == 0. Это ужасно медленно, вместо этого сохраните его в локальной памяти с данными elipses в начале. ПРИМЕЧАНИЕ. Только локальные группы, запущенные ПОСЛЕ того, что группа, получившая совпадение, выиграют от ускорения. Однако это спасет тонны глобальных чтений, что намного лучше, чем сохранение нескольких вычислений. Кроме того, местная группа не может выиграть, поскольку, когда рабочий элемент находит совпадение, все обработанные рабочие элементы уже обработаны.

  __kernel недействительного polygonsIntersectsEllipses (

  __global const float4* pts, 
  __global const float4* ellipses, 
  int ellipseCount, 
  __local float4* localEllipses, 
  __local char* localOuts, 
  __global char* output){ 
  
  // Saving the ellipses to local memory 
  int localId = get_local_id(0); 
  if (localId < eCount) 
      localOuts[localId] = output[localId]; 
  barrier(CLK_LOCAL_MEM_FENCE); 
  if (localId < eCount && localOuts[localId]) // Do not copy elipses if we are not going to check them anyway 
      localEllipses[localId] = ellipses[localId]; 
  
  barrier(CLK_LOCAL_MEM_FENCE); 
  
  // Saving the current polygon into private memory 
  int polygonIdx = get_global_id(0); 
  float2 private_pts[5]; 
  for (int currPtsIdx = 0; currPtsIdx < 4; currPtsIdx++) 
  { 
      private_pts[currPtsIdx] = pts[polygonIdx * 4 + currPtsIdx]; 
  } 
  
  // saving the last point as first so i will not use modulus for cycling, in the intersection algorithm 
  private_pts[4] = private_pts[0]; 
  
  // Run over all the ellipse in the local memory including checking if already there is an output 
  for (int ellipseIdx = 0; ellipseIdx < ellipseCount; ++ellipseIdx) { 
      if (localOuts[ellipseIdx] == 0){ 
       if (<<isIntersects Using localEllipses array and private_pts>>) { 
        localOuts[ellipseIdx] = 1; 
       } 
       barrier(CLK_LOCAL_MEM_FENCE); 
       if(localOuts[ellipseIdx] && localId == 0){ 
        output[ellipseIdx] = 1; 
       } 
      } 
  } 
  

  }