Confuso su profilazione risultati dal mio kernel OpenCL (DSP finestra Function)

Question

Ho completato un kernel Finestra funzione in OpenCL. In sostanza una funzione di finestra proprio applica una serie di coefficienti rispetto ad un altro insieme di numeri pezzo per pezzo (Wikipedia spiega meglio). Sono stato in grado di roba matrice finestra coefficiente galleggiante nella cache costante per la maggior parte dei casi.

mi aspettavo i miei risultati da Compute Prof per dimostrare che l'host per dispositivo e dispositivo per trasferimenti di memoria ospitanti avrebbe preso oltre il 95% del tempo di elaborazione. Per quasi tutti i miei casi è solo l'80% del tempo di elaborazione. Sto scrivendo e leggendo un array mono 4,2 milioni galleggiante e dalla scheda e scrivere un altro array galleggiante che in genere rimane ben al di sotto di un milione.

nulla

fa nel pesce kernel sguardo? Eventuali pareri su se si tratta di un problema che dovrebbe correre più veloce su una GPU di una CPU, in primo luogo (io sono ancora al 100% su questo). Sono un po 'stordito sul motivo per cui il mio gld_efficiency e gst_efficiency hover tra 0,1 e 0,2. Ho fatto questo kernel con G80 memoria globale coalescenza in mente. La mia memoria globale complessivo sembra un throughput bene a 40gbs. Il kernel è abbastanza semplice ed è pubblicato qui di seguito.

__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

Quanti fili (il termine OpenCL è work-elementi, tra l'altro) stai usando? Avete bisogno di almeno qualcosa nell'ordine delle centinaia di caricare un grande GPU in modo efficiente.

Tu dici che vuoi fare uso di accesso alla memoria coalescenza, ma un carico con un offset come

int inputArrayIndex = (gid*primitivesPerDataFrame)+i;

Non lo renderà possibile nella maggior parte dei casi. G80 di Nvidia ha abbastanza severe restrizioni quando si tratta di coalescenza, vedere la "OpenCL Guida Pratica" per ulteriori informazioni. Fondamentalmente, work-elementi da un ordito devono elementi di accesso di un blocco allineato 64 o 128 byte in un certo modo, allo stesso tempo per fare carichi e negozi accadono coalesced.

O, per fare un esempio: se primitivesPerDataFrame è 16, carichi e negozi di un ordito sono fatte a offset distanziati 16 elementi a parte, fare alcun efficiente coalescenza impossibile

.