¿Cómo elijo las dimensiones de cuadrícula y bloque para los núcleos CUDA?

Esta es una pregunta sobre cómo determinar la cuadrícula CUDA, los tamaños de bloque y de rosca. Esta es una pregunta adicional a la publicada aquí .

Siguiendo este enlace, la respuesta de talonmies contiene un fragmento de código (ver más abajo). No entiendo el comentario "valor generalmente elegido por ajuste y restricciones de hardware".

No he encontrado una buena explicación o aclaración que explique esto en la documentación de CUDA. En resumen, mi pregunta es cómo determinar el óptimo blocksize(número de subprocesos) dado el siguiente código:

const int n = 128 * 1024;
int blocksize = 512; // value usually chosen by tuning and hardware constraints
int nblocks = n / nthreads; // value determine by block size and total work
madd<<<nblocks,blocksize>>>mAdd(A,B,C,n);

Hay dos partes en esa respuesta (la escribí). Una parte es fácil de cuantificar, la otra es más empírica.

Restricciones de hardware:

Esta es la parte fácil de cuantificar. El Apéndice F de la guía de programación actual de CUDA enumera una serie de límites estrictos que limitan la cantidad de subprocesos por bloque que puede tener un lanzamiento del kernel. Si excede cualquiera de estos, su kernel nunca se ejecutará. Se pueden resumir a grandes rasgos como:

1. Cada bloque no puede tener más de 512/1024 subprocesos en total ( capacidad de cómputo 1.xo 2.xy posterior respectivamente)
2. Las dimensiones máximas de cada bloque están limitadas a [512,512,64] / [1024,1024,64] (Compute 1.x / 2.xo posterior)
3. Cada bloque no puede consumir más de 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k registros en total (Calcule 1.0,1.1 / 1.2,1.3 / 2.x- / 3.0 / 3.2 / 3.5-5.2 / 5.3 / 6-6.1 / 6.2 / 7.0)
4. Cada bloque no puede consumir más de 16 kb / 48 kb / 96 kb de memoria compartida (Compute 1.x / 2.x-6.2 / 7.0)

Si permanece dentro de esos límites, cualquier kernel que pueda compilar con éxito se ejecutará sin errores.

La optimización del rendimiento:

Esta es la parte empírica. El número de subprocesos por bloque que elija dentro de las restricciones de hardware descritas anteriormente puede afectar y afecta el rendimiento del código que se ejecuta en el hardware. El comportamiento de cada código será diferente y la única forma real de cuantificarlo es mediante una evaluación comparativa y un perfil cuidadosos. Pero de nuevo, muy resumido:

1. El número de hilos por bloque debe ser un múltiplo redondo del tamaño de la deformación, que es 32 en todo el hardware actual.
2. Cada unidad de multiprocesador de transmisión en la GPU debe tener suficientes deformaciones activas para ocultar suficientemente todas las diferentes memorias y la latencia de la canalización de instrucciones de la arquitectura y lograr el máximo rendimiento. El enfoque ortodoxo aquí es intentar lograr una ocupación óptima del hardware (a lo que se refiere la respuesta de Roger Dahl ).

El segundo punto es un tema enorme que dudo que alguien intente cubrirlo en una sola respuesta de StackOverflow. Hay personas que escriben tesis de doctorado en torno al análisis cuantitativo de aspectos del problema (consulte esta presentación de Vasily Volkov de UC Berkley y este artículo de Henry Wong de la Universidad de Toronto para ver ejemplos de cuán compleja es realmente la pregunta).

En el nivel de entrada, debe tener en cuenta que el tamaño de bloque que elija (dentro del rango de tamaños de bloque legales definidos por las restricciones anteriores) puede tener y tiene un impacto en la rapidez con que se ejecutará su código, pero depende del hardware que tiene y el código que está ejecutando. Mediante la evaluación comparativa, probablemente encontrará que la mayoría del código no trivial tiene un "punto óptimo" en el rango de 128-512 subprocesos por bloque, pero requerirá un análisis de su parte para encontrar dónde está. La buena noticia es que debido a que está trabajando en múltiplos del tamaño de la deformación, el espacio de búsqueda es muy finito y la mejor configuración para un fragmento de código dado es relativamente fácil de encontrar.

Las respuestas anteriores señalan cómo el tamaño del bloque puede afectar el rendimiento y sugieren una heurística común para su elección basada en la maximización de la ocupación. Sin querer proporcionar el criterio para elegir el tamaño del bloque, vale la pena mencionar que CUDA 6.5 (ahora en versión Release Candidate) incluye varias funciones de tiempo de ejecución nuevas para ayudar en los cálculos de ocupación y configuración de lanzamiento, ver

Consejo profesional de CUDA: la API de ocupación simplifica la configuración de lanzamiento

Una de las funciones útiles es la cudaOccupancyMaxPotentialBlockSizeque calcula heurísticamente un tamaño de bloque que logra la máxima ocupación. Los valores proporcionados por esa función podrían utilizarse como punto de partida de una optimización manual de los parámetros de lanzamiento. A continuación se muestra un pequeño ejemplo.

#include <stdio.h>

/************************/
/* TEST KERNEL FUNCTION */
/************************/
__global__ void MyKernel(int *a, int *b, int *c, int N) 
{ 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 

    if (idx < N) { c[idx] = a[idx] + b[idx]; } 
} 

/********/
/* MAIN */
/********/
void main() 
{ 
    const int N = 1000000;

    int blockSize;      // The launch configurator returned block size 
    int minGridSize;    // The minimum grid size needed to achieve the maximum occupancy for a full device launch 
    int gridSize;       // The actual grid size needed, based on input size 

    int* h_vec1 = (int*) malloc(N*sizeof(int));
    int* h_vec2 = (int*) malloc(N*sizeof(int));
    int* h_vec3 = (int*) malloc(N*sizeof(int));
    int* h_vec4 = (int*) malloc(N*sizeof(int));

    int* d_vec1; cudaMalloc((void**)&d_vec1, N*sizeof(int));
    int* d_vec2; cudaMalloc((void**)&d_vec2, N*sizeof(int));
    int* d_vec3; cudaMalloc((void**)&d_vec3, N*sizeof(int));

    for (int i=0; i<N; i++) {
        h_vec1[i] = 10;
        h_vec2[i] = 20;
        h_vec4[i] = h_vec1[i] + h_vec2[i];
    }

    cudaMemcpy(d_vec1, h_vec1, N*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_vec2, h_vec2, N*sizeof(int), cudaMemcpyHostToDevice);

    float time;
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);

    cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, MyKernel, 0, N); 

    // Round up according to array size 
    gridSize = (N + blockSize - 1) / blockSize; 

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Occupancy calculator elapsed time:  %3.3f ms \n", time);

    cudaEventRecord(start, 0);

    MyKernel<<<gridSize, blockSize>>>(d_vec1, d_vec2, d_vec3, N); 

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Kernel elapsed time:  %3.3f ms \n", time);

    printf("Blocksize %i\n", blockSize);

    cudaMemcpy(h_vec3, d_vec3, N*sizeof(int), cudaMemcpyDeviceToHost);

    for (int i=0; i<N; i++) {
        if (h_vec3[i] != h_vec4[i]) { printf("Error at i = %i! Host = %i; Device = %i\n", i, h_vec4[i], h_vec3[i]); return; };
    }

    printf("Test passed\n");

}

EDITAR

El cudaOccupancyMaxPotentialBlockSizese define en el cuda_runtime.harchivo y se define de la siguiente manera:

template<class T>
__inline__ __host__ CUDART_DEVICE cudaError_t cudaOccupancyMaxPotentialBlockSize(
    int    *minGridSize,
    int    *blockSize,
    T       func,
    size_t  dynamicSMemSize = 0,
    int     blockSizeLimit = 0)
{
    return cudaOccupancyMaxPotentialBlockSizeVariableSMem(minGridSize, blockSize, func, __cudaOccupancyB2DHelper(dynamicSMemSize), blockSizeLimit);
}

El significado de los parámetros es el siguiente

minGridSize     = Suggested min grid size to achieve a full machine launch.
blockSize       = Suggested block size to achieve maximum occupancy.
func            = Kernel function.
dynamicSMemSize = Size of dynamically allocated shared memory. Of course, it is known at runtime before any kernel launch. The size of the statically allocated shared memory is not needed as it is inferred by the properties of func.
blockSizeLimit  = Maximum size for each block. In the case of 1D kernels, it can coincide with the number of input elements.

Tenga en cuenta que, a partir de CUDA 6.5, es necesario calcular las propias dimensiones del bloque 2D / 3D a partir del tamaño de bloque 1D sugerido por la API.

Tenga en cuenta también que la API del controlador CUDA contiene API funcionalmente equivalentes para el cálculo de ocupación, por lo que es posible utilizarlo cuOccupancyMaxPotentialBlockSizeen el código API del controlador de la misma manera que se muestra para la API de tiempo de ejecución en el ejemplo anterior.

El tamaño del bloque generalmente se selecciona para maximizar la "ocupación". Busque en CUDA Occupancy para obtener más información. En particular, consulte la hoja de cálculo de la Calculadora de ocupación de CUDA.