¿Cómo determinar el rendimiento de CUDA GPU?

Question

Estoy escribiendo un programa CUDA para hacer coincidir cada imagen de entrada de la resolución ~ 180x180, con aproximadamente 10,000 imágenes de plantillas de resolución ~ 128 * 128. El objetivo es lograr el desempeño en tiempo real, es decir, una coincidencia de plantillas de 25 ~ 30 imágenes de entrada (cada una con las 10.000 plantillas) en 1 segundo.

Actualmente estoy usando el siguiente enfoque

1. precargó todas las plantillas en la memoria global de GPU para guardar las operaciones de E / S de tiempo de ejecución.
2. creó un solo kernel para que coincida con una imagen de origen con todas las imágenes de la plantilla y devuelva una matriz para todas las coincidencias positivas.
3. Haciendo todas las operaciones en el dominio del tiempo (no usando FFT). Razón, probé la implementación de Radix-4 FFT, pero requiere una gran cantidad de lecturas globales intermedias y escrituras que terminan tomando más tiempo.

Hasta ahora para 1 imagen de entrada a 10,000 plantillas, se toma alrededor de 2 segundos.

Mis preguntas son:

1. ¿Hay alguna manera de determinar si esta tarea es alcanzable en tiempo real o no? Me refiero con la ayuda de las fragmisas máximas y las limitaciones de ancho de banda de E / S E.T.C.
2. ¿Cómo calcular si la GPU está siendo completamente utilizada en su máximo?
3. Formas posibles de mejorar el rendimiento?

Especificaciones de la máquina: [I7-4770, 8GB, GTX-680]

Explicación del código de kernel actual:

1. Todas las imágenes de la plantilla [el tamaño es de aproximadamente 128x128 en RGB] están por carga en la memoria GPU. Idea es guardar E / S durante la operación de tiempo de ejecución.
2. Cada imagen de entrada está cargada en la memoria de la textura, la motivo de la textura es una buena opción para la dirección 2D.
3. cada "bloque" tiene 1024 hilos.
4. Cada hilo calcula el valor para cada píxel de salida, el tamaño de la salida es [31x31= 961 píxeles].
5. El número de bloques lanzados es igual a la cantidad de imágenes de plantillas que se están haciendo coincidir.

Código 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

Intentaré responder la mayoría de sus preguntas

¿Hay alguna manera de determinar si esta tarea es alcanzable en tiempo real o no? Quiero decir con la ayuda de chapas máximas y limitaciones de ancho de banda I / O E.T.C.

No tengo idea de cómo determinar si el kernel es realizado en tiempo real o no, puede maximizar su kernel de CUDA usando calculadora de ocupación de CUDA , Puede considerar usar la textura, la memoria de la superficie, la memoria constante, la memoria del host pinada y más, aquellos que están a la altura de la implementación de su algoritmo.

¿Cómo calcular si la GPU está siendo completamente utilizada en su máximo?

Puede usar la calculadora de ocupación de CUDA y CUDA VISUAL Profiler. Recomiendo encarecidamente el uso de los perfiles visuales, lo guiaría a través de la comprensión de CUDA.

Formas posibles de mejorar el rendimiento?

Hay varios métodos interesantes que lo hacen, 1º Puede maximizar la llamada de su kernel usando el método anterior, Si eso no es suficiente, intente implementar la tubería usando objetos de STREAM en orden. Copie los trabajos de datos y computación al mismo tiempo.

Si eso no va a funcionar, intente trabajar con la latencia, opere múltiples hilos que acceden a la GPU al mismo tiempo, ya que CC 3.5 CUDA lanzó Hyperq, esto puede ayudarlo a completar varias llamadas en paralelo.

Si eso no va a funcionar, considere usar varios dispositivos GPU.

Por favor, mantennos publicados