Confundido acerca de perfiles de los resultados de mi kernel OpenCL (DSP ventana de función)

Question

Me completó un núcleo ventana de funciones en OpenCL. Básicamente una función de ventana solo se aplica un conjunto de coeficientes sobre otro conjunto de números de pieza a pieza (Wikipedia lo explica mejor). Yo era capaz de rellenar la matriz coeficiente de ventana flotante en la memoria caché constante para la mayoría de los casos.

que esperaba mis resultados de Calcular el profesor para demostrar que el anfitrión al dispositivo y el dispositivo de transferencia de la memoria del host se necesitarían más de 95% del tiempo de procesamiento. Para casi todos mis casos es sólo el 80% del tiempo de procesamiento. Estoy escribiendo y leyendo una gama de 4,2 millones de flotación hacia y desde la pizarra y escribir otra red de flotadores que por lo general se mantiene muy por debajo de un millón.

¿Hay algo a pescado en el aspecto del núcleo? Cualquier opinión sobre si se trata de un problema que debe correr más rápido en una GPU de una CPU en el primer lugar (todavía no estoy 100% en esto). Estoy un poco aturdido por qué mi gld_efficiency y gst_efficiency vuelo estacionario entre 0,1 y 0,2. Hice este núcleo con G80 memoria global de coalescencia en mente. Mi memoria global en general el rendimiento parece bien en 40gbs. El núcleo es bastante simple y se publica a continuación.

__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

El número de threads (el término es OpenCL elementos de trabajo, por cierto) está usando? Se necesita al menos algo en los cientos de cargar una gran GPU de manera eficiente.

Usted dice que quiere hacer uso de acceso a la memoria coalescencia, pero una carga con un desplazamiento como

int inputArrayIndex = (gid*primitivesPerDataFrame)+i;

no va a hacer que esto sea posible en la mayoría de los casos. G80 de Nvidia tiene restricciones muy severas cuando se trata de coalescencia, consulte la "Guía de Buenas Prácticas de OpenCL" para más información. Básicamente, el trabajo a elementos de una urdimbre tienen que elementos de acceso de un bloque alineado 64 o 128 bytes de una cierta manera al mismo tiempo para hacer cargas y tiendas suceden fundieron.

O para darle un ejemplo: si es primitivesPerDataFrame 16, cargas y almacenamientos de una urdimbre se realizan en desplazamientos espaciados 16 elementos separados, por lo que cualquier eficiente coalescencia imposible

.