Tutorial completo de Metal Compute en macOS

Inicio » Mac » Tutorial completo de Metal Compute en macOS

Si trabajas con un Mac Apple Silicon y todavía no estás exprimiendo Metal Compute para tus tareas de GPU, estás dejando mucha potencia encima de la mesa. Metal es la API de bajo nivel de Apple para gráficos y cómputo, y en los chips M1, M1 Pro, M1 Max y posteriores se convierte en una herramienta brutal para renderizar, procesar imágenes o acelerar simulaciones científicas y machine learning.

En las siguientes líneas vamos a ver, con calma pero en profundidad, cómo funciona Metal Compute en macOS, desde el modelo de ejecución (command queues, command buffers y encoders) hasta las mejores prácticas para sacar partido de la memoria unificada, los caches de la GPU y la optimización de kernels. Además, repasaremos cómo se conecta todo esto con el pipeline gráfico, cómo escribir shaders en el lenguaje de Metal basado en C++ y qué dicen las propias herramientas de Apple cuando se perfila un proyecto real.

Qué es Metal Compute y por qué es clave en macOS

Metal es la API moderna y de bajo overhead de Apple para enviar trabajo a la GPU, tanto de gráficos como de cómputo. Nació para iOS y macOS, y hoy es el camino oficial para aprovechar al máximo los chips Apple Silicon, especialmente en portátiles como los MacBook Pro con M1 Pro y M1 Max, donde la CPU y la GPU comparten memoria física.

A diferencia de otras APIs más antiguas, Metal está pensada para ser muy fina, eficiente y multihilo. Permite que varias hebras de la CPU construyan trabajo para la GPU a la vez, y te da flexibilidad para compilar los shaders (funciones de GPU) tanto en tiempo de compilación como en tiempo de ejecución.

Cuando hablamos de Metal Compute, nos referimos a la parte de la API dedicada a programación GPGPU (general purpose GPU): usar la GPU para algo más que dibujar vértices en pantalla. Esto incluye tareas como procesado de imágenes o vídeo, simulaciones físicas, ray tracing, machine learning o transformaciones masivas de datos.

En macOS con Apple Silicon, Metal Compute cobra más importancia gracias a la arquitectura de memoria unificada (UMA). CPU y GPU no tienen memorias separadas: comparten un mismo pool de RAM, lo que elimina copias innecesarias entre “system RAM” y “VRAM” y simplifica mucho el diseño de las aplicaciones.

Los MacBook Pro con M1 Pro y M1 Max pueden montar GPUs de hasta 16 y 32 núcleos respectivamente, con un ancho de banda de memoria mucho mayor que generaciones anteriores y con working sets de GPU de decenas de gigas. Esto abre la puerta a flujos de trabajo que antes estaban restringidos a sobremesas potentes.

Modelo de ejecución de Metal: de la CPU a la GPU

Para entender Metal Compute en macOS hay que interiorizar bien el modelo de ejecución de comandos. Todas las operaciones de GPU que lanzas desde tu app siguen un camino bastante claro: device → command queue → command buffer → encoder → commit.

En la base está el MTLDevice, el objeto raíz que representa una GPU concreta y que te sirve para crear el resto de recursos: buffers, texturas, colas de comandos, pipelines, etc. En la práctica, en un Mac con una única GPU integrada, pedir el “default device” suele ser suficiente para empezar a trabajar.

Sobre el dispositivo se crean una o varias MTLCommandQueue. Cada cola es el canal lógico a través del cual se envían cuerpos de trabajo a la GPU. La cola se encarga de mantener el orden y la sincronización básica entre command buffers.

Un MTLCommandBuffer es un contenedor transitorio que acumula comandos para la GPU: render passes, dispatches de compute, copias de recursos, etc. La CPU construye estos command buffers, los rellena usando encoders y al final los “commitea” para que la GPU los ejecute. La llamada a commit es asíncrona: devuelve al momento y la GPU empezará a trabajar cuando le toque en la cola.

Dentro del command buffer necesitas uno o varios command encoders específicos según el tipo de tarea: un render encoder para draws 3D/2D, un blit encoder para copias/movs de recursos y un compute encoder para despachar kernels de cálculo general. En un mismo command buffer puedes mezclar varios encoders y varios dispatches si lo organizas bien.

Un punto importante es que puedes construir múltiples command buffers en paralelo desde distintas hebras de CPU. Si el orden de ejecución entre ellos importa, puedes usar el método enqueue o simplemente controlar el orden de commit. Y si necesitas saber cuándo ha terminado un cuerpo de trabajo concreto, los command buffers ofrecen:

• Completion handlers asíncronos que se ejecutan cuando la GPU ha acabado ese buffer.
• Métodos síncronos como waitUntilCompleted, menos recomendables porque bloquean la hebra de CPU.

La clave es encontrar un equilibrio entre granularidad de los command buffers y latencia. Si haces command buffers muy pequeños, perderás tiempo en overhead de envío y la GPU puede quedarse “a medio gas”. Si son demasiado grandes, puedes introducir latencias poco deseables. Lo ideal suele ser agrupar suficiente trabajo para mantener la GPU ocupada sin bloquear la lógica de tu aplicación.

Memoria unificada y gestión de recursos en Metal

Con Apple Silicon, CPU y GPU acceden a la misma región física de memoria, lo que simplifica bastante la vida pero también exige pensar en términos de sincronización y working sets más que en copias explícitas.

En Metal, la memoria unificada se refleja en recursos shared (compartidos), que pueden ser leídos y escritos tanto por CPU como por GPU. En lugar de duplicar datos entre “RAM de sistema” y “VRAM”, trabajas sobre una sola instancia del recurso y te preocupas de coordinar quién lo toca y cuándo.

El objetivo es reducir al mínimo las copias y apoyarse en la UMA para ganar ancho de banda efectivo. Si CPU y GPU pelean por el mismo buffer a la vez, lo normal es recurrir a técnicas como el “multi-buffering”: por ejemplo, la CPU escribe en el buffer n mientras la GPU lee el buffer n−1, e intercambias roles en cada frame.

Además del total de memoria que la GPU puede ver (por ejemplo, unos 21 GB en un M1 Pro/Max con 32 GB de RAM, y hasta 48 GB en configuraciones de 64 GB), Apple introduce el concepto de working set: la cantidad máxima de memoria que un solo command encoder puede referenciar a la vez antes de que la plataforma tenga que hacer malabares de residencia.

Este valor se obtiene en tiempo de ejecución mediante recommendedMaxWorkingSetSize en el device. Lo recomendable es usar este límite como referencia para no dar por sentado que puedes tenerlo todo mapeado en un único encoder. Si tu escena o dataset es muy grande, parte el trabajo en varios encoders y command buffers de forma que el working set activo en cada paso se mantenga dentro de valores razonables.

Detrás de bastidores, Metal se encarga de la residencia virtual de recursos GPU: puedes reservar más memoria de GPU de la que cabe en el working set, y el runtime gestionará qué está “dentro” en cada momento, evitando los límites rígidos de la VRAM tradicional. Aun así, si ajustas tu uso de recursos a las recomendaciones, minimizarás page faults y sorpresas en rendimiento.

Cómo mantener la GPU ocupada: batching y concurrencia

En el nivel de command buffer hay un coste fijo de enviar trabajo a la GPU, así que es un error lanzar unidades diminutas de trabajo continuamente. Si te pasas con la fragmentación, la lógica de envío y sincronización te puede comer más tiempo que la ejecución real de los kernels.

Una buena práctica es agrupar varios encoders y dispatches en el mismo command buffer siempre que tenga sentido. Por ejemplo, puedes encadenar varias fases de un pipeline de procesado de imagen o de cómputo científico dentro de un mismo buffer, evitando burbujas en la línea de tiempo de la GPU.

Si tu aplicación depende de resultados de la GPU para decidir qué hacer después en el mismo frame, es fácil introducir huecos donde la GPU se queda esperando. En estos casos suele compensar usar varias hebras de CPU que vayan preparando trabajo futuro, incluso aunque todavía no tengas todas las respuestas del frame actual.

En términos de compute puro, la prioridad es que cada dispatch tenga suficientes hilos totales para ocupar todos los núcleos de la GPU, y que cada hilo haga un trabajo “no ridículo” comparado con el coste de lanzarlo. Un patrón clásico de Metal Compute es asignar un hilo por píxel en un procesado de imagen y lanzar un único dispatch que cubra toda la textura.

Cuando quieras lanzar kernels pequeños en número de hilos, puedes marcar los dispatches como concurrentes en lugar de serializados, permitiendo que la GPU solape varios de ellos y así aproveche mejor sus recursos internos.

En muchos casos se ha visto que apps que van estupendamente en un M1 básico no escalan automáticamente bien en M1 Pro/M1 Max, simplemente porque no alimentan suficientes hilos o command buffers como para ocupar toda la GPU ampliada. Revisar el patrón de submit de trabajo suele dar mejoras considerables sin tocar el código de los kernels.

Texturas vs buffers: sacar partido a los caches de la GPU

Las GPUs de Apple Silicon incluyen caches L1 separados para lecturas desde buffers y desde texturas. Si solo usas buffers como fuente de datos, estás dejando sin usar la parte del hardware optimizada para accesos texturales, que puede marcar la diferencia cuando mueves grandes volúmenes de datos.

Cuando una función de compute accede a un buffer muy grande calentando constantemente un área de memoria distinta, es fácil que “atraviese” el cache, es decir, que cada lectura tenga que irse a RAM. El ancho de banda y la latencia de la RAM, aunque muy buenos en Apple Silicon, son peores que los de los caches integrados en el chip.

En cambio, si conviertes parte de esos datos en texturas de Metal, puedes aprovechar un segundo cache L1 dedicado a texturas, ganando espacio efectivo de cache y reduciendo el tráfico hacia RAM. Metal, además, puede reorganizar internamente (swizzle/twiddle) los texels para que un patrón de acceso 2D o 3D típico se ajuste mejor al cache.

La gracia es que este reordenamiento es transparente para tu kernel: tú sigues leyendo con coordenadas normales, y es el hardware el que se encarga de acceder eficientemente a la memoria. Del mismo modo, Apple Silicon puede aplicar compresión de texturas sin pérdida de forma automática cuando las texturas son privadas de la GPU o cuando llamas a optimizeContentsForGPUAccess desde un blit encoder.

Para que esta compresión lossless esté disponible, la textura debe tener un usage compatible con shaderRead o renderTarget. Si, además, tus datos son imágenes donde un poco de compresión con pérdida es aceptable, puedes optar por formatos como ASTC o BC, que ofrecen ratios de compresión desde 4:1 hasta 36:1, reduciendo tanto el consumo de memoria como el ancho de banda de lectura.

En resumen en este punto, si notas que tu kernel está limitado por memoria y usa solo buffers, plantéate mover parte de los datos a texturas con buenos formatos y usos bien declarados. Es una de las optimizaciones más agradecidas en Apple Silicon.

Metal Compute en código: comandos, hilos y threadgroups

En la práctica, un paso de compute en Metal sigue un patrón bastante constante: crear command buffer, crear compute encoder, fijar pipeline, vincular buffers/texturas y despachar hilos. Todo ello, apoyándose en un pipeline de cómputo precompilado.

Primero obtienes un MTLCommandBuffer de tu cola y a partir de él creas un compute command encoder. Después necesitas un MTLComputePipelineState, que encapsula un kernel de compute (función de “shader” de tipo kernel) ya compilado y listo para ejecutarse.

Ese pipeline se construye a partir de una función de librería Metal, normalmente obtienes la función con makeFunction(name:) sobre la librería por defecto y luego llamas a makeComputePipelineState en el device. Este paso puede ser costoso, así que es aconsejable hacerlo una vez y reutilizar el pipeline.

Una vez fijado el pipeline en el compute encoder, vinculas los buffers y texturas que tu kernel necesita mediante llamadas setBuffer y setTexture en los índices adecuados. A partir de ahí toca decidir cómo organizas los hilos:

• threadsPerGrid: tamaño total del grid de trabajo (1D, 2D o 3D) que quieres procesar.
• threadsPerThreadgroup: tamaño de cada grupo de hilos (threadgroup) que se ejecuta junto y puede compartir memoria local.

Metal te da información útil como threadExecutionWidth y maxTotalThreadsPerThreadgroup del pipeline, que suelen usarse para construir un threadgroup “eficiente”. Un patrón típico en 2D es elegir un bloque de, por ejemplo, 32×16 hilos, comprobar que no excede el máximo total y calcular el número de grupos necesarios para cubrir el grid completo, ajustando con redondeo hacia arriba.

Si el tamaño del grid no es múltiplo exacto del tamaño del threadgroup, puedes usar non-uniform threadgroups, que Metal soporta directamente. Otra alternativa es usar dispatchThreads (en lugar de dispatchThreadgroups) y dejar que el runtime se ocupe de la distribución interna respetando el límite de threadsPerThreadgroup.

Programar la GPU para tareas generales (GPGPU)

Una vez controlas cómo lanzar kernels, la gran pregunta es qué tipo de trabajo merece ir a la GPU. En general, cualquier tarea con un alto grado de paralelismo de datos, donde aplicas la misma operación a muchos elementos independientes, es candidata ideal.

Ejemplos habituales son el procesado de imágenes y vídeo (filtros, correcciones de color, convoluciones), simulaciones físicas o científicas, ray tracing y transformaciones de modelos 3D (como invertir coordenadas de vértices en un modelo que viene en un sistema de referencia diferente).

Un buen ejemplo didáctico es tomar un modelo 3D cargado en CPU (por ejemplo, con Model I/O) y usar la GPU para ajustar sus vértices en paralelo: podrías voltear el eje Z para pasar de un sistema de coordenadas diestro a uno zurdo, o recalcular ciertos atributos por vértice. En CPU basta con recorrer un buffer y modificar cada vértice; en GPU haces exactamente eso, pero con un kernel que recibe un puntero al array y un índice de hilo que indica qué elemento debe tocar.

Además, cuando necesitas operaciones globales sobre el conjunto de datos (como contar el número total de vértices procesados), entran en juego las operaciones atómicas. Metal expone tipos atómicos y funciones específicas para incrementos, sumas, etc., que garantizan ausencia de condiciones de carrera incluso cuando miles de hilos tocan la misma variable.

Eso sí, las atómicas globales no son gratis: abusar de ellas puede convertirse en el nuevo cuello de botella en hardware moderno, justo porque el resto de limitaciones (ALU, ancho de banda) han mejorado tanto que las esperas en atómicas se notan más. Una estrategia típica es agrupar resultados parciales en memoria compartida de threadgroup y solo usar atómicas globales al final del grupo.

Escribir shaders y kernels con el Metal Shading Language

El lenguaje de shaders de Metal está basado en C++ (un subconjunto centrado en C++11), extendido con tipos vectoriales y matriciales, atributos de función y espacio de direcciones para ajustarse a la GPU. Si ya sabes C++ y has tocado GLSL o HLSL, el cambio es bastante natural.

En lugar de variables globales para los inputs y outputs, Metal prefiere que todos los datos entren como parámetros de función y salgan como valor de retorno. Un vertex shader es una función declarada con el atributo vertex, un fragment shader con fragment y un kernel de compute con kernel.

Las estructuras que describen vértices, outputs interpolados o uniformes las defines tú mismo. Para decirle al compilador qué campo es la posición, cuál va a un color attachment o cómo enlazar un argumento con un buffer o textura de la API, se usan atributos y qualifiers como [[position]], [[color(n)]], [[buffer(i)]] o [[thread_position_in_grid]].

Metal define varios espacios de direcciones (address spaces) para punteros y referencias, fundamentales para que el compilador optimice bien la memoria:

• device/global: memoria general accesible por cualquier hilo (backed por RAM).
• constant: datos inmutables para un dispatch, que el hardware puede prefetch y cachear agresivamente.
• threadgroup: memoria compartida entre hilos de un mismo grupo, útil para reducciones y patrones cooperativos.
• thread: memoria privada de cada hilo, normalmente mapeada a registros.

Una recomendación importante es identificar bien qué datos son constantes y compartidos entre muchos hilos (por ejemplo, matrices de vista/proyección o descriptores de luces) y pasarlos por referencia en el espacio constant. Eso permite al compilador reducir el número de registros usados y mejorar la ocupación.

La librería estándar de Metal sustituye a la de C++ en código GPU e incluye funciones matemáticas vectoriales, trigonométricas, de interpolación, etc. Por defecto, las operaciones se compilan en modo “fast math”, que prioriza el rendimiento aunque pierdas algo de precisión en casos extremos (NaNs, rangos súper amplios de trigonometría, etc.). Para trozos concretos donde necesites precisión total, puedes llamar explícitamente a las variantes “precise” o ajustar las opciones de compilación.

Optimización de kernels: índices, registros y ocupación

Una vez que tu código funciona, llega el momento de exprimirlo. Las herramientas de Xcode, en particular el GPU Frame Debugger y el profiler, te dan acceso a contadores de hardware y estadísticas de compilación que permiten ver dónde se atasca cada kernel.

Entre los indicadores más útiles están la utilización de ALU (unidad aritmético-lógica) y su rol como “limiter” (cuánto tiempo se está limitado por ALU en lugar de por memoria, texturas, etc.), así como la ocupación, que mide qué porcentaje del máximo teórico de hilos activos se está usando.

Si la utilización de ALU es moderada pero el limiter para ALU es alto, puede que estés haciendo operaciones complejas (por ejemplo, log, trigonometría pesada o formatos de textura caros) que introducen stalls internos o pipelines infrautilizados. En ese caso, revisar la matemática, precalcular parte en CPU o usar versiones más baratas puede dar buen fruto.

La ocupación baja es otro foco de atención. Un valor reducido puede ser normal si el problema es pequeño o si el kernel está claramente limitado por ALU o ancho de banda. Sin embargo, una ocupación baja con limiters también bajos suele indicar que el compilador ha tenido que reservar muchos recursos por hilo: muchos registros, mucha memoria de threadgroup, etc.

Cuando la presión de registros es alta, el compilador empieza a hacer “spill”: algunos valores que cabrían en registros se derraman a memoria, lo que empeora la latencia y reduce aún más la capacidad de tener muchos hilos simultáneos. Herramientas de Xcode muestran el número de bytes “spilled” por kernel, que junto con la ocupación dan pistas claras de qué hay que atacar.

Algunas tácticas concretas de optimización a nivel de kernel son:

• Preferir tipos de 16 bits (half, short) cuando tenga sentido, liberando registros y permitiendo vectorizaciones más compactas.
• Evitar arrays grandes en la pila o estructuras enormes que acaben ocupando muchos registros por hilo.
• Diseñar los datos de entrada para explotar el espacio constant, reduciendo el uso de registros generales.
• Evitar indexar arrays en stack o en constant con índices dinámicos no conocidos en compilación, que fuerzan spills a memoria.

También es importante elegir bien el tamaño máximo de hilos por threadgroup: puedes fijarlo con maxThreadsPerThreadgroup en el pipeline descriptor o mediante atributos en el propio kernel (max_total_threads_per_threadgroup). Ajustar este valor al mínimo múltiplo útil del ancho de ejecución de hilos de la GPU suele facilitar que el compilador gestione mejor los registros y mejore la ocupación.

Con todo esto, Metal Compute en macOS se convierte en una herramienta tremendamente flexible: permite desde el clásico “triángulo giratorio” con uniforms animados y sincronizados por semáforos, hasta pipelines sofisticados de render híbrido y cómputo donde se concatenan kernels, se explota la memoria unificada y se mide cada milisegundo con los contadores de Xcode. Dominar el modelo de ejecución, la UMA, el uso combinado de buffers y texturas y las técnicas de profiling y optimización es lo que marca la diferencia entre una app que simplemente funciona y otra que realmente exprime la GPU de tu Mac al máximo.