Compreendendo as dimensões da grade CUDA, dimensões de bloco e organização de threads (explicação simples) [fechado]

Como os threads são organizados para serem executados por uma GPU?

Hardware

Se um dispositivo GPU tiver, por exemplo, 4 unidades de multiprocessamento e eles puderem executar 768 threads cada: então, em um determinado momento, não mais do que 4 * 768 threads estarão em execução paralelamente (se você planejou mais threads, eles estarão aguardando a sua vez).

Programas

threads são organizados em blocos. Um bloco é executado por uma unidade de multiprocessamento. Os encadeamentos de um bloco podem ser identificados (indexados) usando os índices 1Dimension (x), 2Dimensions (x, y) ou 3Dim (x, y, z), mas em qualquer caso x y z <= 768 para o nosso exemplo (outras restrições se aplicam para x, y, z, consulte o guia e a capacidade do seu dispositivo).

Obviamente, se você precisar de mais do que os threads 4 * 768, precisará de mais de 4 blocos. Os blocos também podem ser indexados em 1D, 2D ou 3D. Há uma fila de blocos aguardando para entrar na GPU (porque, em nosso exemplo, a GPU possui 4 multiprocessadores e apenas 4 blocos estão sendo executados simultaneamente).

Agora, um caso simples: processando uma imagem de 512x512

Suponha que queremos que um thread processe um pixel (i, j).

Podemos usar blocos de 64 threads cada. Então precisamos de 512 * 512/64 = 4096 blocos (para ter segmentos de 512x512 = 4096 * 64)

É comum organizar (para facilitar a indexação da imagem) os encadeamentos em blocos 2D com blockDim = 8 x 8 (os 64 encadeamentos por bloco). Eu prefiro chamá-lo threadsPerBlock.

dim3 threadsPerBlock(8, 8);  // 64 threads

e gridDim 2D = 64 x 64 blocos (são necessários os 4096 blocos). Eu prefiro chamá-lo numBlocks.

dim3 numBlocks(imageWidth/threadsPerBlock.x,  /* for instance 512/8 = 64*/
              imageHeight/threadsPerBlock.y); 

O kernel é lançado assim:

myKernel <<<numBlocks,threadsPerBlock>>>( /* params for the kernel function */ );       

Finalmente: haverá algo como "uma fila de 4096 blocos", em que um bloco aguarda a atribuição de um dos multiprocessadores da GPU para executar seus 64 threads.

No kernel, o pixel (i, j) a ser processado por um thread é calculado da seguinte maneira:

uint i = (blockIdx.x * blockDim.x) + threadIdx.x;
uint j = (blockIdx.y * blockDim.y) + threadIdx.y;

Suponha uma GPU 9800GT:

• possui 14 multiprocessadores (SM)
• cada SM possui 8 processadores de threads (AKA stream-processadores, SP ou núcleos)
• permite até 512 threads por bloco
• warpsize é 32 (o que significa que cada um dos processadores de thread de 14x8 = 112 pode agendar até 32 threads)

https://www.tutorialspoint.com/cuda/cuda_threads.htm

Um bloco não pode ter threads mais ativos que 512, portanto, __syncthreadspode sincronizar apenas um número limitado de threads. ou seja, se você executar o seguinte com 600 threads:

func1();
__syncthreads();
func2();
__syncthreads();

então o kernel deve rodar duas vezes e a ordem de execução será:

1. func1 é executado para os primeiros 512 threads
2. func2 é executado para os primeiros 512 threads
3. func1 é executado para os threads restantes
4. func2 é executado para os threads restantes

Nota:

O ponto principal __syncthreadsé uma operação de todo o bloco e não sincroniza todos os threads.

Não tenho certeza sobre o número exato de threads que __syncthreadspodem ser sincronizados, pois você pode criar um bloco com mais de 512 threads e deixar o warp manipular o agendamento. No meu entender, é mais preciso dizer: func1 é executado pelo menos nos primeiros 512 threads.

Antes de editar esta resposta (em 2010), medi os threads de 14x8x32 que eram sincronizados usando __syncthreads.

Eu apreciaria muito se alguém testasse isso novamente para obter uma informação mais precisa.