Wie kann ich die lokalen Speicher in OpenCL nutzen?

Question

Ich habe vor kurzem mit OpenCL gespielt, und ich bin in der Lage einfachen Kernel zu schreiben, die nur die globalen Speicher verwenden. Nun würde Ich mag mit lokalen Speichern starten, aber ich kann nicht scheinen, um herauszufinden, wie get_local_size() und get_local_id() zu berechnen eine „chunk“ zu verwenden, des Ausgangs zu einem Zeitpunkt.

Zum Beispiel, sagen wir, ich wollte Beispiel konvertieren Apple OpenCL Hallo Welt Kernel etwas die uses lokalen Speicher. Wie würdest du es machen? Hier ist die Original-Kernel Quelle:

__kernel square(
    __global float *input,
    __global float *output,
    const unsigned int count)
{
    int i = get_global_id(0);
    if (i < count)
        output[i] = input[i] * input[i];
}

Wenn dieses Beispiel nicht leicht in etwas umgewandelt werden, dass zeigt, wie die Verwendung von lokalen Speichern zu machen, anderes einfaches Beispiel tun wird.

Answer 1

Überprüfen Sie die Proben in der NVIDIA oder AMD SDKs, sollten sie Sie in die richtige Richtung weisen. Matrix transponieren wäre zum Beispiel lokale Speicher verwenden.

Arbeiten mit dem Quadratur-Kernel, können Sie die Daten in einem Zwischenpufferstufe. Denken Sie daran, in den zusätzlichen Parameter zu übergeben.

__kernel square(
    __global float *input,
    __global float *output,
    __local float *temp,
    const unsigned int count)
{
    int gtid = get_global_id(0);
    int ltid = get_local_id(0);
    if (gtid < count)
    {
        temp[ltid] = input[gtid];
        // if the threads were reading data from other threads, then we would
        // want a barrier here to ensure the write completes before the read
        output[gtid] =  temp[ltid] * temp[ltid];
    }
}

Answer 2

Es gibt eine andere Möglichkeit, dies zu tun, wenn die Größe des lokalen Speichers konstant ist. Ohne die Verwendung eines Zeigers in den Kernel-Parameterliste kann der lokale Puffer innerhalb des Kernels deklariert werden nur durch Deklaration __local:

__local float localBuffer[1024];

Dieser Code entfernt aufgrund weniger clSetKernelArg Anrufe.

Answer 3

In OpenCL lokale Speicher werden in allen Arbeitspositionen in einer Arbeitsgruppe zum Austausch von Daten gemeint. Und es erfordert in der Regel eine Barriere Anruf zu tun, bevor die lokalen Speicherdaten können verwendet werden (zum Beispiel ein Workitem will einen lokalen Speicher, Daten lesen, die von den anderen Arbeitselemente geschrieben). Barriere ist teuer in der Hardware. Denken Sie daran, sollten die lokalen Speicher für wiederholte Daten verwendet werden Lese- / Schreib. Bankkonflikt sollte so weit wie möglich vermieden werden.

Wenn Sie nicht vorsichtig mit lokalen Speicher sind, können Sie mit schlechter Leistung einige Zeit am Ende als globale Speicher.