Kann diese Parallelität in OpenCL implementiert werden

Question

Dies ist mein erster Beitrag. Ich werde versuchen, es kurz zu halten, weil ich deine Zeit schätze. Diese Gemeinschaft war unglaublich für mich.

Ich lerne OpenCL und möchte ein wenig Parallelität aus dem folgenden Algorithmus extrahieren. Ich werde dir nur den Teil zeigen, an dem ich arbeite, den ich auch vereinfacht habe, soweit ich kann.

1) Eingänge: Zwei 1D-Arrays der Länge (n): A, B und Wert von n. Auch Werte C [0], D [0].

2) Ausgänge: Zwei 1D-Arrays der Länge (n): C, D.

C[i] = function1(C[i-1]) 
D[i] = function2(C[i-1],D[i-1]) 

So sind diese rekursive Definitionen, jedoch ist die Berechnung von C & D i für einen gegebenen Wert kann in geschehen parallel (sie sind offensichtlich komplizierter, um Sinn zu ergeben). Ein naiver Gedanke wäre zwei Arbeitselemente für die folgende Kernel erschaffen:

__kernel void test (__global float* A, __global float* B, __global float* C, 
        __global float* D, int n, float C0, float D0) { 
    int i, j=get_global_id(0); 

    if (j==0) { 
     C[0] = C0; 
     for (i=1;i<=n-1;i++) { 
      C[i] = function1(C[i-1]); 
      [WAIT FOR W.I. 1 TO FINISH CALCULATING D[i]]; 
     } 
     return; 
    } 
    else { 
     D[0] = D0; 
     for (i=1;i<=n-1;i++) { 
      D[i] = function2(C[i-1],D[i-1]); 
      [WAIT FOR W.I. 0 TO FINISH CALCULATING C[i]]; 
     } 
     return; 
    } 
} 

würde im Idealfall einen ersten Vergleich tut jeder der beiden Arbeitspositionen (Nummern 0,1) und dann ihre jeweilige Schleife eingeben, für jede Iteration zu synchronisieren . Angesichts der SIMD-Implementierung von GPUs gehe ich davon aus, dass dies NICHT funktionieren wird (Arbeitsaufgaben würden auf den gesamten Kernel-Code warten), aber ist es möglich, diese Art von Arbeit zwei CPU-Kernen zuzuordnen und funktioniert es wie erwartet? Was wird die Barriere in diesem Fall sein?

Answer 1

Abhängigkeit in Ihrem Fall ist vollständig linear/rekursiv (ich brauche i-1). Nicht einmal logarithmisch wie andere Probleme (Reduktion, Summe, Sortierung etc.). Und deshalb passt dieses Problem nicht gut in ein SIMD-Gerät.

Das Beste, was Sie tun können, ist ein 2-Threads-Ansatz in der CPU. Thread 1 wird "produzieren" Daten (C-Wert) für das Gewinde 2.

Ein sehr naiver Ansatz zum Beispiel:

Thread 1: 
for(){ 
    ProcessC(i); 
    atomic_inc(counter); //This function should unlock 
} 

Thread 2: 
for(){ 
    atomic_dec(counter); //This function should lock 
    ProcessD(i); 
} 

Wo atomic_incatomic_dec und können mit der Zählung von Semaphoren beispielsweise implementiert werden.

Answer 2

Dies kann in opencl implementiert werden, aber wie die andere Antwort sagt, werden Sie auf maximal 2 Threads begrenzt sein.

Meine Version Ihrer Funktion sollte mit einer einzigen Arbeitsgruppe mit zwei Arbeitsaufgaben aufgerufen werden.

__kernel void test (__global float* A, __global float* B, __global float* C, __global float* D, int n, float C0, float D0) 
{ 
    int i; 
    int gid = get_global_id(0); 

    local float prevC; 
    local float prevD; 

    if (gid == 0) { 
     C[0] = prevC = C0; 
     D[0] = prevD = D0; 
    } 

    barrier(CLK_LOCAL_MEM_FENCE); 

    for (i=1;i<=n-1;i++) { 
     if(gid == 0){ 
      C[i] = function1(prevC); 
     }else if (gid == 1){ 
      D[i] = function2(prevC, prevD); 
     } 

     barrier(CLK_LOCAL_MEM_FENCE); 
     prevC = C[i]; 
     prevD = D[i]; 
    } 
} 

Dies sollte auf jeder opencl-Hardware ausgeführt werden. Wenn Sie nicht alle C- und D-Werte speichern möchten, können Sie prevC und prevD einfach in zwei Gleitkommazahlen anstelle der gesamten Liste zurückgeben. Dies würde es auch viel schneller machen, da es bei jedem Lesen und Schreiben der Zwischenwerte auf einer niedrigeren Cache-Ebene (dh einem lokalen Speicher) bleibt. Die lokale Speicherverstärkung sollte auch für alle Opencl-Hardware gelten.

Gibt es also einen Punkt auf einer GPU zu laufen? Nicht für die Parallelität. Sie sind mit 2 Threads festgefahren. Aber wenn Sie nicht alle zurückgegebenen Werte von C und D benötigen, würden Sie aufgrund des viel schnelleren Speicherbedarfs von GPUs wahrscheinlich eine deutliche Beschleunigung sehen.

All dies setzt voraus, dass function1 und function2 nicht übermäßig komplex sind. Wenn das der Fall ist, bleiben Sie einfach bei CPUs - und wahrscheinlich einer anderen Multiprocessing-Technik wie OpenMP.