¿Cómo utilizo la memoria local en OpenCL?

Question

He estado jugando con OpenCL hace poco, y yo soy capaz de escribir núcleos simples que utilizan memoria única mundial. Ahora me gustaría empezar a utilizar la memoria local, pero me parece que no puede encontrar la manera de utilizar get_local_size() y get_local_id() a uno de cómputo "trozo" de salida a la vez.

Por ejemplo, digamos que quería convertir ejemplo kernel OpenCL Hello World algo de Apple para los usos memoria local. ¿Como lo harias? Aquí está la fuente del núcleo original de:

__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];
}

Si este ejemplo no puede ser fácilmente convertida en algo que muestra cómo hacer uso de la memoria local, cualquier otro ejemplo sencillo es suficiente.

Answer 1

Vea las muestras en el NVIDIA o AMD SDK, deben apuntar en la dirección correcta. Transposición de la matriz sería utilizar la memoria local, por ejemplo.

El uso de su núcleo cuadratura, se podría organizar los datos en un buffer intermedio. Recuerde que debe pasar en el parámetro adicional.

__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

Hay otra posibilidad de hacer esto, si el tamaño de la memoria local es constante. Sin necesidad de utilizar un puntero en los núcleos de los parámetros de lista, el búfer local puede declararse dentro del kernel que acaba declarando que __local:

__local float localBuffer[1024];

Este código quita debido a menos llamadas clSetKernelArg.

Answer 3

En OpenCL memoria local está destinado a compartir datos a través de todos los elementos de trabajo en un grupo de trabajo. Y por lo general requiere hacer una llamada barrera antes de que los datos de la memoria local se puede utilizar (por ejemplo, un elemento de trabajo quiere leer un dato de la memoria local que está escrito por los demás elementos de trabajo). Barrera es costoso en hardware. Tenga en cuenta, la memoria local se debe utilizar para los datos repetidos de lectura / escritura. conflicto banco debe evitarse tanto como sea posible.

Si usted no es cuidadoso con la memoria local, usted puede terminar con un peor rendimiento algún tiempo que el uso de la memoria global.