Como determinar o desempenho da GPU CUDA?

Question

Estou escrevendo um programa cuda para combinar cada imagem de entrada com resolução ~ 180X180, com cerca de 10.000 imagens de modelo com resolução ~ 128 * 128.O objetivo é alcançar desempenho em tempo real, ou seja,Correspondência de modelos de 25 a 30 imagens de entrada (cada uma com todos os 10.000 modelos) em 1 segundo.

atualmente estou usando a seguinte abordagem

1. Pré-carregou todos os modelos na memória global da GPU para salvar operações de E/S em tempo de execução.
2. Criado um único kernel para combinar uma imagem de origem com todas as imagens de modelo e retornar uma matriz para todas as correspondências positivas.
3. Fazendo todas as operações no domínio do tempo (sem usar FFT).por isso, tentei a implementação do Radix-4 fft, mas requer muitas leituras e gravações globais intermediárias, acabando por levar mais tempo.

até agora, para 1 imagem de entrada para 10.000 modelos, leva cerca de 2 segundos.

Minhas perguntas são:

1. Existe uma maneira de determinar se esta tarefa pode ser realizada em tempo real ou não?Quero dizer, com a ajuda de FLOPS máximos e limitações de largura de banda de E/S, etc.
2. Como calcular se a GPU está sendo totalmente utilizada no máximo?
3. Possíveis maneiras de melhorar o desempenho?

Especificações da máquina:[i7-4770, 8GB, GTX-680]

Explicação do código do kernel atual:

1. todas as imagens de modelo [o tamanho é cerca de 128X128 em RGB] são carregadas na memória da GPU.A ideia é salvar E/S durante a operação em tempo de execução.
2. Cada imagem de entrada é carregada na memória de textura, por isso as texturas são uma boa opção para endereçamento 2D.
3. Cada "Bloco" possui 1.024 threads.
4. Cada thread calcula o valor de cada pixel de saída, o tamanho da saída é [31X31 = 961 pixels].
5. O número de blocos lançados é igual ao número de imagens de modelo correspondentes.

Código do kernel:

__global__ void cudaMatchTemplate(TemplateArray *templates, uchar *Match)
{
    int global = blockIdx.x*blockDim.x + threadIdx.x;

    __shared__ int idx[TEMPLATE_MATCH_DIM];
    __shared__ float out_shared[TEMPLATE_MATCH_DIM];

    //halving the template size....
    int rows = (templates[blockIdx.x].nHeight)/2;
    int cols = (templates[blockIdx.x].nWidth)/2;

    int fullCol = templates[blockIdx.x].nWidth;

    int x = templates[blockIdx.x].nMatchLeft;
    int y = templates[blockIdx.x].nMatchTop;

    int offset_y =  (threadIdx.x/TEMPLATE_MATCH_SIZE);
    int offset_x =  (threadIdx.x - offset_y*TEMPLATE_MATCH_SIZE);

    // *************** Performing match in time domain *****************************//
    int sum = 0;
    float temp;
    int idxXFactor = 3*(2*(offset_x) + x);
    int idxYFactor = 2*(offset_y) + y ;

    for (int i = 0; i < rows; i++)
    {
        int I=3*i*fullCol;
        int sourceIdxY = idxYFactor + 2*i;
        for (int j = 0; j < cols; j++)
        {
            int J=3*j;
            int sourceIdxX = idxXFactor + 2*J;          
            int templateIdx = 2*I+2*J;
            //**** R *****//
            temp = float(tex2D(SourceImgColorTex,sourceIdxX,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx]);
            sum = sum + temp*temp;
            //**** G *****//
            temp = float(tex2D(SourceImgColorTex,sourceIdxX+1,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +1]);
            sum = sum + temp*temp;
            //**** B *****//
            temp = float(tex2D(SourceImgColorTex,sourceIdxX+2,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +2]);
            sum = sum + temp*temp;
        }
    }

    __syncthreads();

//placing all values in shared memory for comparison.
    if(threadIdx.x < TEMPLATE_MATCH_DIM)
    {
        idx[threadIdx.x] = threadIdx.x;
        out_shared[threadIdx.x] = sum;
    }
    __syncthreads();


// //computing the Min location.....//

#pragma unroll
    for(int s=512; s>0; s>>=1) 
    {
        if ((threadIdx.x < s) &&((threadIdx.x + s)<TEMPLATE_MATCH_DIM))
        {
            idx[threadIdx.x] = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? idx[threadIdx.x] : idx[threadIdx.x + s];
            out_shared[threadIdx.x]  = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? out_shared[threadIdx.x] : out_shared[threadIdx.x + s];           
        }

    }

    __syncthreads();

    if(threadIdx.x <1)
    {
        int half_Margin = MARGIN_FOR_TEMPLATE_MATCH/2;
        int matchY = idx[0]/TEMPLATE_MATCH_SIZE ;
        int matchX = idx[0] - matchY * TEMPLATE_MATCH_SIZE;

        int diff = absolute(half_Margin - matchX) + absolute(half_Margin - matchY);
        if(diff < THRESHOLD)
        {
            Match[blockIdx.x] = 1;
        }
        else
            Match[blockIdx.x] = 0;

    }
}

Answer 1

Tentarei responder a maioria das suas perguntas

Existe uma maneira de determinar se esta tarefa pode ser realizada em tempo real ou não?Quero dizer, com a ajuda de FLOPS máximos e limitações de largura de banda de E/S, etc.

Não tenho ideia de como determinar se o kernel é alcançável em tempo real ou não, você pode maximizar seu kernel CUDA usando Calculadora de ocupação CUDA, Você pode considerar o uso de textura, memória de superfície, memória constante, memória de host fixada e muito mais, eles estão à sua implementação de algoritmo.

Como calcular se a GPU está sendo totalmente utilizada no máximo?

Você pode usar a Calculadora de Ocupação CUDA e o criador de perfil visual CUDA.Eu recomendo fortemente o uso do criador de perfil visual, pois ele o guiará pela compreensão do CUDA.

Possíveis maneiras de melhorar o desempenho?

Existem vários métodos interessantes que, 1º, você pode maximizar sua chamada de kernel usando o método acima, se isso não for suficiente, tente implementar o pipeline usando objetos de fluxo para copiar os trabalhos de dados e computação ao mesmo tempo.

se não der certo, tente trabalhar com latência, opere vários threads acessando a GPU ao mesmo tempo, já que o CC 3.5 CUDA lançou o HyperQ, isso pode te ajudar a completar várias chamadas em paralelo.

Se isso não funcionar, considere usar vários dispositivos GPU.

Por favor, mantenha-nos informados