Brechen und Fortfahren im OpenCL-Kernel

Question

Wie implementiert man diese im OpenCL-Kernel? Ist "Rückkehr" entspricht "brechen"?

Ich bin mit OpenCL 1.2

Ich mag dies für Schleifen, die Sie durch einen typedef struct verschachtelte Arrays verschachtelt mit 3 implementieren.

EDIT

Realisiert Ich brauche einige Code zu erhalten besser verstanden

in kernel ...

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; 
    } 
} 

IN HOST zu sein (der Code, den ich im Kernel zu replizieren bin versucht, für Effizienz)

  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

Das Abbrechen aller anderen Threads von opencl würde zu undefiniertem Verhalten führen, da viele von ihnen mitten im Schreiben/Lesen von globalem/lokalem Speicher sein könnten und das den laufenden Thread löschen könnte (alle anderen/verbleibenden Kernel/Threads stoppen). Wahrscheinlich gibt es so etwas in opencl nicht.

Aber Sie können ein Ausgabe-Array hinzufügen, dass jeder Thread seinen letzten Zustand schreibt. Wenn ein Element einen "Return" -Code hat, sollten Sie nach "after_return" -Code suchen, um die Ergebnisse dieser Ergebnisse zu löschen und "before_return" -Einträge zu akzeptieren. Dies würde auch atomare Operationen in der Ausgangsstufe erfordern, wird also langsamer, was schlecht ist.

Aber man von einzelnen Kernen sicher zurückkehren kann:

Below Code gut und früh verlassen (beendet Kernel-Ausführung für einige Threads aber nicht alle), zusammengestellt von der ‚Rückkehr‘ auf einer HD7870 und einen R7-240 ohne Fehler, weil 'return' keine der Einschränkungen ist, die von OpenCL angewendet werden.

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

      ... 
      } 

Verwendet Open 1.2 Header von C++.

Aber, wenn Sie noch fake-Rückkehr und einen Faden braucht nicht beeinflussen andere Threads' Ausgänge/Eingänge, dann so etwas wie dies helfen würde:

// 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 
    } 

dann in Host-Seite, Sie überprüfen/Filter-in nur Ergebnisse von "BEFORE_RETURNED" und "RETURNED" und Ergebnisse von "AFTER_RETURNED".

In OpenCL 2.0 können Sie dies versuchen:

• Start mit 1 Faden
• Rückkehr? nein = spawn 2 threads, ja = stop
• rekursiv berechnen: zurück? no = Spawn 2 Fäden, ja = Stopp
• rekursiv alle Arbeiten fertig

könnte dies die Hälfte der Fäden mindestens (oder 1/4 oder 1/8 ... oder 1/N) sparen, sondern wäre langsam, da nur 2 Threads ineffizient sind.