Aceleración Paralela de Convolución de Imágenes con C++/SYCL y oneAPI

Este proyecto se enfoca en la implementación de la operación de convolución de imágenes utilizando SYCL (un modelo de programación paralela de alto rendimiento basado en C++). El objetivo es aprovechar las capacidades de cómputo paralelo de GPUs y otros aceleradores de hardware para mejorar significativamente la eficiencia en el procesamiento de imágenes.

Análisis del Problema: Convolucción de Imágenes

La convolución de imágenes es una operación fundamental en el procesamiento de imágenes, utilizada para aplicar filtros y detectores de características. Conceptualmente, implica deslizar una pequeña matriz (el kernel o filtro) sobre la imagen, multiplicando los elementos del kernel por los píxeles de la imagen correspondientes y sumando los resultados para obtener el valor del píxel de salida. La fórmula general es:

S(i, j) = ∑<sub>k</sub> ∑<sub>l</sub> I(i + k, j + l) ⋅ K(k, l)

Donde S(i, j) es el valor del píxel en la posición (i, j) de la imagen resultante, I(i + k, j + l) es el valor del píxel en la imagen original, y K(k, l) es el peso del kernel en la posición (k, l).

Tradicionalmente, esta operación puede ser computacionalmente intensiva, especialmente con imágenes de alta resolución o grandes volúmenes de datos. El procesamiento secuencial resulta ineficiente. El cómputo paralelo, particularmente en GPUs, ofrece una vía para acelerar significativamente este proceso, dividiendo la imagen en porciones y procesándolas simultáneamente.

Introducción a oneAPI

oneAPI es un modelo de programación unificado, abierto y basado en estándares desarrollado por Intel. Su propósito es simplificar el desarrollo de aplicaciones de alto rendimiento en arquitecturas de cómputo heterogéneas, que incluyen CPUs, GPUs, FPGAs y otros aceleradores de hardware. La premisa principal de oneAPI es "escribe una vez, ejecuta en cualquier lugar", permitiendo a los desarrolladores crear código que puede ser portado y optimizado para diversas plataformas de hardware.

Beneficios de usar oneAPI:

  • Portabilidad entre Hardware: Escribe código una vez y ejecútalo en diferentes arquitecturas de hardware, incluyendo hardware no Intel.
  • Optimización de Rendimiento: Proporciona bibliotecas y herramientas optimizadas, especialmente para hardware Intel, para mejorar el rendimiento de las aplicaciones.
  • Simplificación del Desarrollo: Unifica la API y las cadenas de herramientas, reduciendo la complejidad y la curva de aprendizaje para la programación heterogénea.
  • Soporte para Lenguajes Modernos: Compatible con lenguajes como C++ y Python.
  • Fomento de la Innovación: Facilita la exploración e implementación de nuevos algoritmos y aplicaciones mediante modelos de programación de alto nivel.

Arquitectura de oneAPI:

La especificación abierta de oneAPI incluye:

  • Data Parallel C++ (DPC++): Un lenguaje de programación basado en C++ que es compatible con la especificación SYCL. DPC++ permite escribir código fuente único donde el código del host y los kernels del acelerador pueden coexistir.
  • SYCL: Un estándar abierto para programación paralela C++ de fuente única, desarrollado por el Khronos Group. SYCL permite programra CPUs, GPUs, FPGAs y otros aceleradores de manera eficiente. SYCL 2020 introduce características como Unified Shared Memory (USM), algoritmos de grupo de trabajo y subgrupo, y una interoperabilidad mejorada con diversas APIs de back-end.
  • oneAPI Level Zero: Una interfaz de hardware de bajo nivel que permite a los desarrolladores interactuar directamente con el hardware.

oneAPI se complementa con varios toolkits, como el Intel® oneAPI Base Toolkit (fundamental para la construcción de aplicaciones) y el Intel® AI Analytics Toolkit (para acelerar flujos de trabajo de IA y Machine Learning).

Programación Paralela con oneAPI

Modelo de Plataforma y Ejecución:

oneAPI adopta un modelo de plataforma basado en SYCL, donde un host (generalmente la CPU) coordina y controla la ejecución del cómputo en uno o más dispositivos (aceleradores como GPUs). La ejecución se gestiona a través de colas de comandos, a las cuales se envían grupos de comandos que contienen operaciones como llamadas a kernels y accesos a memoria.

El cómputo en el dispositivo se distribuye jerárquicamente a través de ND-ranges, grupos de trabajo (work-groups) y elementos de trabajo (work-items). Cada work-item ejecuta el código del kernel.

Modelo de Memoria:

El modelo de memoria de oneAPI, también basado en SYCL, abstrae la gestión de la memoria entre el host y los dispositivos. Se utilizan buffers (para datos generales) e imágenes (para datos con estructura espacial) para encapsular los datos. Los accesores (accessors) definen cómo y cuándo los kernels o el host acceden a estos buffers, especificando el modo de acceso (lectura, escritura) y el orden de ejecución para gestionar dependencias.

DPC++ simplifica la gestión de datos al encargarse de las transferencias implícitas cuando se usan buffers y accesores. Sin embargo, también permite un control explícito sobre el movimiento de datos para optimizaciones de rendimiento.

Flujo de Programación DPC++:

  1. Asignación de Memoria Host: Reservar memoria en el sistema principal.
  2. Creación de Buffers SYCL: Crear buffers SYCL a partir de la memoria host y definir accesores para especificar cómo se accederá a los datos desde el host y los dispositivos.
  3. Creación de Colas: Instanciar una cola SYCL para enviar comandos de ejecución al dispositivo seleccionado (CPU, GPU, etc.). Se pueden usar selectores para elegir el disposiitvo.
  4. Llamada a Kernels: Definir y ejecutar kernels paralelos (generalmente como funciones lambda C++) utilizando parallel_for dentro de un grupo de comandos enviado a la cola.
  5. Lectura de Resultados: Una vez completada la ejecución en el dispositivo, leer los datos modificados de los buffers SYCL de vuelta a la memoria del host.

Implementación de Convolución de Imágenes

Carga de Imágenes:

Se utilizan las bibliotecas stb\_image y stb\_image\_write para cargar y guardar imágenes. El código carga una imagen, la almacena en un buffer SYCL y se prepara para el procesamiento.


#include <CL/sycl.hpp>
#include <cmath>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>

#define STB_IMAGE_WRITE_IMPLEMENTATION
#include "stb_image_write.h"

#define STB_IMAGE_IMPLEMENTATION
#include "stb_image.h"

// ...

int width, height, channels;
// Carga la imagen desde la ruta especificada. '3' indica que queremos 3 canales (RGB).
uint8_t* imageData = stbi_load(imagePath, &width, &height, &channels, 3);
// Crea un buffer SYCL para almacenar los datos de la imagen.
cl::sycl::buffer<uint8_t, 1> imageBuffer{ imageData, static_cast<size_t>(width * height * channels) };

Inicialización de la Cola del Dispositivo:

Se crea una cola SYCL para interactuar con el dispositivo de cómputo.


// Crea una cola SYCL. default_selector intentará seleccionar el mejor dispositivo disponible (ej. GPU).
cl::sycl::queue deviceQueue{ cl::sycl::default_selector{} };

Conversión a Escala de Grises:

Se implementa una función luminance para calcular el brillo perceptivo de un píxel RGB y se aplica en paralelo para convertir la imagen a escala de grises.


// Función para calcular la luminancia perceptiva de un píxel RGB.
float calculateLuminance(uint8_t r, uint8_t g, uint8_t b) {
    // Normaliza los valores de color a [0, 1].
    float r_lin = static_cast<float>(r) / 255.0f;
    float g_lin = static_cast<float>(g) / 255.0f;
    float b_lin = static_cast<float>(b) / 255.0f;
    // Fórmula de luminancia perceptiva (CIE 1931).
    return 0.2126f * r_lin + 0.7152f * g_lin + 0.0722f * b_lin;
}

// Buffer para la imagen en escala de grises.
cl::sycl::buffer<float, 1> greyScaleBuffer{ static_cast<size_t>(width * height) };

// Envía un comando de ejecución a la cola.
deviceQueue.submit([&](cl::sycl::handler& h) {
    // Accesor para escribir en el buffer de escala de grises. discard_write es eficiente ya que no preserva datos antiguos.
    auto outAccessor = greyScaleBuffer.get_access<cl::sycl::access::mode::discard_write>(h);
    // Accesor para leer de forma segura el buffer de la imagen original.
    auto inAccessor = imageBuffer.get_access<cl::sycl::access::mode::read>(h);

    // Ejecuta un kernel paralelo sobre cada píxel de la imagen.
    h.parallel_for(cl::sycl::range<1>(width * height), [=](cl::sycl::id<1> idx) {
        // Calcula el desplazamiento en el buffer de la imagen original (RGB).
        int pixelOffset = 3 * idx[0];
        // Convierte el píxel RGB a escala de grises y lo almacena.
        outAccessor[idx[0]] = calculateLuminance(inAccessor[pixelOffset],
                                                 inAccessor[pixelOffset + 1],
                                                 inAccessor[pixelOffset + 2]);
    });
});

Implementación de la Convolución:

Se define un kernel de convolución y se aplica a la imagen en escala de grises. El kernel, un array float kernel[9], representa el filtro que se desliza sobre la imagen.


// Define el kernel de convolución (ejemplo: detector de bordes Sobel).
const float convolutionKernel[9] = {
    -1.0f, -2.0f, -1.0f,
     0.0f,  0.0f,  0.0f,
     1.0f,  2.0f,  1.0f
};

// Buffer para el resultado de la convolución.
cl::sycl::buffer<float, 1> convolutionResultBuffer{ static_cast<size_t>(width * height) };

// Envía el comando de convolución a la cola.
deviceQueue.submit([&](cl::sycl::handler& h) {
    // Accesor para leer la imagen en escala de grises.
    auto inAccessor = greyScaleBuffer.get_access<cl::sycl::access::mode::read>(h);
    // Accesor para escribir el resultado de la convolución.
    auto outAccessor = convolutionResultBuffer.get_access<cl::sycl::access::mode::discard_write>(h);

    // Kernel paralelo para realizar la convolución. Se usa un rango 2D para las dimensiones de la imagen.
    // Nota: Es crucial manejar los bordes de la imagen para evitar accesos fuera de límites.
    h.parallel_for(cl::sycl::range<2>(width, height), [=](cl::sycl::id<2> idx) {
        float convolutionSum = 0.0f;
        int kernelSize = 3; // El kernel es 3x3.

        // Itera sobre el kernel de convolución.
        for (int ky = 0; ky < kernelSize; ++ky) {
            for (int kx = 0; kx < kernelSize; ++kx) {
                // Calcula las coordenadas del píxel de la imagen correspondientes al kernel.
                int imageY = idx[1] + ky - kernelSize / 2;
                int imageX = idx[0] + kx - kernelSize / 2;

                // Verifica si las coordenadas están dentro de los límites de la imagen.
                if (imageY >= 0 && imageY < height && imageX >= 0 && imageX < width) {
                    // Obtiene el valor del píxel y del kernel.
                    float pixelValue = inAccessor[imageY * width + imageX];
                    float kernelValue = convolutionKernel[ky * kernelSize + kx];
                    // Realiza la multiplicación y suma.
                    convolutionSum += pixelValue * kernelValue;
                }
            }
        }
        // Almacena el resultado de la convolución para el píxel actual.
        outAccessor[idx[0] * width + idx[1]] = convolutionSum;
    });
});

Salida de Resultados:

Una vez completadas las operaciones, se espera la finalización de la cola y se guarda la imagen resultante.


// Espera a que todos los comandos en la cola se completen.
deviceQueue.wait();

// Obtiene acceso de lectura al buffer de resultados para guardarlo.
auto resultAccessor = convolutionResultBuffer.get_access<cl::sycl::access::mode::read>();

// Guarda la imagen resultante como un archivo PNG. '1' indica 1 canal (escala de grises).
stbi_write_png("convolution_output.png", width, height, 1, resultAccessor.get_pointer(), width);

// Libera la memoria cargada originalmente.
stbi_image_free(imageData);
// Libera explícitamente la memoria asignada por SYCL en el dispositivo (si es necesario).
// cl::sycl::free(resultAccessor.get_pointer(), deviceQueue); // Usar con precaución, puede que no sea necesario si el buffer se destruye automáticamente.

Verificación del Algoritmo

Para validar la implementación, se puede comparar los resultados con un conjunto de datos de referencia conocido, como el utilizado en los Hackathon Benchmark Datasets. Esto implica leer matrices de entrada y kernels desde archivos y comparar la salida generada por el código SYCL con los resultados esperados.


// Ejemplo conceptual de lectura de datos (simplificado)
// En una implementación real, esta función leería la imagen y el kernel
// y los prepararía para ser cargados en buffers SYCL.

// void loadData(const char* filePath, std::vector<float>& matrix, std::vector<float>& kernel) {
//     // Lógica para leer el archivo y poblar matrix y kernel
//     // ...
// }
//
// int main() {
//     // ... inicialización SYCL ...
//     std::vector<float> inputMatrix;
//     std::vector<float> convolutionKernel;
//     loadData("input_data.txt", inputMatrix, convolutionKernel);
//
//     // ... crear buffers SYCL a partir de inputMatrix y convolutionKernel ...
//     // ... ejecutar kernel de convolución ...
//     // ... obtener resultados y compararlos ...
// }

La verificación implicaría comparar los píxeles de la imagen de salida generada con los valores esperados para el mismo conjunto de entrada y kernel. Cualquier discrepancia señalaría la necesidad de depuración en la lógica del kernel o en el manejo de datos.

Etiquetas: SYCL OneAPI C++ convolución de imágenes procesamiento de imágenes

Publicado el 10-16 10:57

Etiquetas populares

©2026 Friki Work