Quand votre PC ne vous sert pas à jouer, la carte graphique est un peu surdimensionnée. Alors qu'en faire ? La revendre est un peu laborieux, mais vous pouvez l'utiliser pour faire du calcul général: c'est le GPGPU (General-purpose computing on graphics processing units).
Attention! Je ne suis pas expert. Je découvre le sujet et écris ces lignes en vue de les améliorer au fil de ce que j'apprendrai sur le sujet. Tout commentaire bienvenu! Et si vous êtes anglophone vous trouverez plein de jolies choses ici: http://people.maths.ox.ac.uk/gilesm/cuda/
Pour que ceci soit accessible au plus grand nombre, je mets un peu de lexique en bas de page.
Comment coder sur GPU ?
CUDA est le langage le plus classique pour ça. Très proche du C, vous faites très vite un code qui écrit "hello world" et après vous pouvez décoller vers de vraies applications très rapidement.
Cela dit, c'est plus dur que de coder en C ou autre: en particulier, on ne peut pas faire un hello world super facilement (il faut une couche CPU et une couche GPU mise en place par le couche CPU), et le code produit est forcément modérément portable, puisqu'il ne marchera pas (sauf en émulation) sur une machine qui n'a pas de GPU...
La compilation est assez simple:
nvcc -o monBinaire moncode.cu
où ".cu" est l'extension des programmes cuda.
Quelques morceaux de code ?
Voilà un exemple:
- cudaMalloc(): pour allouer de la mémoire sur GPU.
- cudaMemcpy(): pour copier de la mémoire du GPU vers le CPU.
- cudaFree(): libérer un bloc mémoire sur le GPU.
- maFonction<<<nbBlocs, nbThreads>>>( mes arguments);
execute maFonction sur nbThreads dans chacun des nbBlocs. - __shared__ int maVariable: pour une variable partagée entre tous les threads d'un bloc.
- __syncThreads: on attend que tous les threads soient arrivés à cette ligne pour pouvoir continuer.
La page http://llpanorama.wordpres
Et pour débugguer / optimiser un peu ?
Pour trouver les erreurs, avoir des printf est drôlement commode.
Or, pas de "printf" sur GPU. Ouf! Compiler avec
nvcc --device-emulation -o monexecutable moncode.cu
permet d'avoir un executable émulé, et là les printf sont autorisés.
Votre premier programme devrait être lent, c'est normal.
Comment accélérer ?
- rendre local à chaque thread toutes les lectures mémoires possibles. C'est là que l'on gagne le plus vite et le plus facilement.
- rendre local à chaque bloc toues les lectures mémoires possibles. Ce n'est pas pareil que ci-dessus! Il faut veiller à lire "en interne au thread", dès que possible, et en interne au bloc sinon, et seulement en dernier recours revenir à la mémoire du CPU.
Pour aller plus loin, il y a des profileurs CUDA qui trouveront où vous perdez du temps.
Et l'installation de CUDA pour coder sur GPU ?
Ca implique
d'avoir récupéré les ".run" (pour Linux) ici:
http://developer.nvidia.co
et de les installer:
chmod +x cudadriver_2.3_linux_64_19
(ou cf http://doc.ubuntu-fr.org/c
Forcément, c'est moins "tout inclus" que du C ou C++ ou java
où tout est par défaut dans linux, mais j'imagine que ça viendra.
Est-ce un outil générique de calcul parallèle ?
http://people.maths.ox.ac.
Vu comme ça ça fait envie, mais ça n'est pas si puissant que ça pour tout.
- La mémoire est d'une part sur le GPU central et d'autre part sur chacun des blocs.
- L'accès à la mémoire est:
- assez lent dans tous les cas; au sein d'un bloc, il faut éviter d'avoir trop de threads qui accèdent simultanément à la mémoire d'un bloc (limite aux alentours de 16, sinon il y a ralentissement).
- très lent pour ce qui est du transfert de la mémoire "centrale" de l'ordinateur vers la mémoire du GPU.
- limité aussi de chaque bloc vers la mémoire centrale.
- Le GPU est composé d'environ 128 blocs qui peuvent calculer indépendamment (mais avec un accès à la mémoire globale limitée).
- Au sein d'un bloc, on travaille essentiellement en "stream processing", sur flux: c'est-à-dire qu'on ne fait pas (pas trop) de branchements qui ne donnent pas la même branche. Idéalement on peut avoir parait-il 512 threads par bloc, mais en pratique ils vont forcément se marcher un peu dessus.
- Petite mémoire sur chaque bloc; de l'ordre de 16Ko; ça va vite!
- Le calcul flottant est (sur beaucoup de GPU) en 32bits, alors que sur les CPU le 64bits se répand. Cela dit pour beaucoup de gens, 32bits ça suffit.
Méfiez-vous des comparatifs de puissance de calcul qui oublient souvent le temps de transfert vers la mémoire du GPU!
Alors, qu'est-ce-qui se programme bien sur GPU ?
Deux niveaux de parallélisme: 128 blocs, avec des threads plutôt SIMD sur chaque bloc.
Le cas idéal: beaucoup de donner à traiter, sans avoir des branchements sur chaque donnée.
Et ce que j'en ai vu ?
Les gens trichent parfois! Ils disent qu'ils ont un speed-up X ou Y mais ils ne prennent pas en compte les transferts mémoire. Obtenir des super-grosses
accélérations n'est pafs trivial. Sur le papier on peut atteindre des chiffres mirifiques avec 128 blocs de 512 threads, mais les threads d'un bloc ont un accès très limité à la mémoire et il est donc rare que l'on puisse atteindre des chiffres pareils. Les blocs sont supposés pouvoir travailler à peu près indépendamment, mais les premiers essais que l'on a fait ici n'étaient pas fantastiques.
Les jeunes centraliens qui par ici se sont lancés dans le codage de MCTS
(http://teytaud.over-blog.c
Glossaire.
- Cuda: le langage le plus classique pour coder sur GPU.
- GPU: processeur graphique! le truc qui fait que les jeux vidéos vont vite.
- GPGPU: le fait d'utiliser le GPU pour faire autre chose que du graphique.
-
Parallélisme: le fait d'effectuer un grand nombre d'opérations en même temps.
- MIMD: parallélisme permettant d'effectuer un grand nombre d'instructions différentes à un grand nombre de données différentes ("multiple instructions multiple data").
- SIMD: parallélisme basé sur le fait d'appliquer une même instruction à plein de données différentes en même temps ("single instruction multiple data"). Le GPU peut faire du MIMD, mais au prix d'un très fort ralentissement; il convient d'être "aussi SIMD" que possible dans du code GPU.
Un échantillon provenant de http://llpanorama.wordpres
// example1.cpp : Defines the entry point for the console application.
//
#include "stdafx.h"
#include <stdio.h>
#include <cuda.h>
// Kernel that executes on the CUDA device
__global__ void square_array(float *a, int N)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx<N) a[idx] = a[idx] * a[idx];
}
// main routine that executes on the host
int main(void)
{
float *a_h, *a_d; // Pointer to host & device arrays
const int N = 10; // Number of elements in arrays
size_t size = N * sizeof(float);
a_h = (float *)malloc(size); // Allocate array on host
cudaMalloc((void **) &a_d, size); // Allocate array on device
// Initialize host array and copy it to CUDA device
for (int i=0; i<N; i++) a_h[i] = (float)i;
cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
// Do calculation on device:
int block_size = 4;
int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
square_array <<< n_blocks, block_size >>> (a_d, N);
// Retrieve result from device and store it in host array
cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
// Print results
for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
// Cleanup
free(a_h); cudaFree(a_d);
}
