Comparativa: Decodificación Tradicional vs CUDA-GL Interop Zero-Copy

1. Caso A: Decodificación Acelerada Tradicional (Con Copia PCIe Redundante)

En la decodificación acelerada por GPU clásica (la que usa la mayoría de reproductores estándar), el video se procesa en el chip decodificador físico de la GPU (NVDEC). Sin embargo, al no estar enlazado con la API gráfica, la CPU tiene que realizar una descarga forzada de los píxeles a la RAM del sistema (`av_hwframe_transfer_data`). Posteriormente, el hilo de renderizado vuelve a subir la imagen a la GPU como una textura de OpenGL (`glTexSubImage2D`).

Este flujo cruza el bus PCIe dos veces por cada fotograma, lo que sobrecarga la CPU en un 15% - 30% a altas resoluciones y limita la fluidez.

2. Caso B: Decodificación con Interoperabilidad CUDA-GL (Zero-Copy)

En la solución optimizada, el video se decodifica en VRAM por NVDEC. Luego, las texturas de OpenGL se mapean directamente en CUDA. Realizamos una transferencia interna **dentro del mismo bus de memoria de la tarjeta gráfica** mediante la función `cudaMemcpy2DToArray` con la instrucción `cudaMemcpyDeviceToDevice`.

Los píxeles nunca abandonan la VRAM. El bus PCIe queda libre al 100%, la CPU no procesa copias de datos (0% de consumo) y el renderizado se realiza a la velocidad física de la memoria del procesador gráfico (Giga/Terabytes por segundo).

3. Detalles Técnicos de Implementación

Punto 1: Inicialización del Decodificador de Hardware NVIDIA (NVDEC)

En el archivo videodecoder.cpp, configuramos FFmpeg para crear un contexto de hardware forzado a utilizar la API de CUDA:

AVHWDeviceType devType = AV_HWDEVICE_TYPE_CUDA;
m_hwPixFmt = AV_PIX_FMT_CUDA;
av_hwdevice_ctx_create(&m_hwDeviceCtx, devType, nullptr, nullptr, 0);

El decodificador físico de NVIDIA (NVDEC) descomprime la imagen directamente en VRAM en formato **NV12**, retornando un fotograma (`AVFrame`) que contiene punteros de CUDA en VRAM.

Punto 2: Creación y Registro de Texturas OpenGL

En el archivo glwidget.cpp, creamos las texturas vacías en OpenGL. Como el formato del video en VRAM es NV12, se crean dos texturas: una para la Luminancia ($Y$) y otra para la Crominancia entrelazada ($UV$). Reservamos el tamaño de almacenamiento en VRAM usando `nullptr` como puntero de datos inicial:

// Reservar textura Y (Luminancia - 1 canal)
glBindTexture(GL_TEXTURE_2D, m_texY);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RED, w, h, 0, GL_RED, GL_UNSIGNED_BYTE, nullptr);

// Reservar textura UV (Crominancia - 2 canales)
glBindTexture(GL_TEXTURE_2D, m_texUV);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RG, w / 2, h / 2, 0, GL_RG, GL_UNSIGNED_BYTE, nullptr);

Luego, registramos los identificadores de estas texturas en CUDA a través de la API CUDA-GL Interop:

cudaGraphicsGLRegisterImage(&m_cudaResourceY, m_texY, GL_TEXTURE_2D, cudaGraphicsRegisterFlagsWriteDiscard);
cudaGraphicsGLRegisterImage(&m_cudaResourceUV, m_texUV, GL_TEXTURE_2D, cudaGraphicsRegisterFlagsWriteDiscard);

Punto 3: Transferencia Directa VRAM-to-VRAM (Device-to-Device)

En el método `paintGL()`, el hilo de pintado de OpenGL bloquea las texturas de lectura externa y delega temporalmente el control a CUDA. Luego, se realiza una copia directa de memoria entre las superficies del decodificador y las texturas:

// Bloqueo de textura para OpenGL
cudaGraphicsMapResources(1, &m_cudaResourceY, 0);
cudaGraphicsMapResources(1, &m_cudaResourceUV, 0);

// Copia ultra-rápida interna en GPU
cudaMemcpy2DToArray(arrayY, 0, 0, (void*)frame->data[0], frame->linesize[0], w, h, cudaMemcpyDeviceToDevice);
cudaMemcpy2DToArray(arrayUV, 0, 0, (void*)frame->data[1], frame->linesize[1], w, h / 2, cudaMemcpyDeviceToDevice);

// Devolución del control de las texturas a OpenGL
cudaGraphicsUnmapResources(1, &m_cudaResourceY, 0);

Punto 4: Conversión Cromática y Filtros en el Fragment Shader

Para corregir el origen invertido entre OpenGL y CUDA, las coordenadas del Quad están configuradas con el eje vertical invertido ($V = 1.0 - V$). En el fragment shader (shader.frag), leemos ambas texturas en paralelo y realizamos la conversión matemática de NV12 a RGB:

float y = texture(u_texY, uv).r;
vec2 uvVal = texture(u_texUV, uv).rg;
float u = uvVal.r - 0.5;
float v = uvVal.g - 0.5;

float r = y + 1.402 * v;
float g = y - 0.344136 * u - 0.714136 * v;
float b = y + 1.772 * u;
vec3 color = vec3(r, g, b);

Por último, sobre esta señal de color RGB, se aplican los filtros de curvatura 3D de la lente, líneas de barrido (scanlines), aberración cromática y viñeteado.

4. Resolución de Bugs Críticos en el Proyecto

Efecto Espejo

Inversión de Coordenadas Y (OpenGL vs CUDA)

OpenGL lee las texturas desde la esquina inferior izquierda, mientras que CUDA genera los frames con su origen en la esquina superior izquierda. La imagen aparecía invertida verticalmente. Esto se corrigió invirtiendo directamente la coordenada vertical de textura en el Quad de OpenGL en C++, evitando añadir costo aritmético al Fragment Shader.

Tono Verde

Copia de Crominancia UV en bytes

El formato de video NV12 almacena el color agrupando U y V de forma entrelazada (2 bytes por píxel). Aunque el número de píxeles horizontales es de $W/2$, el ancho en bytes es de **$2 \times (W / 2) = W$ bytes**. La API de CUDA `cudaMemcpy2DToArray` exige el ancho en bytes.

Al indicarle previamente `w / 2`, solo se copiaba el lado izquierdo de los datos de color de la imagen. La mitad derecha del buffer se mantenía en `0.0`. Al procesar la conversión matemática YUV-a-RGB en el fragment shader, una crominancia nula se desvía fuertemente hacia una tonalidad verde. Se corrigió modificando el ancho de transferencia de la copia a `w` bytes en `glwidget.cpp`.

Flujos de Arquitectura Gráfica