Como os blocos CUDA são divididos em warps?

Se eu iniciar meu kernel com uma grade cujos blocos tenham dimensões:

dim3 block_dims(16,16); 

Como os blocos de grade agora são divididos em deformações? As duas primeiras linhas desse bloco formam uma distorção, ou as duas primeiras colunas, ou isso é arbitrariamente ordenado?

Assuma uma capacidade de computação de GPU de 2,0.

Os threads são numerados em ordem dentro dos blocos, de modo que threadIdx.x varie o mais rápido, em seguida, threadIdx.y o segundo mais rápido variando e threadIdx.z o mais lento, variando. Isso é funcionalmente o mesmo que ordenar a coluna principal em matrizes multidimensionais. Warps são sequencialmente construídos a partir de threads nesta ordem. Então, o cálculo para um bloco de 2d é

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

Isso é abordado no guia de programação e no guia PTX.