150 lines
9.9 KiB
Markdown
150 lines
9.9 KiB
Markdown
# Flujos de Arquitectura Gráfica
|
|
|
|
### Decodificación Tradicional vs Interoperabilidad CUDA-GL
|
|
|
|
---
|
|
|
|
## 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](file:///home/sebastian/repos/zero_copy_video_player/videodecoder.cpp), configuramos FFmpeg para crear un contexto de hardware forzado a utilizar la API de **CUDA**:
|
|
|
|
```cpp
|
|
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](file:///home/sebastian/repos/zero_copy_video_player/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:
|
|
|
|
```cpp
|
|
// 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:
|
|
|
|
```cpp
|
|
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:
|
|
|
|
```cpp
|
|
// 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 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](file:///home/sebastian/repos/zero_copy_video_player/shader.frag), leemos ambas texturas en paralelo y realizamos la conversión matemática de NV12 a RGB:
|
|
|
|
```glsl
|
|
float y = texture(u_texY, v_texCoord).r;
|
|
vec2 uvVal = texture(u_texUV, v_texCoord).rg;
|
|
|
|
// El formato NV12 almacena U y V entrelazados. Restamos 0.5 para centrar la crominancia.
|
|
float u = uvVal.r - 0.5;
|
|
float v = uvVal.g - 0.5;
|
|
|
|
// Matriz estándar de conversión YUV (NV12) a RGB
|
|
float r = y + 1.402 * v;
|
|
float g = y - 0.344136 * u - 0.714136 * v;
|
|
float b = y + 1.772 * u;
|
|
|
|
fragColor = vec4(r, g, b, 1.0);
|
|
```
|
|
|
|
---
|
|
|
|
## 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 mantí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`.
|
|
|
|
---
|
|
|
|
## 5. Conceptos Arquitectónicos de Alto Nivel
|
|
|
|
Para comprender mejor cómo esta solución maximiza el rendimiento del hardware gráfico, aquí se detallan los principios de diseño y concurrencia implementados:
|
|
|
|
### A. Representación de Formatos de Video en OpenGL (¿Por qué dos texturas?)
|
|
Las pantallas físicas (monitores) funcionan proyectando subpíxeles de color **RGB** (Rojo, Verde y Azul). Sin embargo, por cuestiones de compresión y eficiencia, los videos se almacenan en formatos **YUV (como NV12)**, separando la información de brillo (Luminancia / $Y$) y la de color (Crominancia / $UV$ submuestreada a la mitad de resolución).
|
|
|
|
Dado que OpenGL estándar no cuenta con un formato nativo para "NV12" en una única textura:
|
|
1. **Separación física**: Dividimos la información del fotograma en dos texturas de OpenGL: una para luminancia $Y$ (con formato `GL_RED` a resolución completa) y otra para crominancia $UV$ (con formato `GL_RG` a la mitad de la resolución en píxeles, pero con 2 bytes por píxel que contienen $U$ y $V$ entrelazados).
|
|
2. **Conversión en GPU**: El Fragment Shader ([shader.frag](file:///home/sebastian/repos/zero_copy_video_player/shader.frag)) muestrea ambas texturas simultáneamente en la misma coordenada normalizada y ejecuta la ecuación matemática de conversión YUV-a-RGB en el chip gráfico a la velocidad del hardware.
|
|
|
|
---
|
|
|
|
### B. Sincronización Concurrente mediante Mutex
|
|
La reproducción interactiva y fluida de video se basa en un diseño multihilo. Tenemos dos hilos que acceden concurrentemente a la variable compartida del fotograma activo (`m_currentFrame`):
|
|
- **Hilo de Decodificación (FFmpeg/CUDA)**: Escribe un nuevo fotograma en memoria tan pronto como se descomprime del archivo.
|
|
- **Hilo de Pintado (OpenGL/Qt)**: Lee la referencia del fotograma para actualizar las texturas en cada ciclo de dibujo (`paintGL`).
|
|
|
|
Si estos accesos ocurrieran al mismo tiempo sin control, se produciría una condición de carrera (*Race Condition*), provocando que OpenGL intente renderizar un fotograma a medio escribir. Esto resultaría en corrupciones visuales (*tearing*) o caídas del programa.
|
|
|
|
Para resolverlo, implementamos un cerrojo de exclusión mutua (`m_frameMutex`):
|
|
- El decodificador bloquea el mutex justo cuando va a guardar el frame (`av_frame_ref`).
|
|
- El componente OpenGL bloquea el mutex brevemente mediante `getCurrentFrameRef(...)` para clonar la referencia de manera atómica.
|
|
- El bloqueo es instantáneo (solo se copian punteros y metadatos, no los píxeles en sí), por lo que no degrada el rendimiento de la aplicación.
|
|
|
|
---
|
|
|
|
### C. Control de Framerate y Reloj de Referencia (Master Clock)
|
|
Para evitar que el video se reproduzca desordenado o excesivamente rápido (a la velocidad máxima del decodificador de hardware), el hilo de decodificación sincroniza su bucle con un reloj en tiempo real:
|
|
|
|
1. **La Base de Tiempo del Video**: Cada fotograma decodificado contiene un indicador de tiempo llamado **PTS (Presentation Time Stamp)**. Este número viene expresado en "ticks" de reloj internos del archivo multimedia. Multiplicamos el PTS por la base de tiempo fraccionaria del video para convertirlo a segundos, y luego a **milisegundos reales**:
|
|
$$\text{PTS en ms} = \text{PTS en ticks} \times \text{Base de tiempo} \times 1000$$
|
|
2. **Sincronización contra el reloj del sistema**: Mantenemos un reloj real (`QElapsedTimer`) que mide el tiempo transcurrido desde que pulsamos "Play".
|
|
3. **Control de Demora**: Si el fotograma decodificado tiene un PTS del futuro con respecto al reloj (`ptsMs > playbackTime`), calculamos la diferencia y ponemos a dormir al hilo decodificador (`QThread::msleep(delay)`) para ralentizarlo. Si el fotograma viene atrasado o al tiempo justo, el hilo no duerme y lo entrega a la GPU inmediatamente para evitar saltos o pérdida de sincronización.
|
|
|
|
---
|
|
|
|
*Guía de Integración Técnica • Debian 13 • GPU NVIDIA*
|