Como determinar automaticamente o tamanho do bloco CUDA e o tamanho da grade para um array 2D?

Question

como determinar o tamanho do bloco e tamanho da grade automaticamente para array 2D (por exemplo, processamento de imagem) em CUDA?

CUDA tem cudaOccupancyMaxPotentialBlockSize()função para calcular o tamanho do bloco para funções do kernel cuda automaticamente. veja aqui . Nesse caso, funciona bem para array 1D.

Para o meu caso, tenho uma imagem de 640x480.

Como determinar o tamanho do bloco / grade? Eu uso:

////image size: 640x480


int x_min_grid_size, x_grid_size, x_block_size;
int y_min_grid_size, y_grid_size, y_block_size;

cudaOccupancyMaxPotentialBlockSize
(
    &x_min_grid_size, &x_block_size,
    my_cuda_kernel,
    0, image.width()
);
cudaOccupancyMaxPotentialBlockSize
(
    &y_min_grid_size, &y_block_size,
    my_cuda_kernel,
    0, image.height()
);

x_grid_size = (image.width()  + x_block_size - 1) / x_block_size;
y_grid_size = (image.height() + y_block_size - 1) / y_block_size;

dim3 grid_dim(x_grid_size, y_grid_size);
dim3 block_dim(x_block_size, y_block_size);

my_cuda_kernel<<<grid_dim, block_dim>>>(<arguments...>)

////check cuda kernel function launch error
cudaError_t error = cudaGetLastError();
if(cudaSuccess != error)
{
    std::cout<<"CUDA Error! "<<cudaGetErrorString(error)<<std::endl;
    exit(1);
}
cudaDeviceSynchronize();

Pergunta 1 Posso calcular o tamanho do bloco / grade usando este método?

Para este código, recebi um erro após o lançamento da função kernel.

CUDA Error! invalid configuration arguments

Se eu configurar x_block_size = 32; y_block_size = 32manualmente, funciona e não tem erro.

Posso perguntar por que CUDA recebe invalid configuration argumentsuma mensagem de erro? Parece que não posso usar cudaOccupancyMaxPotentialBlockSize()diretamente para array 2D?

Solução potencial Tive uma ideia sobre a solução potencial:

E se eu calcular o número da rosca primeiro e, em seguida, usar o cudaOccupancyMaxPotentialBlockSize()cálculo do tamanho do bloco para a matriz 2D:

////total_thread_num = 640x480 = 307200
int total_thread_num = image.width * image.height;

////compute block/grid size
int min_grid_size, grid_size, block_size;
cudaOccupancyMaxPotentialBlockSize
(
    &min_grid_size, &block_size,
    my_cuda_kernel,
    0, total_thread_num
);

grid_size = (total_thread_num + block_size - 1) / block_size;

//launch CUDA kernel function
my_cuda_kernel<<<grid_size, block_size>>>(<arguments...>);

Em my_cuda_kernel, ele calcula o índice correspondente com base no tamanho da imagem:

__global__ void my_cuda_kernel()
{
    //compute 2D index based on 1D index;
    unsigned int idx = BlockIdx.x * blockDim.x + threadIdx.x;
    unsigned int row_idx = idx / image.width;
    unsigned int col_idx = idx % image_width;

    /*kernel function code*/

}

Questão 2 Se o método da Questão 1 não for viável, posso usar o método acima?

Answer 1

Pergunta 1 Posso calcular o tamanho do bloco / grade usando este método?

Não.

É importante lembrar que essas chamadas de API fornecem o número que maximiza a ocupação de threads por bloco e não as dimensões do bloco. Se você executar a API duas vezes em cada direção, provavelmente obterá um tamanho de bloco ilegal quando os dois valores forem combinados. Por exemplo, se a ocupação que maximiza a contagem de encadeamentos para um kernel era 256, você poderia terminar com um tamanho de bloco de 256 x 256, que é muito maior do que 1024 encadeamentos totais por bloco, portanto, a falha de inicialização.

Questão 2 Se o método da Questão 1 não for viável, posso usar o método acima?

Em princípio, isso deve funcionar, embora você esteja sofrendo uma pequena penalidade de desempenho porque a operação do módulo inteiro não é particularmente rápida na GPU. Alternativamente, você pode calcular um tamanho de bloco 2D que satisfaça suas necessidades a partir do máximo de threads por bloco retornado pela API.

Por exemplo, se você deseja apenas blocos com 32 threads na dimensão do bloco que mapeará para a ordem principal de seus dados (para coalescência de memória), divida a contagem de threads por 32 (observando que a API sempre retornará uma rodada múltiplo de 32 threads por bloco porque esse é o tamanho do warp). Portanto, por exemplo, se o retorno de threads por bloco da API fosse 384, o tamanho do seu bloco seria 32 x 12.

Se você realmente quer algum tipo de esquema de mosaico que usa blocos quadrados, então é muito fácil descobrir que apenas 64 (8 x 8), 256 (16 x 16), 576 (24 x 24) e 1024 (32 x 32 ) são os tamanhos de bloco viáveis, que são números quadrados e múltiplos redondos de 32. Nesse caso, você provavelmente deseja selecionar o tamanho de bloco maior que seja menor ou igual à contagem total de threads retornada pela API.

Em última análise, como você escolhe fazer isso dependerá dos requisitos de seu código de kernel. Mas certamente é possível conceber um esquema para o dimensionamento de blocos 2D que seja compatível com as APIs de dimensionamento de blocos que o CUDA expõe atualmente