CUDA & OpenCV en C++

Pasar imagen a escala de grises

Este artículo forma parte de una serie de documentos en castellano orientados a mejorar algoritmos de análisis de movimiento en vídeo mediante computación paralela.

Imagen 1 — Lienzo de Leonid Afremov pasado a escala de grises mediante las técnicas de este documento

Primero, al grano: Si estás en Ubuntu Línux y has seguido las guías para instalar OpenCV y Cuda de este blog ya puedes descargar y ejecutar el código:

git clone https://github.com/jesusdario/cudacv-bw.git
cd cudacv-bw
# Crear el Makefile que enlaza a tus librerias CUDA y CV
cmake .
# Compila el codigo fuente
make
# Convierte la imagen a escala de grises
./to_bw image.jpg

Si no tienes git, puedes bajarte el código en un zip igualmente con el siguiente enlace. El código fichero a fichero está más abajo.

Si has tenido problemas al ejecutar cmake aquí dejo un Makefile simplificado que me ha servido para trastear y superar algunos problemas:

NVCC=nvcc
OPENCV_LIBPATH=/usr/lib
OPENCV_INCLUDEPATH=/usr/include
OPENCV_LIBS=-lopencv_core -lopencv_imgproc -lopencv_highgui
CUDA_INCLUDEPATH=/usr/local/cuda-6.5/include
NVCC_OPTS=-O3 -arch=sm_11 -Xcompiler -Wall -Xcompiler -Wextra -m64
GCC_OPTS=-O3 -Wall -Wextra -m64
executable: main.o rgba_to_grey.o Makefile
 $(NVCC) -o to_bw main.o rgba_to_grey.o -L $(OPENCV_LIBPATH) \ $(OPENCV_LIBS) $(NVCC_OPTS)
main.o: main.cpp preprocess.cpp utils.h
 g++ -c main.cpp $(GCC_OPTS) -I $(CUDA_INCLUDEPATH) -I \ $(OPENCV_INCLUDEPATH)
rgba_to_grey.o: rgba_to_grey.cu utils.h
 nvcc -c rgba_to_grey.cu $(NVCC_OPTS)
clean:
 rm -f *.o *.png hw

Para otra GPU probablemente haya que cambiar el parámetro de NVCC_OPTS -archi=sm_11 a otro superior. La mía es una GeForce 210.

Fichero a fichero

El código se divide en varios ficheros, aquí traduzco los comentarios al español. El código es exactamente el mismo que hay en github.

rgba_to_gray.cu

Contiene el kernel que va a lanzar en la GPU

/* rgba_to_gray.cu */
#include "utils.h"
#include <stdio.h>
#include <math.h>       /* ceil */
// Max Threads per block in GeForce 210
#define TxB 512
__global__
void rgba_to_grey_kernel(
const uchar4* const rgbaImage, 
unsigned char* const greyImage, 
int numRows, int numCols) {
// El mapeo de los componentes uchar4 aRGBA es:
// .x -> R ; .y -> G ; .z -> B ; .w -> A
//La salida debe ser resultado de aplicar la siguiente formula //resultado = .299f * R + .587f * G + .114f * B;
//Nota: Ignoramos el canal alfa
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  uchar4 px = rgbaImage[i]; // pixel que procesa este hilo
  greyImage[i] = .299f * px.x +
                 .587f * px.y +
                 .114f * px.z;
}
void rgba_to_grey(uchar4 * const d_rgbaImage, unsigned char* const d_greyImage, size_t numRows, size_t numCols)
{
  // Dado que no importa la posicion relativa de los pixels
  // en este algoritmo, la estrategia para asignar hilos a
  // bloques y rejillas sera sencillamente la de cubrir
  // a todos los pixeles con hebras en el eje X
  long long int total_px = numRows * numCols;  // total pixels
  long int grids_n = ceil(total_px / TxB); // grids numer
  const dim3 blockSize(TxB, 1, 1);
  const dim3 gridSize(grids_n, 1, 1);
  rgba_to_grey_kernel<<<gridSize, blockSize>>>(
                      d_rgbaImage, d_greyImage, numRows, numCols);
  cudaDeviceSynchronize(); 
  checkCudaErrors(cudaGetLastError());
}

preprocess.cpp

Lee la imagen y reserva espacio en la tarjeta gráfica para copiarla:

/* preprocess.cpp */
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/opencv.hpp>
#include "utils.h"
#include <cuda.h>
#include <cuda_runtime.h>
#include <string>
cv::Mat imageRGBA;
cv::Mat imageGrey;
uchar4        *d_rgbaImage__;
unsigned char *d_greyImage__;
size_t numRows() { return imageRGBA.rows; }
size_t numCols() { return imageRGBA.cols; }
// Devuelve un puntero de la version RGBA de la imagen de entrada
// y un puntero a la imagend e canal unico de la salida
// para ambos huesped y GPU
void preProcess(uchar4 **inputImage, unsigned char **greyImage,
                uchar4 **d_rgbaImage, unsigned char **d_greyImage,
                const std::string &filename) {
  //Comprobar que el contexto se inicializa bien
  checkCudaErrors(cudaFree(0));
  cv::Mat image;
  image = cv::imread(filename.c_str(), CV_LOAD_IMAGE_COLOR);
  if (image.empty()) {
    std::cerr << "Couldn't open file: " << filename << std::endl;
    exit(1);
  }
  cv::cvtColor(image, imageRGBA, CV_BGR2RGBA);
  
 // Reserva memoria para el output
  imageGrey.create(image.rows, image.cols, CV_8UC1);
*inputImage = (uchar4 *)imageRGBA.ptr<unsigned char>(0);
*greyImage  = imageGrey.ptr<unsigned char>(0);
const size_t numPixels = numRows() * numCols();
  //Reserva memoria en el dispositivo
  checkCudaErrors(
cudaMalloc(d_rgbaImage, sizeof(uchar4) * numPixels));
  checkCudaErrors(
cudaMalloc(d_greyImage, sizeof(unsigned char) * numPixels));
  checkCudaErrors(
cudaMemset(*d_greyImage, 0, numPixels * sizeof(unsigned char))); 
// Asegurate de que no queda memoria sin liberar
// Copia el input en la GPU
  checkCudaErrors(
cudaMemcpy(*d_rgbaImage, *inputImage,
 sizeof(uchar4) * numPixels, cudaMemcpyHostToDevice));
  
  d_rgbaImage__ = *d_rgbaImage;
  d_greyImage__ = *d_greyImage;
}

utils.h

Define una función para comprobar fallos en CUDA —esto siempre puede ser un poco tedioso—

/* utils.h */
#ifndef UTILS_H__
#define UTILS_H__
#include <iostream>
#include <iomanip>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>
#include <cassert>
#include <cmath>
#define checkCudaErrors(val) check( (val), #val, __FILE__, __LINE__)
template<typename T>
void check(T err, const char* const func, const char* const file, const int line) {
  if (err != cudaSuccess) {
    std::cerr << "CUDA error at: " << file << ":" << line << std::endl;
    std::cerr << cudaGetErrorString(err) << " " << func << std::endl;
    exit(1);
  }
}
#endif

main.cpp

Organiza el código y va llamando a las distintas partes.

/* main.cpp */
#include <iostream>
#include "utils.h"
#include <string>
#include <stdio.h>
// Declaramos la funcion que invoca al kernel
void rgba_to_grey(uchar4 * const d_rgbaImage,
                  unsigned char* const d_greyImage, 
                  size_t numRows, size_t numCols);
// Incluye las definiciones del fichero de mas arriba
#include "preprocess.cpp"
int main(int argc, char **argv) {
  uchar4        *h_rgbaImage, *d_rgbaImage;
  unsigned char *h_greyImage, *d_greyImage;
  std::string input_file;
  std::string output_file;
switch (argc)
  {
 case 2:
   input_file = std::string(argv[1]);
   output_file = "output.png";
   break;
 default:
      std::cerr << "Usage: ./to_bw input_file [output_filename]" << std::endl;
      exit(1);
  }
  // Carga la imagen y nos prepara los punteros para la entrada y  
  // salida de datos
  preProcess(&h_rgbaImage, &h_greyImage, 
        &d_rgbaImage, &d_greyImage, input_file);
  // Invoca al kernel
  rgba_to_grey(d_rgbaImage, d_greyImage, numRows(), numCols());
  size_t numPixels = numRows()*numCols();
  checkCudaErrors( 
    cudaMemcpy(h_greyImage, d_greyImage, 
      sizeof(unsigned char) * numPixels, cudaMemcpyDeviceToHost));
/* Saca la imagen en escala de grises */
  cv::Mat output(numRows(), numCols(), CV_8UC1, (void*)h_greyImage);
  // Abre la ventana
  cv::namedWindow("to_bw");
  // Pasa la imagen a la ventana anterior
  cv::imshow("to_bw", output);
  cvWaitKey (0);
  cvDestroyWindow ("to_bw");
  // Imprime a fichero
  cv::imwrite(output_file.c_str(), output);
/* Libera memoria*/
  cudaFree(d_rgbaImage__);
  cudaFree(d_greyImage__);
return 0;
}

CMakeLists.txt

Este fichero busca automáticamente las librerías que hemos instalado de CUDA y OpenCV (y otras cualquiera) y genera un Makefile sin que tengamos que escribirlo nosotros.

# Collect source files
cmake_minimum_required(VERSION 2.8)
find_package(CUDA)
find_package(OpenCV REQUIRED)
file( GLOB  hdr *.hpp *.h )
file( GLOB  cu  *.cu)
SET (to_bw_files main.cpp)
CUDA_ADD_EXECUTABLE(to_bw ${to_bw_files} ${hdr} ${cu})
target_link_libraries(to_bw ${OpenCV_LIBS})

Para compilar y ejecutar hay que repetir los pasos del principio del artículo:

# Buscar librerias
cmake .
# Compilar
make
# Ejecutar
./to_bw image.jpg

Siendo image.jpg un cuadro de Leonid Afremov:

Imagen 2 — image.jpg Lienzo de Leonid Afremov

Casi todos los conocimientos y materiales de éste artículo los aprendí en el curso de Udacity de “Introducción a la computación en paralelo”. El material está en inglés, pero han hecho un trabajo brutal de síntesis y de compilación de ejercicios que creo que se debe comentar. De hecho este artículo es acorde con la primera lección del curso.