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

Esta es una pregunta sobre cómo determinar la cuadrícula CUDA, bloques y tamaños de rosca. Esta es una pregunta adicional a la que se publica aquí:

https://stackoverflow.com/a/5643838/1292251

Siguiendo este enlace, la respuesta de talonmies contiene un fragmento de código (ver a continuación). No entiendo el comentario “valor generalmente elegido por afinación 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 tamaño de bloques óptimo (= 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<<>>mAdd(A,B,C,n); 

Por cierto, comencé mi pregunta con el enlace de arriba porque en parte responde mi primera pregunta. Si esta no es una forma adecuada de hacer preguntas sobre Stack Overflow, discúlpeme o asesíneme.

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 progtwigción actual de CUDA enumera una serie de límites estrictos que limitan la cantidad de hilos por bloque que puede tener un lanzamiento de kernel. Si supera cualquiera de estos, su núcleo nunca se ejecutará. Se pueden resumir aproximadamente como sigue:

  1. Cada bloque no puede tener más de 512/1024 subprocesos en total ( Compute Capability 1.xo 2.x y posterior respectivamente)
  2. Las dimensiones máximas de cada bloque están limitadas a [512,512,64] / [1024,1024,64] (Compute 1.x / 2.x o posterior)
  3. Cada bloque no puede consumir más de 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k de registros totales (Compute 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 16kb / 48kb / 96kb de memoria compartida (Compute 1.x / 2.x-6.2 / 7.0)

Si se mantiene dentro de esos límites, cualquier kernel que pueda comstackr exitosamente se lanzará sin error.

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 el rendimiento del código que se ejecuta en el hardware. Cómo se comporta cada código será diferente y la única manera real de cuantificarlo es mediante una comparación cuidadosa y un perfil. Pero de nuevo, muy resumido:

  1. El número de hilos por bloque debe ser un múltiplo redondo del tamaño de urdimbre, que es 32 en todo el hardware actual.
  2. Cada unidad multiprocesador de transmisión en la GPU debe tener suficientes warps activos para ocultar suficientemente toda la memoria diferente y la latencia de la tubería de instrucciones de la architecture 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 gran tema que dudo que alguien vaya a tratar de cubrir en una sola respuesta de StackOverflow. Hay personas que escriben tesis de doctorado en torno al análisis cuantitativo de aspectos del problema (ver esta presentación de Vasily Volkov de UC Berkley y este documento de Henry Wong de la Universidad de Toronto para ejemplos de cuán compleja es en realidad la pregunta).

En el nivel de entrada, debes saber que el tamaño de bloque que elijas (dentro del rango de tamaños de bloques legales definidos por las restricciones anteriores) puede tener un impacto en la velocidad de ejecución del código, pero depende del hardware usted tiene y el código que está ejecutando. Al hacer un benchmarking, probablemente descubras que la mayoría del código no trivial tiene un “punto óptimo” en el rango de 128-512 hilos por bloque, pero requerirá un análisis de tu parte para encontrar dónde está. La buena noticia es que, debido a que está trabajando en múltiplos del tamaño de la urdimbre, el espacio de búsqueda es muy finito y la mejor configuración para una pieza determinada de código 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 en función de la maximización de la ocupación. Sin querer proporcionar el criterio para elegir el tamaño de bloque, vale la pena mencionar que CUDA 6.5 (ahora en la versión Candidate Release) incluye varias nuevas funciones de tiempo de ejecución para ayudar en los cálculos de ocupación y la configuración de inicio, ver

Sugerencia para CUDA Pro: Occupancy API simplifica la configuración de inicio

Una de las funciones útiles es cudaOccupancyMaxPotentialBlockSize que calcula heurísticamente un tamaño de bloque que alcanza la ocupación máxima. Los valores proporcionados por esa función podrían usarse entonces como el punto de partida de una optimización manual de los parámetros de inicio. A continuación hay un pequeño ejemplo.

 #include  /************************/ /* 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>>(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 

EDITAR

El cudaOccupancyMaxPotentialBlockSize se define en el archivo cuda_runtime.h y se define de la siguiente manera:

 template __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, uno necesita calcular las dimensiones propias de un bloque 2D / 3D del tamaño de bloque 1D sugerido por la API.

Tenga en cuenta también que la API de controlador CUDA contiene API funcionalmente equivalentes para el cálculo de ocupación, por lo que es posible usar cuOccupancyMaxPotentialBlockSize en el código de la 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 de bloques se suele seleccionar para maximizar la “ocupación”. Busque en CUDA Occupancy para más información. En particular, vea la hoja de cálculo Calculadora de ocupación CUDA.