Verwirrt über Profilierungs Ergebnisse von meinem OpenCL-Kernel (DSP Fensterfunktion)

Question

absolvierte ich eine Fensterfunktion Kernel in OpenCL. Im Grunde eine Fensterfunktion gilt nur einen Satz von Koeffizienten über einen anderen Satz von Zahlen Stück für Stück (Wikipedia erklärt es besser). Ich konnte in den meisten Fällen den Fensterkoeffizienten float-Array in konstanten Cache stopfen.

erwartete ich meine Ergebnisse von Compute Prof zu zeigen, dass die Host-Gerät und Gerät zu Host-Speicherübertragungen würde mehr als 95% der Bearbeitungszeit. Für fast alle meine Fällen ist es nur 80% der Bearbeitungszeit. Ich schreibe und einen 4.200.000 Schwimmer Array Lese aus dem Brett und anderen Schwimmer Array zu schreiben, dass im Allgemeinen deutlich unter einer Million bleibt.

Gibt es etwas in den Kernel Blick fischig? Alle Meinungen auf, wenn es ein Problem, das auf einer GPU als eine CPU in erster Linie schneller laufen sollte (ich bin immer noch nicht zu 100% auf dieser). Ich bin ein wenig erstaunt, warum mein gld_efficiency und gst_efficiency schwebt zwischen 0,1 und 0,2. Ich habe diesen Kernel mit G80 globalem Speicher im Auge Koaleszenz. Mein globaler Speicher scheint Gesamtdurchsatz in Ordnung bei 40gbs. Der Kernel ist ziemlich einfach und ist unten geschrieben.

__kernel void window(__global float* inputArray, // first frame to ingest starts at 0.  Sized to nFramesToIngest*framesize samples
    __constant float* windowArray, // may already be partly filled
    int windowSize, // size of window frame, in floats
    int primitivesPerDataFrame, //amount of primitives in each frame of inputArray parameter
    int nInFramesThisCall, //each thread solves a frame, so this integer represent how many threads this kernel launches
    int isRealNumbers //0 for complex, non-zero for real 
)
{
int gid = get_global_id(0) + get_global_size(0) * get_global_id(1);

if(gid < nInFramesThisCall) //make sure we don't execute unnecessary threads
{
    if(isRealNumbers)
    {
        for(int i = 0; i < primitivesPerDataFrame; i++)
        {
            int inputArrayIndex = (gid*primitivesPerDataFrame)+i;
            inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize];
        }
    }
    else //complex
    {
        for(int i = 0; i < primitivesPerDataFrame; i++)
        {
            int inputArrayIndex = (gid*primitivesPerDataFrame)+i;
            inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize/2];
        }
    }
}

}

Answer 1

Wie viele Threads (der OpenCL Begriff ist arbeits Artikel, übrigens) verwenden Sie? Sie müssen zumindest etwas in den Hunderten effizient eine große GPU laden.

Sie sagen, Sie Verwendung von verschmolzenen Speicherzugriff machen wollen, aber eine Last mit einem Versatz wie

int inputArrayIndex = (gid*primitivesPerDataFrame)+i;

wird nicht machen es möglich in den meisten Fällen. NVidia G80 hat ziemlich strenge Beschränkungen, wenn es um Koaleszenz kommt, finden Sie in der „OpenCL Best Practices Guide“ für weitere Informationen. Grundsätzlich Arbeits Elemente von einem Kett haben Zugriff Elemente eines 64- oder 128-Byte-Block ausgerichtet in einer bestimmten Art und Weise in der gleichen Zeit zu machen, Lade- und Speicher koalesziert passiert.

Oder Ihnen ein Beispiel geben: Wenn primitivesPerDataFrame 16 ist, Ladungen und Speicherungen einer Kette an Versetzungen erfolgen im Abstand 16 Elemente auseinander, jeden effiziente Koaleszenz unmöglich machen

.