Le calcul GPU

En un mot#

BASIC peut lancer des compute shaders WebGPU : on dépose des nombres dans un tampon sur le GPU, on fait tourner un petit programme WGSL sur tous à la fois, puis on relit le résultat dans un memblock. Quatorze commandes, dans la catégorie Compute, font tout le travail.

⚠️ Engine uniquement. Le calcul GPU marche quand vous lancez Play dans l'éditeur, sur un GPU et un navigateur qui gèrent le calcul WebGPU. L'export HTML et les compilations Bevy ne l'ont pas. Écrivez le repli CPU avant d'en dépendre.

Quand il est disponible#

Appelez d'abord ComputeSupported(). Elle répond sur toutes les cibles, et toutes les autres commandes Compute supposent qu'elle a répondu 1.

Il n'y a pas d'émulation silencieuse : une commande Compute tourne sur le GPU ou s'arrête sur une erreur qui nomme la cible. Le schéma est toujours le même : tester ComputeSupported(), et prendre un chemin CPU quand elle vaut 0 (voir le schéma de repli plus bas).

Les pièces#

Les étapes 6 et 7 rendent des tâches de plateforme (jobs), les mêmes que celles des autres commandes asynchrones. On les interroge avec PlatformJobDone, PlatformJobError et PlatformJobResultNumber.

Les tampons de stockage#

buffer = CreateComputeBuffer(<nombreElements>, <type$>)
CreateComputeBuffer(<idTampon>, <nombreElements>, <type$>)
ComputeBufferFromMemblock(<tampon>, <memblock>)

• Trois types d'élément : "f32", "i32", "u32". Les trois font 4 octets. Un tampon de n éléments fait n * 4 octets.
• Le memblock doit avoir exactement la taille du tampon. Un tampon de 16 octets prend un memblock de 16 octets, ni plus ni moins. Toute autre taille arrête le programme.
• Les octets sont copiés tels quels. Écrivez les valeurs f32 avec MemblockFloat, les valeurs i32 et u32 avec MemblockInt, aux décalages 0, 4, 8….
• Un tampon contient au plus 128 Mio. C'est le minimum courant des GPU mobiles.
• La seconde forme prend l'identifiant que vous choisissez ; il ne doit pas être déjà pris.

Les pipelines WGSL#

shader = CreateComputeShader(<wgsl$>, <liaisonsJson$> [, <pointEntree$>])
CreateComputeShader(<idShader>, <wgsl$>, <liaisonsJson$> [, <pointEntree$>])
ComputeShaderBuffer(<shader>, <nom$>, <tampon>)

• Le noyau est du WGSL brut. Ce n'est pas du GLSL, et il n'est pas traduit. Il doit contenir un point d'entrée @compute, nommé main sauf si vous en passez un autre.

• Le JSON des liaisons associe chaque variable WGSL à son emplacement, exactement comme le WGSL la déclare :

  {"data": {"group": 0, "binding": 0}}
  

  Une entrée par variable de stockage, au moins une, aucun emplacement utilisé deux fois.

• ComputeShaderBuffer lie un tampon à une variable par son nom WGSL.

• Les tampons de stockage sont les seules liaisons. Pas d'uniforms, pas de textures. Pour passer un paramètre, comme un pas de temps, mettez-le dans un petit tampon de stockage à lui.

• CreateComputeShader ne vérifie que la forme : la source contient @compute, le JSON est valide, les emplacements sont distincts. Le WGSL lui-même compile en arrière-plan, et une erreur WGSL revient sur la tâche du dispatch, par PlatformJobError.

Dans une chaîne BASIC, les guillemets du JSON s'écrivent \" :

bindings$ = "{\"data\":{\"group\":0,\"binding\":0}}"
Console(bindings$)          ' {"data":{"group":0,"binding":0}}

Le dispatch, et sa place dans l'image#

job = ComputeDispatchAsync(<shader>, <groupesX> [, <groupesY> [, <groupesZ>]])

Les arguments sont des nombres de groupes de travail, pas des nombres d'éléments. Avec @workgroup_size(64) et n éléments, demandez (n + 63) DIV 64 groupes, et laissez le noyau ignorer les indices au-delà de la fin.

Un dispatch ne s'exécute pas quand vous l'appelez. Il est mis en file, et la file part vers le GPU juste avant que l'image suivante soit dessinée.

• La file garde son ordre. Deux dispatchs lancés l'un après l'autre s'exécutent dans cet ordre, même quand le pipeline du second est prêt le premier. Un second noyau peut lire ce que le premier a écrit.
• La tâche est terminée quand le travail est parti. Attendez-la avec PlatformJobDone, puis vérifiez PlatformJobError (vide en cas de succès).
• Un pipeline toujours pas prêt au bout de 10 secondes fait échouer sa tâche.
• Attendre exige un programme qui peut attendre. La boucle habituelle est WHILE PlatformJobDone(job) = 0 … Sync() … ENDWHILE. Son Sync() doit vraiment céder : le programme a besoin d'un DO … LOOP ou d'un END, ou la boucle doit être dans une étiquette. Sinon l'onglet se fige, voir La boucle de jeu.

La relecture dans un memblock#

job = ComputeBufferReadAsync(<tampon>)
job = ComputeBufferReadAsync(<tampon>, <idNouveauMemblock>)
memblock = PlatformJobResultNumber(job)

• La lecture crée un nouveau memblock avec les octets du tampon. Le nombre de la tâche est son identifiant.
• La seconde forme choisit cet identifiant, et il doit être libre. Relire dans un memblock qui existe arrête le programme : supprimez-le d'abord.
• ⚠️ La lecture n'est pas mise en file derrière les dispatchs. Ne la lancez qu'une fois la tâche du dispatch terminée, sinon vous lisez les données d'avant le passage du noyau.
• Un memblock contient au plus 100 000 000 octets : un tampon que vous relisez doit donc rester sous 25 000 000 éléments, en dessous des 128 Mio qu'un tampon peut contenir.
• Supprimez les memblocks dont vous n'avez plus besoin, et libérez les tâches terminées avec PlatformJobRelease.

Un exemple complet#

Deux noyaux sur un même tampon : le premier double chaque valeur, le second ajoute un. Les deux sont mis en file d'un coup, et l'ordre de la file garantit le résultat.

' 1, 2, 3, 4 deviennent 3, 5, 7, 9 sur le GPU.
IF ComputeSupported() = 0
  Console("Pas de calcul WebGPU sur cette cible")
  END
ENDIF

' Les données : quatre valeurs f32, de 4 octets chacune.
input = CreateMemblock(16)
FOR i = 0 TO 3
  MemblockFloat(input, i * 4, i + 1)
NEXT

' Un tampon de stockage de même taille en octets, rempli depuis le memblock.
buffer = CreateComputeBuffer(4, "f32")
ComputeBufferFromMemblock(buffer, input)

' Deux noyaux qui partagent leurs premières lignes. Un appel par élément.
head$ = "@group(0) @binding(0) var<storage, read_write> data: array<f32>; "
head$ = head$ + "@compute @workgroup_size(64) "
head$ = head$ + "fn main(@builtin(global_invocation_id) id: vec3<u32>) { "
head$ = head$ + "if (id.x >= arrayLength(&data)) { return; } "
bindings$ = "{\"data\":{\"group\":0,\"binding\":0}}"
doubler = CreateComputeShader(head$ + "data[id.x] = data[id.x] * 2.0; }", bindings$)
adder = CreateComputeShader(head$ + "data[id.x] = data[id.x] + 1.0; }", bindings$)
ComputeShaderBuffer(doubler, "data", buffer)
ComputeShaderBuffer(adder, "data", buffer)

' Les deux en file. Ils passent dans cet ordre, avant la prochaine image.
groups = (4 + 63) DIV 64
job1 = ComputeDispatchAsync(doubler, groups)
job2 = ComputeDispatchAsync(adder, groups)
WHILE PlatformJobDone(job2) = 0
  Sync()
ENDWHILE
IF PlatformJobError(job1) <> "" OR PlatformJobError(job2) <> ""
  Console("échec du dispatch : " + PlatformJobError(job1) + PlatformJobError(job2))
  END
ENDIF

' Relecture. Le nombre de la tâche est l'identifiant d'un nouveau memblock.
readJob = ComputeBufferReadAsync(buffer)
WHILE PlatformJobDone(readJob) = 0
  Sync()
ENDWHILE
result = PlatformJobResultNumber(readJob)
FOR i = 0 TO 3
  Console(Str(MemblockFloat(result, i * 4)))   ' 3, puis 5, 7, 9
NEXT

' Tout libérer.
DeleteMemblock(result)
DeleteMemblock(input)
DeleteComputeShader(doubler)
DeleteComputeShader(adder)
DeleteComputeBuffer(buffer)
END

Ce qui le fait marcher :

• Les lignes END. Elles permettent à Sync() de céder dans les boucles WHILE. Sans elles, le programme n'a ni boucle DO ni END, et les boucles figeraient l'onglet.
• Attendre job2 suffit. La file passe dans l'ordre : quand le second dispatch est parti, le premier l'est aussi.
• La lecture commence après la fin des dispatchs.
• arrayLength(&data) garde la fin du tableau. Un groupe de 64 lance 64 appels ; seuls les 4 premiers ont un élément.

Le schéma de repli#

Gardez les données dans un memblock sur les deux chemins. Le chemin GPU le téléverse et le relit ; le chemin CPU travaille dessus en place. Tout ce qui suit l'embranchement est le même code.

' Doubler 1024 valeurs : sur le GPU si possible, sinon sur le CPU.
count = 1024
data = CreateMemblock(count * 4)
FOR i = 0 TO count - 1
  MemblockFloat(data, i * 4, i)
NEXT

IF ComputeSupported()
  GOSUB DoubleOnGpu
ELSE
  DoubleOnCpu(data, count)
ENDIF
Console(Str(MemblockFloat(data, 10 * 4)))   ' 20, sur les deux chemins
END

FUNCTION DoubleOnCpu(mem, n)
  FOR k = 0 TO n - 1
    MemblockFloat(mem, k * 4, MemblockFloat(mem, k * 4) * 2)
  NEXT
ENDFUNCTION

DoubleOnGpu:
  buf = CreateComputeBuffer(count, "f32")
  ComputeBufferFromMemblock(buf, data)
  wgsl$ = "@group(0) @binding(0) var<storage, read_write> data: array<f32>; "
  wgsl$ = wgsl$ + "@compute @workgroup_size(64) "
  wgsl$ = wgsl$ + "fn main(@builtin(global_invocation_id) id: vec3<u32>) { "
  wgsl$ = wgsl$ + "if (id.x >= arrayLength(&data)) { return; } "
  wgsl$ = wgsl$ + "data[id.x] = data[id.x] * 2.0; }"
  kernel = CreateComputeShader(wgsl$, "{\"data\":{\"group\":0,\"binding\":0}}")
  ComputeShaderBuffer(kernel, "data", buf)
  job = ComputeDispatchAsync(kernel, (count + 63) DIV 64)
  WHILE PlatformJobDone(job) = 0
    Sync()
  ENDWHILE
  ' Relire sous le même identifiant : le libérer, puis laisser la lecture le recréer.
  DeleteMemblock(data)
  readJob = ComputeBufferReadAsync(buf, data)
  WHILE PlatformJobDone(readJob) = 0
    Sync()
  ENDWHILE
  DeleteComputeShader(kernel)
  DeleteComputeBuffer(buf)
RETURN

• Le chemin GPU est une étiquette, pas une FUNCTION : une FUNCTION ne peut pas attendre. Un GOSUB vers une étiquette qui attend exige une boucle DO ou un END dans le programme.
• Le chemin CPU est une FUNCTION ordinaire : il n'attend pas.
• Dans un vrai programme, vérifiez aussi PlatformJobError après chaque attente, comme dans l'exemple complet.

Pour tenir le code GPU tout à fait hors des exports, placez-le entre #IF ENGINE et #ENDIF : l'éditeur définit ENGINE, les compilations HTML et Bevy non. Voir le préprocesseur.

Durée de vie#

• Stop détruit tout. Tampons et shaders sont libérés quand le programme s'arrête, et leurs identifiants repartent de 1 au lancement suivant.
• Les dispatchs encore en file au Stop échouent avec Compute session ended before queued work was dispatched.
• Supprimez ce dont vous n'avez plus besoin pendant l'exécution : la mémoire GPU n'est pas récupérée pour vous tant que le programme tourne.

Ce qui change par rapport aux dialectes classiques#

AGK, DarkBASIC et PureBasic n'ont pas de calcul GPU. C'est un ajout d'aukimi, bâti sur WebGPU. Deux conséquences pour qui vient de ces dialectes :

• Les noyaux sont en WGSL, le langage de shaders de WebGPU. Les compute shaders GLSL ne se chargent pas.
• Tout est asynchrone. Un dispatch et une relecture sont des tâches qu'on attend d'une image à l'autre, jamais des appels qui bloquent jusqu'à ce que le GPU ait fini.

Les commandes Compute#

Limites et pannes courantes#

• Testez ComputeSupported() d'abord, et écrivez le chemin CPU.
• Le memblock téléversé a exactement la taille du tampon.
• Les nombres du dispatch sont des groupes. Gardez la fin du tableau dans le WGSL.
• Attendez la tâche du dispatch avant de lancer la lecture.
• Un identifiant de relecture doit être libre. Supprimez d'abord l'ancien memblock.
• Le Sync() des boucles d'attente doit céder : DO … LOOP, END, ou une étiquette.
• Tampons de stockage seulement. Les paramètres vont dans un petit tampon à eux.