Aller au contenu

9 · Portable : SYCL, OpenCL, Vulkan, WebGPU, Metal

Les options qui ne dépendent pas d'un seul fournisseur, leurs domaines de pertinence réels, et le prix mesuré de la portabilité.


9.1 Le prix de la portabilité

Commençons par le chiffre, parce qu'il structure tout le reste.

Le code portable est 10 à 30 % plus lent au niveau du noyau sur du silicium équivalent.

— comparaison SYCL / OpenCL / Vulkan Compute, 2026

Ce chiffre mérite d'être qualifié.

Pourquoi cet écart existe. Un noyau portable ne peut pas utiliser TMA, wgmma, tcgen05, les clusters, ni les particularités du wavefront. Il vise le dénominateur commun.

Pourquoi il n'est pas rédhibitoire. Si la portabilité vous permet d'acheter le meilleur rapport performance/prix du marché à chaque cycle, une flexibilité d'approvisionnement de 30 % couvre largement 30 % de performance par noyau.

Le calcul dépend de votre situation, pas d'un principe. Un laboratoire avec un parc mixte n'a pas le même arbitrage qu'un service dont toute l'infrastructure est sur H100.


9.2 SYCL

Le standard Khronos de programmation hétérogène en C++ moderne, source unique.

#include <sycl/sycl.hpp>

int main() {
    sycl::queue q;
    const int n = 1024;

    float* a = sycl::malloc_shared<float>(n, q);   // USM
    float* b = sycl::malloc_shared<float>(n, q);
    float* c = sycl::malloc_shared<float>(n, q);

    q.parallel_for(sycl::range<1>(n), [=](sycl::id<1> i) {
        c[i] = a[i] + b[i];
    }).wait();

    sycl::free(a, q);
    return 0;
}

Les points forts :

  • source unique : hôte et périphérique dans le même fichier, en C++ standard (pas d'extension de langage comme __global__) ;
  • USM (Unified Shared Memory) : des pointeurs utilisables des deux côtés, sans gestion explicite de tampons — une amélioration ergonomique majeure sur OpenCL ;
  • lambdas et templates : le C++ moderne fonctionne normalement ;
  • des back-ends pour Intel, NVIDIA, AMD et CPU.

Les implémentations :

Implémentation Portée
Intel oneAPI DPC++ Intel natif, + NVIDIA et AMD via des plugins Codeplay
AdaptiveCpp (ex-hipSYCL) NVIDIA, AMD, CPU, Intel

Quand SYCL est le bon choix

Quand une même base de code doit tourner sur AMD, Intel et NVIDIA, et que vous écrivez du C++ moderne. C'est le cas typique du calcul scientifique en centre de calcul académique, où le parc change à chaque appel d'offres.

C'est aussi le choix de plusieurs supercalculateurs européens, ce qui garantit un investissement soutenu.

Alternative à considérer dans le même créneau : Kokkos (Sandia National Laboratories), une bibliothèque C++ de portabilité de performance qui cible CUDA, HIP, SYCL et OpenMP depuis une seule source. Très utilisée en HPC américain, avec un modèle d'abstraction (vues, espaces mémoire, politiques d'exécution) souvent jugé plus pratique que SYCL pour du code scientifique existant.


9.3 OpenCL

Le standard historique (2009). Il reste pertinent pour une raison : la portée matérielle.

OpenCL tourne sur des GPU, des CPU, des FPGA, des DSP, des puces embarquées de tous les fabricants. Aucun autre standard n'a cette couverture.

// Le noyau, dans un fichier séparé ou une chaîne de caractères
__kernel void add(__global const float* a,
                  __global const float* b,
                  __global float* c) {
    int i = get_global_id(0);
    c[i] = a[i] + b[i];
}

Les inconvénients sont sérieux :

  • API C verbeuse : une centaine de lignes pour un « hello world » ;
  • compilation à l'exécution du code du noyau, depuis une chaîne ;
  • pas de source unique : hôte en C/C++, noyau en C OpenCL ;
  • support inégal : NVIDIA maintient OpenCL au minimum syndical (version 3.0, sans enthousiasme).

OpenCL en 2026

Pour un nouveau projet GPU, SYCL est presque toujours préférable — il est d'ailleurs souvent implémenté au-dessus d'OpenCL.

OpenCL reste justifié pour l'embarqué, les FPGA, et les cibles exotiques où rien d'autre n'existe.


9.4 Vulkan Compute

Vulkan est une API graphique et de calcul. Ses compute shaders sont un moyen légitime de faire du GPGPU.

#version 450
layout(local_size_x = 256) in;
layout(binding = 0) buffer A { float a[]; };
layout(binding = 1) buffer B { float b[]; };
layout(binding = 2) buffer C { float c[]; };

void main() {
    uint i = gl_GlobalInvocationID.x;
    c[i] = a[i] + b[i];
}

Compilé en SPIR-V, une représentation intermédiaire binaire portable.

Le cas d'usage propre de Vulkan Compute : quand le calcul est co-localisé avec du graphique. Rendu avec budget de temps par image, vision par ordinateur alimentant une passe de rendu, moteurs de réalité étendue mêlant calcul et graphique dans la même file de soumission.

Dans ces contextes, faire du calcul en CUDA exigerait une interopérabilité coûteuse ; en Vulkan, c'est la même API, les mêmes ressources, la même synchronisation.

Vulkan est aussi la meilleure option sur mobile (Android), où CUDA n'existe pas.

Extensions utiles pour le calcul : VK_KHR_cooperative_matrix (accès aux tensor cores), VK_KHR_shader_float16_int8, VK_KHR_buffer_device_address.

La courbe d'apprentissage

Vulkan est explicite jusqu'à l'excès : descriptor sets, pipeline layouts, command buffers, memory barriers, gestion manuelle de la mémoire. Un « hello world » de calcul fait plusieurs centaines de lignes.

Ne choisissez Vulkan pour du calcul pur que si vous y êtes déjà, ou si vous ciblez le mobile.


9.5 WebGPU et WGSL

L'API GPU du navigateur, successeur de WebGL. En 2026, elle est mûre : 82 % de support navigateur mondial.

État par navigateur :

Navigateur État en 2026
Chrome / Edge activé par défaut depuis la version 113 (avril 2023)
Safari activé
Firefox désactivé par défaut, pour des raisons de sécurité (empreinte numérique) et de stabilité des pilotes

Le langage de shaders est WGSL, d'inspiration Rust :

@group(0) @binding(0) var<storage, read>       a: array<f32>;
@group(0) @binding(1) var<storage, read>       b: array<f32>;
@group(0) @binding(2) var<storage, read_write> c: array<f32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
    let i = gid.x;
    if (i < arrayLength(&a)) {
        c[i] = a[i] + b[i];
    }
}

Le modèle est reconnaissable : workgroup_size est la taille de bloc, global_invocation_id l'indice global, var<workgroup> la mémoire partagée.

Ce que WebGPU permet en pratique :

  • inférence de modèles dans le navigateur (transformers.js, ONNX Runtime Web, WebLLM) ;
  • visualisation de données à grande échelle ;
  • rendu 3D temps réel ;
  • calcul scientifique interactif sans installation.

Ses limites :

  • pas d'accès direct aux tensor cores (des extensions sont en discussion) ;
  • pas de contrôle fin de la mémoire ;
  • des limites de sécurité sur les tailles d'allocation ;
  • performance typiquement 2 à 5× en dessous du natif.

Le vrai argument de WebGPU

Ce n'est pas la performance, c'est la distribution. Un modèle qui tourne dans le navigateur ne nécessite ni installation, ni serveur, ni transfert de données vers un tiers. Pour beaucoup d'applications — démonstrations, outils internes, traitement de données sensibles — cela vaut plus qu'un facteur 3 de performance.


9.6 Metal

Sur Apple, c'est la seule API GPU native.

#include <metal_stdlib>
using namespace metal;

kernel void add(device const float* a [[buffer(0)]],
                device const float* b [[buffer(1)]],
                device float*       c [[buffer(2)]],
                uint i [[thread_position_in_grid]]) {
    c[i] = a[i] + b[i];
}

Les particularités d'Apple Silicon :

Trait Conséquence
Mémoire unifiée pas de transfert CPU↔GPU, la contrainte PCIe disparaît
SIMD-group de 32 analogue au warp
Threadgroup memory analogue à la mémoire partagée
Bande passante élevée 400 à 800 Go/s sur les puces Max/Ultra
Capacité importante jusqu'à 512 Go de mémoire unifiée sur certaines configurations

La mémoire unifiée est décisive pour l'inférence LLM locale : un modèle de 70 milliards de paramètres quantifié tient dans la mémoire d'un Mac Studio, ce qui exigerait deux cartes de centre de données ailleurs. La bande passante limite le débit, mais la faisabilité change tout.

Les couches supérieures :

  • MPS (Metal Performance Shaders) : bibliothèque de primitives optimisées ;
  • MLX : le cadre d'apprentissage automatique d'Apple, conçu pour la mémoire unifiée, avec une API proche de NumPy/PyTorch ;
  • PyTorch MPS backend : support partiel mais fonctionnel.

Point pratique : la plupart des Mojo GPU Puzzles fonctionnent sur Apple Silicon via Metal, ce qui en fait une plateforme d'apprentissage viable sans carte NVIDIA.


9.7 Le tableau de décision

Situation Choix
Calcul scientifique, parc multi-fournisseur SYCL ou Kokkos
Portage d'un code Fortran/C++ existant OpenMP target ou OpenACC
Deep learning multi-fournisseur Triton (pas les APIs de ce chapitre)
Calcul intégré à du rendu graphique Vulkan Compute
Mobile Android Vulkan
Navigateur WebGPU / WGSL
Apple Metal ou MLX
FPGA, DSP, embarqué exotique OpenCL
NVIDIA seul, performance maximale CUDA (ce chapitre ne vous concerne pas)

Résumé de la partie 5

Les cinq idées à emporter

  1. Il n'y a pas un langage, il y a un gradient. Descendez d'un cran seulement quand le profileur le justifie.
  2. La première question est toujours : existe-t-il une bibliothèque ?
  3. Triton est le point d'entrée réaliste : ~90 % de la performance experte pour ~10 % de l'effort. Ses limites (layouts, synchronisation inter-blocs) définissent quand descendre vers Gluon, CuTe DSL ou CUDA C++.
  4. Le modèle par tuiles a gagné. Triton, Helion, cuTile, CuTe DSL, ThunderKittens, Mojo : cinq projets indépendants, la même idée.
  5. La portabilité coûte 10 à 30 % au niveau du noyau. Que ce soit trop cher dépend de votre parc, pas d'un principe.

Vérifiez que vous avez compris

Pourquoi Triton est-il un meilleur choix que SYCL pour du deep learning multi-fournisseur ?

Parce que les deux ne résolvent pas le même problème.

SYCL porte du C++ général : bon pour des solveurs, des simulations, des algorithmes irréguliers. Mais il n'a aucune abstraction de tuile ni d'accès portable aux unités matricielles.

Triton est spécialisé pour les tenseurs : tl.dot émet l'instruction matricielle appropriée sur chaque cible (wgmma sur Hopper, MFMA sur CDNA). C'est exactement ce dont le deep learning a besoin, et c'est ce que SYCL ne fait pas.

À l'inverse, Triton serait un mauvais choix pour un solveur de mécanique des fluides à maillage non structuré.

WebGPU est 2 à 5× plus lent que le natif. Dans quel cas est-ce acceptable ?

Chaque fois que l'alternative n'est pas « la même chose en natif » mais « rien du tout ».

Exemples concrets : une démonstration interactive qu'un utilisateur ouvre en un clic ; un outil traitant des données sensibles qui ne doivent pas quitter la machine ; une visualisation intégrée à une documentation ; un modèle d'inférence embarqué dans une application web sans coût de serveur.

Dans tous ces cas, le facteur 3 est le prix d'une distribution sans friction — et c'est un bon prix.

Pourquoi la mémoire unifiée d'Apple Silicon change-t-elle le raisonnement sur l'inférence LLM locale ?

Parce qu'elle supprime la contrainte capacité, qui est la contrainte binaire.

Sur une carte discrète, un modèle qui ne tient pas en VRAM ne tourne pas (ou tourne à la vitesse du PCIe, c'est-à-dire pas). Sur Apple Silicon, la mémoire du GPU est celle du système : un Mac Studio à 512 Go peut charger un modèle qu'aucune carte unique ne peut contenir.

La bande passante (400-800 Go/s) reste 4 à 10× inférieure à celle d'un H100, donc le débit est bien plus faible. Mais pouvoir le faire lentement est qualitativement différent de ne pas pouvoir le faire.

C'est aussi pourquoi la partie IA · Entraînement vs inférence insiste sur le fait que l'inférence à petit lot est un problème de mémoire, pas de calcul.


Partie suivante : Domaines d'application


Sources de ce chapitre