diff --git a/CUDA_GL_INTEROP.md b/CUDA_GL_INTEROP.md deleted file mode 100644 index 047246b..0000000 --- a/CUDA_GL_INTEROP.md +++ /dev/null @@ -1,89 +0,0 @@ -# Guía Técnica: Interoperabilidad CUDA-GL para Renderizado Zero-Copy (GPU NVIDIA) - -Esta guía explica la arquitectura y detalles técnicos de la solución de **Zero-Copy real (GPU-to-GPU)** implementada en este proyecto C++/Qt6 (Debian 13 con tarjeta gráfica NVIDIA). - ---- - -## 1. El Desafío Arquitectónico en NVIDIA - -En un reproductor de video acelerado tradicional, el flujo de datos suele sufrir un cuello de botella en el bus PCIe: - -``` -[Disco Duro] ──> [CPU / RAM] ──> [GPU (Decodifica)] ──> [CPU / RAM (Descarga)] ──> [GPU (Textura OpenGL)] -``` - -Este viaje de ida y vuelta a la memoria RAM principal satura la CPU debido al copiado de millones de píxeles por segundo (especialmente en resoluciones 4K). - -Para evitar esto en tu máquina con GPU NVIDIA, implementamos **Interoperabilidad CUDA-GL** con **FFmpeg nativo (NVDEC)**, logrando que el flujo sea puramente en la GPU: - -``` -[Disco Duro] ──> [CPU (Demuxing de comandos)] - │ - ▼ - [GPU NVDEC (Decodifica en VRAM)] ──(Copia interna CUDA-GL)──> [Texturas OpenGL (VRAM)] ──> [Shader CRT] ──> [Monitor] -``` - ---- - -## 2. Los 4 Pasos Clave del Código C++ - -La interoperabilidad se logra conectando las APIs de CUDA y OpenGL a través de la librería `libcuda.so` de NVIDIA. - -### Paso A: Registro del Recurso Compartido -En la inicialización, creamos dos texturas OpenGL vacías (`m_texY` para luminancia y `m_texUV` para crominancia) y las registramos en CUDA usando la función `cudaGraphicsGLRegisterImage`: -```cpp -cudaGraphicsGLRegisterImage(&m_cudaResourceY, m_texY, GL_TEXTURE_2D, cudaGraphicsRegisterFlagsWriteDiscard); -cudaGraphicsGLRegisterImage(&m_cudaResourceUV, m_texUV, GL_TEXTURE_2D, cudaGraphicsRegisterFlagsWriteDiscard); -``` -*Esto le dice al driver de NVIDIA que reserve una "dirección de memoria compartida" que tanto CUDA como OpenGL pueden entender.* - -### Paso B: Mapeo y Bloqueo de Recursos -Durante el renderizado de cada fotograma (`paintGL()`), debemos asegurar que OpenGL no intente leer de las texturas mientras CUDA escribe en ellas. Bloqueamos los recursos mediante: -```cpp -cudaGraphicsMapResources(1, &m_cudaResourceY, 0); -cudaGraphicsMapResources(1, &m_cudaResourceUV, 0); -``` - -### Paso C: Copia VRAM-to-VRAM (Device-to-Device) -Extraemos los manejadores de arrays internos de CUDA vinculados a nuestras texturas OpenGL y ejecutamos una copia directa por hardware: -```cpp -cudaArray_t arrayY, arrayUV; -cudaGraphicsSubResourceGetMappedArray(&arrayY, m_cudaResourceY, 0, 0); -cudaGraphicsSubResourceGetMappedArray(&arrayUV, m_cudaResourceUV, 0, 0); - -// Copias internas de memoria física en la VRAM -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); -``` -*La bandera `cudaMemcpyDeviceToDevice` garantiza que la GPU use su controlador DMA interno en VRAM. La CPU solo envía las coordenadas de inicio y fin, pero no toca un solo píxel.* - -### Paso D: Desmapeo -Liberamos las texturas para que OpenGL vuelva a tomar el control y dibuje la escena: -```cpp -cudaGraphicsUnmapResources(1, &m_cudaResourceY, 0); -cudaGraphicsUnmapResources(1, &m_cudaResourceUV, 0); -``` - ---- - -## 3. Resolviendo tus Problemas Específicos - -Durante el desarrollo de esta integración en tu máquina, solucionamos dos bugs críticos de bajo nivel: - -### Caso A: La imagen cabeza abajo (Inversión vertical de origen) -* **Problema:** OpenGL posiciona el origen $(0,0)$ en la esquina inferior izquierda. CUDA lo posiciona arriba a la izquierda. El video se veía al revés. -* **Solución:** Modificamos el Vertex Buffer (VBO) en C++ para invertir el eje vertical ($V$) de las coordenadas de textura sin costo de procesamiento: - * Fila inferior: Coordenada de textura $V = 1.0$ (en lugar de $0.0$) - * Fila superior: Coordenada de textura $V = 0.0$ (en lugar de $1.0$) - -### Caso B: La mitad derecha de la pantalla verde (Ancho de plano UV en NV12) -* **Problema:** El formato decodificado es **NV12**. El plano UV (color) tiene la mitad del ancho de píxeles del video ($W/2$). Pero al estar $U$ y $V$ entrelazados (2 bytes por píxel), el ancho físico en bytes de cada fila del buffer de color es **$2 \times (W / 2) = W$ bytes**. -* **El bug:** Originalmente le pasábamos `w / 2` (el conteo de píxeles) a `cudaMemcpy2DToArray`. Al exigir esta función el ancho en **bytes**, la GPU solo copiaba la mitad izquierda de la información de color. La mitad derecha quedaba vacía (`0.0`), lo cual en la matriz de conversión YUV-a-RGB de tu fragment shader se traduce en un **tono verde intenso**. -* **Solución:** Cambiamos el parámetro de ancho a `w` (ancho completo en bytes), rellenando la textura de croma completa. - ---- - -## 4. Beneficios Obtenidos en tu Sistema -1. **Carga del Procesador (CPU) de 0%:** El procesador de tu sistema queda totalmente libre para otras tareas de tu aplicación principal. -2. **Latencia Cero:** Al evitar el cuello de botella del bus PCIe, no hay retraso (*stuttering*) al reproducir archivos de altos bitrates o resoluciones 4K/8K. -3. **Procesado de Shaders Directo:** Dado que la textura se actualiza directamente en OpenGL, los filtros CRT y la curvatura de lente se aplican instantáneamente al renderizar sobre la pantalla del reproductor. diff --git a/guide.html b/guide.html deleted file mode 100644 index a586305..0000000 --- a/guide.html +++ /dev/null @@ -1,250 +0,0 @@ - - - - -Comparativa: Decodificación Tradicional vs CUDA-GL Interop Zero-Copy - - - -
- -
-

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. -

- -
- Diagrama de flujo tradicional -
- -

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). -

- -
- Diagrama de flujo de interoperabilidad -
- -

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`. -

-
- -
- Guía de Integración Técnica • Debian 13 • GPU NVIDIA -
- -
- -