OpenCL 1.2: mem_fence() oder barrier() oder beide

Question

Ich habe gerade OpenCL C-Programmierung gestartet. Alle Arbeitselemente einer Arbeitsgruppe aktualisieren eindeutige Speicherorte des lokalen Speichers. Später wird eine private Variable eines Arbeitselements basierend auf lokalen Daten aktualisiert, die von zwei anderen Arbeitselementen aktualisiert werden. Etwas wie folgt aus:

__kernel MyKernel(__global int *in_ptr) 
     { 
      /* Define a variable in private address space */ 
      int priv_data; 
      /* Define two indices in private address space */ 
      int index1, index2; 

      /* index1 and index2 are legitimate local work group indices */ 
      index1 = SOME_CORRECT_VALUE; 
      index2 = ANOTHER_CORRECT_VALUE; 

      /* Define storage in local memory large enough to cater to all work items of this work group */ 
      __local int tempPtr[WORK_GROUP_SIZE]; 
      tempPtr[get_local_id(0)] = SOME_RANDOM_VALUE; 

      /* Do not proceed until the update of tempPtr by this WI has completed */ 
      mem_fence(CLK_LOCAL_MEM_FENCE); 

      /* Do not proceed until all WI of this WG have updated tempPtr */ 
      barrier(CLK_LOCAL_MEM_FENCE); 

      /* Update private data */ 
      priv_data = tempPtr[index1] + tempPtr[index2]; 
     } 

Obwohl das Snippet oben ist konservativ, hätte nicht barriere den Job zu erledigen, wie es intern Fechten tut?

Answer 1

Ja, Barriere bereits Fechten.

Eine Barriere synchronisiert die Ausführung in diesem Punkt. Daher müssen alle vorherigen Anweisungen ausgeführt werden, daher ist der Speicher an diesem Punkt konsistent. Ein Zaun wird nur sicherstellen, dass alle Lese-/Schreibvorgänge abgeschlossen sind, bevor weitere Lese-/Schreibvorgänge ausgeführt werden, aber die Arbeiter führen möglicherweise unterschiedliche Anweisungen aus.

In einigen Fällen können Sie mit einem einzigen Zaun gehen. Wenn Sie sich nicht darum kümmern, dass lokale Mitarbeiter nicht mehr synchron sind, und Sie möchten, dass die vorherigen Schreib-/Lesevorgänge abgeschlossen werden. In deinem Fall würde ein Zaun ausreichen. (es sei denn, dieser Code läuft in einer Schleife und es gibt zusätzlichen Code, den Sie nicht in das Beispiel eingegeben haben).