¿Cómo se dividen los bloques CUDA en warps?

Si comienzo mi kernel con una grilla cuyos bloques tienen dimensiones:

dim3 block_dims(16,16); 

¿Cómo se dividen ahora los bloques de malla en warps? ¿Las primeras dos filas de dicho bloque forman una urdimbre, o las dos primeras columnas, o se ordena arbitrariamente?

Supongamos una capacidad de cálculo GPU de 2.0.

Los subprocesos están numerados en orden dentro de los bloques para que threadIdx.x varíe más rápido, luego threadIdx.y el segundo más rápido que varía, y threadIdx.z el más lento que varía. Esto es funcionalmente el mismo que el ordenamiento de columnas principales en matrices multidimensionales. Las deformaciones se construyen secuencialmente a partir de hilos en este orden. Entonces, el cálculo para un 2do bloque es

 unsigned int tid = threadIdx.x + threadIdx.y * blockDim.x; unsigned int warpid = tid / warpSize; 

Esto está cubierto tanto en la guía de progtwigción como en la guía PTX.