Da otto anni Apple usa un’unica API — Metal — per il rendering grafico e per il calcolo general-purpose su GPU, con lo stesso linguaggio shader in entrambi i casi. Alla WWDC di oggi, 6 giugno 2022, ha annunciato Metal 3 per macOS Ventura e iOS 16. Sotto il nome c’è una struttura tecnica precisa, e da quella dipende cosa si riesce davvero a fare su un Mac Apple Silicon usato come ambiente di calcolo.

Contesto

Fino al 2014 le piattaforme Apple esponevano OpenGL e OpenGL ES per la grafica, OpenCL per il calcolo su GPU. Sono API disegnate tra gli anni Novanta e i primi Duemila, con un modello a stato implicito: il driver tiene una macchina a stati globale, deduce quando ricompilare le pipeline e serializza la sottomissione dei comandi. Su hardware mobile il costo in CPU di ogni draw call pesava in modo sproporzionato rispetto al lavoro effettivo della GPU.

Apple ha presentato Metal alla WWDC del 2 giugno 2014, per iOS 8 sui system-on-chip A7 e A8. L’obiettivo dichiarato era tagliare l’overhead in CPU della sottomissione e rendere esplicita la gestione di stato e risorse. La stessa direzione — controllo esplicito, driver più sottile — è poi tornata in Direct3D 12 (2015) e in Vulkan (2016).

Architettura

Metal descrive il lavoro con un insieme ristretto di oggetti espliciti. MTLDevice è la GPU. MTLCommandQueue e MTLCommandBuffer reggono le code e i batch di comandi. Tre encoder coprono i tipi di workload: MTLRenderCommandEncoder per la grafica, MTLComputeCommandEncoder per il calcolo, MTLBlitCommandEncoder per le copie di memoria. Gli stati di pipeline (MTLRenderPipelineState, MTLComputePipelineState) si compilano in anticipo e a runtime restano immutabili.

La sottomissione è esplicita: l’applicazione costruisce un command buffer, lo registra (commit) sulla coda e si occupa da sé della sincronizzazione, con fence, event e shared event. Il driver non prova a indovinare quando uno stato è cambiato. Le risorse (MTLBuffer, MTLTexture, MTLHeap) dichiarano uno storage mode esplicito — shared, managed, private, memoryless — che fissa dove risiedono e come la CPU vi accede.

Il modello sposta complessità dal driver all’applicazione. In cambio elimina un’intera categoria di costi nascosti: validazioni implicite, ricompilazioni di pipeline al momento sbagliato, copie non richieste tra spazi di memoria.

Un solo linguaggio shader

Gli shader si scrivono in Metal Shading Language (MSL), derivato da C++14 con estensioni per la GPU. Lo stesso linguaggio copre gli stadi vertex, fragment, tile e compute; con Metal 3 si aggiungono gli object e mesh shader. Il compilatore metal, costruito su LLVM, produce un binario intermedio in formato AIR (Apple Intermediate Representation) impacchettato in un metallib. Solo a runtime il driver traduce l’AIR nel codice macchina della GPU concreta. In pratica la stessa sorgente di calcolo gira sulle GPU integrate iOS e su Apple Silicon desktop senza riscrivere niente.

Un kernel di calcolo minimo in MSL — un SAXPY, y = a·x + y — fa vedere la struttura:

kernel void saxpy(
    device const float* x [[buffer(0)]],
    device float*       y [[buffer(1)]],
    constant float&     a [[buffer(2)]],
    uint                i [[thread_position_in_grid]])
{
    y[i] = a * x[i] + y[i];
}

Gli attributi tra doppie parentesi quadre legano i parametri agli slot di binding e alla posizione del thread nella griglia. Non c’è un linguaggio a parte per il calcolo: lo stesso compilatore e la stessa pipeline servono grafica e compute.

Memoria unificata su Apple Silicon

Con il passaggio ad Apple Silicon — M1, novembre 2020 — la memoria è fisicamente unificata: CPU, GPU e Neural Engine attingono allo stesso pool di DRAM. In Metal una risorsa con storage mode shared vive in questa memoria comune, senza copie esplicite tra host e device. È una differenza strutturale rispetto al modello a GPU discreta, dove un buffer va trasferito sul bus PCI Express prima che la GPU possa leggerlo.

Per il calcolo questo riduce o azzera il costo del trasferimento dei dati, che su pipeline con molti tensori intermedi arriva a dominare il tempo totale. Il limite si sposta sulla banda di memoria e sulla capienza del pool unificato, condiviso con tutto il resto del sistema.

Lo stack di calcolo e machine learning

Sopra Metal, Metal Performance Shaders (MPS) offre kernel ottimizzati per ciascuna famiglia di GPU Apple: convoluzioni, prodotti tra matrici, FFT, filtri immagine, primitive per reti neurali. MPSGraph aggiunge un’API a grafo orientata al machine learning, con un ruolo vicino a quello di cuDNN affiancato a un compilatore di grafi.

Su questo stack poggiano i framework di livello più alto. Core ML usa MPS come backend per l’esecuzione su GPU. Alla WWDC di questa settimana Apple e Meta hanno mostrato in preview il device mps per PyTorch, che mappa le operazioni su MPSGraph e richiede macOS 12.3 o successivo. C’è un backend tensorflow-metal con lo stesso ruolo. Il Neural Engine, come unità a sé, divide il lavoro: i kernel generici girano su GPU via MPS, mentre i modelli compatibili passano al Neural Engine attraverso Core ML.

Limiti

Metal è legato alle piattaforme Apple. Non esiste un’implementazione cross-vendor: il codice MSL non gira su GPU non Apple, e portare un kernel scritto per CUDA vuol dire riscriverlo, non ricompilarlo. Progetti come MoltenVK traducono Vulkan su Metal in user space, ma nel verso opposto — far girare codice Vulkan su macOS — non in questo.

Sul calcolo pesante Apple non offre hardware da server. Non ci sono configurazioni data-center con GPU Apple, né interconnessioni tra nodi paragonabili a quelle usate per il training distribuito su GPU discrete. La memoria unificata è capiente per gli standard di una workstation, ma resta condivisa con CPU e sistema, e la banda — pur larga su M1 Max e Ultra — non arriva a quella della VRAM dedicata di fascia alta. Apple Silicon regge quindi come ambiente di prototipazione e di inferenza locale per modelli di taglia media; per il training su larga scala lo stack di riferimento resta quello a GPU discrete.

Il compilatore metal e la documentazione di MSL e MPS sono pubblici nella Apple Developer Documentation; il linguaggio è specificato, anche se l’implementazione del driver e l’AIR a basso livello non lo sono.

https://developer.apple.com/metal/ https://developer.apple.com/videos/play/wwdc2022/10066/ https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf https://developer.apple.com/documentation/metalperformanceshaders https://developer.apple.com/metal/pytorch/ https://pytorch.org/blog/introducing-accelerated-pytorch-training-on-mac/ https://www.noze.it/insights/apple-metal-api/

Immagine di copertina: Primo piano del chip Apple M1 con il dissipatore integrato rimosso: si vede il die di silicio quadrato al centro del package e i due… — foto di Sonic8400, CC0 — https://commons.wikimedia.org/wiki/File:Apple_M1_-_APL1102_-_IHS_removed.jpg