Como escolho as dimensões da grade e do bloco para os kernels CUDA?

Esta é uma questão sobre como determinar os tamanhos de grade, bloco e rosca CUDA. Esta é uma pergunta adicional àquela postada aqui .

Seguindo este link, a resposta de talonmies contém um trecho de código (veja abaixo). Não entendo o comentário "valor normalmente escolhido por ajustes e restrições de hardware".

Não encontrei uma boa explicação ou esclarecimento que explique isso na documentação CUDA. Em resumo, minha pergunta é como determinar o ideal blocksize(número de threads) dado o seguinte 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);

Essa resposta tem duas partes (eu a escrevi). Uma parte é fácil de quantificar, a outra é mais empírica.

Restrições de hardware:

Essa é a parte fácil de quantificar. O Apêndice F do guia de programação CUDA atual lista uma série de limites rígidos que limitam a quantidade de threads por bloco que uma inicialização do kernel pode ter. Se você exceder qualquer um deles, seu kernel nunca será executado. Eles podem ser resumidos aproximadamente como:

1. Cada bloco não pode ter mais de 512/1024 threads no total ( Compute Capability 1.x ou 2.xe posterior, respectivamente)
2. As dimensões máximas de cada bloco são limitadas a [512,512,64] / [1024,1024,64] (Computar 1.x / 2.x ou posterior)
3. Cada bloco não pode consumir mais de 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k no total de registros (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 bloco não pode consumir mais de 16kb / 48kb / 96kb de memória compartilhada (Compute 1.x / 2.x-6.2 / 7.0)

Se você permanecer dentro desses limites, qualquer kernel que possa compilar com êxito será iniciado sem erros.

Ajuste de desempenho:

Esta é a parte empírica. O número de threads por bloco que você escolhe dentro das restrições de hardware descritas acima pode afetar e afeta o desempenho do código em execução no hardware. O comportamento de cada código será diferente e a única maneira real de quantificá-lo é por meio de benchmarking e perfis cuidadosos. Mas, novamente, resumido de forma muito aproximada:

1. O número de threads por bloco deve ser um múltiplo redondo do tamanho do warp, que é 32 em todo o hardware atual.
2. Cada unidade de multiprocessador de streaming na GPU deve ter warps ativos suficientes para ocultar suficientemente toda a memória diferente e latência de pipeline de instrução da arquitetura e atingir o rendimento máximo. A abordagem ortodoxa aqui é tentar alcançar a ocupação ideal de hardware (a que se refere a resposta de Roger Dahl ).

O segundo ponto é um tópico enorme que duvido que alguém vá tentar cobri-lo com uma única resposta StackOverflow. Existem pessoas escrevendo teses de doutorado em torno da análise quantitativa de aspectos do problema (veja esta apresentação de Vasily Volkov da UC Berkley e este artigo de Henry Wong da Universidade de Toronto para exemplos de quão complexa a questão realmente é).

No nível de entrada, você deve estar ciente de que o tamanho do bloco que você escolher (dentro da faixa de tamanhos de bloco legais definidos pelas restrições acima) pode e tem um impacto na velocidade de execução do seu código, mas depende do hardware você tem e o código que está executando. Fazendo o benchmarking, você provavelmente descobrirá que a maioria dos códigos não triviais tem um "ponto ideal" na faixa de 128-512 threads por bloco, mas será necessário alguma análise de sua parte para descobrir onde isso está. A boa notícia é que, como você está trabalhando em múltiplos do tamanho do warp, o espaço de busca é muito finito e a melhor configuração para um determinado código é relativamente fácil de encontrar.

As respostas acima apontam como o tamanho do bloco pode afetar o desempenho e sugerem uma heurística comum para sua escolha com base na maximização da ocupação. Sem querer fornecer o critério para escolher o tamanho do bloco, vale a pena mencionar que CUDA 6.5 (agora na versão Release Candidate) inclui várias novas funções de tempo de execução para auxiliar nos cálculos de ocupação e configuração de lançamento, consulte

Dica CUDA Pro: API de ocupação simplifica a configuração de inicialização

Uma das funções úteis é cudaOccupancyMaxPotentialBlockSizecalcular heuristicamente um tamanho de bloco que atinge a ocupação máxima. Os valores fornecidos por essa função podem ser usados ​​como o ponto de partida de uma otimização manual dos parâmetros de lançamento. Abaixo está um pequeno exemplo.

#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

O cudaOccupancyMaxPotentialBlockSizeé definido no cuda_runtime.harquivo e é definido da seguinte forma:

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);
}

Os significados dos parâmetros são os seguintes

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.

Observe que, a partir do CUDA 6.5, é necessário calcular as próprias dimensões do bloco 2D / 3D a partir do tamanho do bloco 1D sugerido pela API.

Observe também que a API do driver CUDA contém APIs funcionalmente equivalentes para cálculo de ocupação, portanto, é possível usar cuOccupancyMaxPotentialBlockSizeno código da API do driver da mesma maneira mostrada para a API de tempo de execução no exemplo acima.

O tamanho do bloco é normalmente selecionado para maximizar a "ocupação". Pesquise em Ocupação CUDA para obter mais informações. Em particular, consulte a planilha Calculadora de ocupação CUDA.