Перерыв и продолжение в ядре OpenCL

Question

Как реализовать их в ядре OpenCL? Это «возвращение»; эквивалентно «перерыву»?

Я использую OpenCL 1.2

Я хочу осуществить это с 3 вложено для петель, цикл через ЬурейуЮ структуру вложенных массивов.

EDIT

Реализованный мне нужно показать код, чтобы быть лучше поняты

в ядре ...

typedef struct tag_sfextras 
{ 
    float *high; 
    float *low; 
}sfextras; 

typedef struct tag_sdirection 
{ 
    int time; 
    float result; 
    sfextras *fextras; 
}sdirection; 

__kernel void Call(sdirection *_direction, 
        int _index, 
        int _start, 
        int _stop, 
        __global float *_result) 
{ 

    float _sum = 0.0f; 

    if (_index > 1) 
    { 
     _result[0] = 0.0f; 

     int i = get_global_id(0); 

     if (_direction[i].time >= _stop) 
     { 
      break;//or return?...      
     } 

     if (_direction[i].time < _start) 
     { 
      continue;// what to put here?...   
     } 
     else 
     { 
      _start = _direction[i].time + (1440 * 60); 
     } 

     int d = get_global_id(1); 
     int f = get_global_id(2); 

     float _fextras_weight = 0.0f;// need to zeroize on each inner loop (for f) 

     _fextras_weight += (float)pow(_direction[_index - 1].fextras[d].high[f] - _direction[i].fextras[d].high[f], 2.0f); 
     _fextras_weight += (float)pow(_direction[_index - 1].fextras[d].low[f] - _direction[i].fextras[d].low[f], 2.0f); 

     _result[0] += _fextras_weight*_direction[i].result; 
     _sum += _fextras_weight; 

    } 

    if (_sum > 0.0f) 
    { 
     _result[0] /= _sum; 
    } 
} 

В HOST (код, который я пытаюсь повторить в ядре для эффективности)

  if(_direction_index > 1) 
      { 
       _fextras = 0.0f; 
       for(int i=0;i<_direction_index-1;i++) 
       { 
        if(_direction[i].time >= _stop) 
         { 
          break; 
         } 

        if(_direction[i].time < _start) 
         { 
          continue; 
         } 
        else 
         { 
          _direction_start = _direction[i].time + (1440*60); 
         } 

        for(int d=0;d<_DIRECTION;d++) 
         { 
          for(int f=0;f<_FEXTRAS;f++) 
          { 
           float _fextras_weight = 0.0f; 

           _fextras_weight += (float)pow(_direction[_direction_index-1].fextras[d].high[f]-_direction[i].fextras[d].high[f],2.0f); 
           _fextras_weight += (float)pow(_direction[_direction_index-1].fextras[d].low[f]-_direction[i].fextras[d].low[f],2.0f); 

           _fextras += _fextras_weight*_direction[i].result; 
           _sum += _fextras_weight; 
          } 
         } 
       } 

       if(_sum > 0.0f) 
       { 
        _fextras /= _sum; 
       } 
      } 

Answer 1

Отмена всех остальных потоков opencl сделает их неопределенным поведением, потому что многие из них могут быть в середине записи/чтения глобальной/локальной памяти, а также могут очищать поток (останавливать все остальные/оставшиеся ядра/потоки). Вероятно, именно поэтому в opencl нет такой вещи.

Но вы можете добавить выходной массив, чтобы каждый поток записывал свое последнее состояние. Если у элемента есть код «return», вы должны проверить код «after_return», чтобы опустить результаты результатов этих результатов и принять «before_return». Это также потребует атомных операций на выходном каскаде, поэтому становится медленнее, что плохо.

Но вы можете спокойно вернуться из отдельных ядер:

Ниже код, составленный хорошо и рано вышел (закончилось исполнение ядра для некоторых потоков, но не все) в «возвращении» на HD7870 и R7-240 без потому что «return» не является одним из ограничений, применяемых OpenCL.

__kernel void rarToVideo(__global int * p,__global char * c) 
      { 
       ... 
          if (tmp) 
          { 
           foo=1; 
          } 
          else 
          { 
           return; 
          } 

      ... 
      } 

Использованы opencl 1.2 заголовки C++.

Но, если вам все еще нужно поддельное возвращение и нить не влияют на выходах других потоков/входы, то что-то вроде этого помогло бы:

// beginning phase of this thread 
    if(globalAtomicElement[0]>=RETURNED) 
    { 
     // finished this thread so it doesn't waste ALU/LD-ST/.... 
     // leaves room for other wavefronts at least 
     outputState[threadId]=NOT_STARTED; 
     return; 
    } 
    ... 
    ... 
    // ending phase of this thread 
    // localState has information if this thread needed a "return" 
    // 0=NOT_RETURNED 
    // 1=RETURNED 
    // 2=NOT_STARTED 
    lastResult=atomic_add(globalAtomicElement,localState); 
    if(lastResult>=RETURNED) 
    { 
     outputState[threadId]=AFTER_RETURNED; // you ommit 
              // this thread's result 
              // because an other thread 
              // pretends to stop all 

     // so this thread wasted cycles but dont worry, 
     // it would always waste even if you don't use 
     // a core for GCN 1.0 - GCN 3.0 architectures 
     // a core always spin within a compute unit if a 
     // core/shader is working on something. 
     // polaris architecture will have ability 
     // to shut down unused cores so that will not be 
     // a problem of power consumption either. 
    } 
    else if(lastResult==NOT_RETURNED && thisThreadReturned) 
    { 
     outputState[threadId]=RETURNED; // this is returning 
             // thread 
             //(finishing,pretending to stop all) 
    } 
    else if(lastResult==NOT_RETURNED && !thisThreadReturned) 
    { 
     outputState[threadId]=BEFORE_RETURNED; // you accept this thread's 
               // results because no thread 
               // has ever stopped and this 
               // thread surely computed 
               //everything before that 
    } 

затем в принимающем стороны, вы проверяете/фильтр в только результаты «BEFORE_RETURNED» и «RETURNED» и результаты эллиминации «AFTER_RETURNED».

В OpenCL 2.0, вы можете попробовать это:

• начало с 1 нить
• возвращение? no = spawn 2 threads, yes = stop
• рекурсивно вычислять: return? нет = респауна 2 нитки, да = остановка
• рекурсивно закончить все работы

это может сэкономить половину нитей, по крайней мере (или 1/4 или 1/8 ... или 1/N), но будет медленный, поскольку только 2 потока неэффективны.