Notes De Révision Du Cours De Calcul Parallèle

1 Plan de révision

Chapitre 1 Loi d’Amdahl Compréhension de la loi (augmentation de la vitesse, rapport d’accélération, limite d’accélération)

Problème d’application 6’*5

Disposition des grilles et des blocs de threads, calcul de l’ID global Parallèle, concurrent, warp, ID global, CPU multicœur et GPU multicœur

Analyse de programme 10*2

Écrire le résultat du code, analyser pourquoi ce résultat se produit

CPU multicœur 10*2

Division des données : préciser la plage de données traitée par chaque partie Parallélisme des tâches : expérimentation du pool de threads

Programmation CUDA 15*2

Problème spécifique, conception de la grille et des blocs de threads, ou si les blocs de threads sont donnés, il suffit de concevoir la grille ; Processus fixe dans la fonction principale ; l’essentiel est d’écrire la fonction noyau ;

2 Calcul parallèle

2.1 Concurrence et parallélisme

Série : machine unique, cœur unique, exécution séquentielle des instructions.

Concurrence : machine unique, cœur unique, exécution parallèle des instructions dans le temps, se produisant dans le même intervalle de temps.

Parallélisme : machine unique, multicœur, machine multiple, mono/multicœur, exécution parallèle des instructions dans l’espace, se produisant au même moment.

Le calcul parallèle est un supercalcul effectué sur des ordinateurs parallèles ou des systèmes distribués, etc., des systèmes de calcul haute performance. Le calcul parallèle peut réduire le temps de résolution d’un problème unique, augmenter l’échelle et la précision de la résolution, améliorer le débit, etc.

Trois classifications :

• Mode de calcul : parallélisme temporel (pipeline), parallélisme spatial (multiprocesseur)
• Logique du programme : parallélisme des tâches, parallélisme des données
• Angle d’application : intensif en calcul, intensif en données, intensif en réseau

2.2 Classification de Flynn

Une méthode de classification de l’architecture des ordinateurs parallèles basée sur le mode d’exécution des flux d’instructions et de données. Comprend SISD (machine sérielle précoce), SIMD (ordinateur monocœur), MISD (rarement utilisé), MIMD (ordinateur multicœur, parallèle);

2.3 Loi d’Amdahl

Suppose un nombre de tâches constant, révèle qu’une partie d’un programme qui ne peut pas être parallélisée limitera l’amélioration des performances globales du programme. $$S=\frac{W_{s}+W_{p}}{W_{s}+W_{p}/p}$$ où $W_{s}$ est le nombre de tâches sérielles, $W_{p}$ est le nombre de tâches parallèles, $p$ est le nombre de processeurs, $S$ est le rapport d’accélération. En fonction de la proportion de la composante sérielle $f=W_{s}/W$, en divisant l’équation ci-dessus par $W$, on obtient l’équation suivante : $$S=\frac{f+(1-f)}{f+\frac{1-f}{p}} =\frac{p}{1+f(p-1)}$$ $\lim_{x\rightarrow \infty}S=1/f$, lorsque le nombre de processeurs augmente indéfiniment, le rapport d’accélération que le système peut atteindre est limité par la partie sérielle du programme.

1
2
3

1. Dans une application sérielle, 20% doit être exécuté en série. Pour réaliser une amélioration de performance de 3 fois, combien de CPU sont nécessaires ? Pour réaliser un rapport d'accélération de 5 fois, combien de CPU sont nécessaires ?
2. Un programme parallèle fonctionnant sur 5 ordinateurs a 10% de partie parallèle. Par rapport à l'exécution sérielle sur un ordinateur, quel est le rapport d'accélération ? Si nous voulons doubler le rapport d'accélération, combien de CPU sont nécessaires ?
3. Modifier un programme dont la partie non parallélisable représente 5% en un programme parallèle. Il y a deux types d'ordinateurs parallèles sur le marché : l'ordinateur X a 4 CPU, chaque CPU peut exécuter le programme en 1 heure ; l'ordinateur Y a 16 CPU, chaque CPU peut exécuter le programme en 2 heures. Si vous devez minimiser le temps d'exécution, quel ordinateur devriez-vous acheter ?

3 Vue d’ensemble de CUDA

3.1 Calcul hétérogène

Le calcul parallèle sur GPU est un calcul hétérogène, divisé en côté hôte (CPU) et côté appareil (GPU), leur relation n’a jamais été égale, CUDA doit même spécifier où le code doit être exécuté.

3.2 Différences entre CPU et GPU

Intuitivement, le CPU utilise plus de ressources pour le cache et le flux de contrôle, tandis que le GPU se concentre davantage sur le calcul des données.

1. Dans un environnement GPU, le cœur du GPU est responsable de l’exécution de toutes les tâches de calcul, mais les instructions de travail proviennent toujours du CPU.
2. Dans le cas du GPU, le cœur du GPU ne récupère jamais les données par lui-même, les données proviennent toujours du côté CPU, et les résultats de calcul sont renvoyés au côté CPU. Par conséquent, le GPU joue en arrière-plan le rôle d’accélérateur de calcul, effectuant certaines tâches externalisées pour le CPU.
3. Ce type d’architecture de système n’est très efficace que lorsqu’il y a un grand nombre d’unités de traitement parallèles, et non seulement 2 ou 4.
4. Le concept de warp a un impact significatif sur l’architecture du GPU. Les données doivent être entrées dans le GPU par blocs de données de même taille, un bloc de données est un demi-warp, soit 16 éléments.
5. Le fait que les données doivent être transférées au cœur du GPU par blocs de taille d’un demi-warp signifie que le système de stockage responsable de l’entrée des données dans le GPU doit entrer 16 données à chaque fois. Cela nécessite un sous-système de stockage parallèle capable de transférer 16 nombres à la fois. C’est pourquoi la mémoire DRAM du GPU est constituée de DDR5, car c’est une mémoire parallèle.
6. Étant donné que le cœur du GPU et le cœur du CPU sont des unités de traitement complètement différentes, on peut prévoir qu’ils ont des ISA (architecture de jeu d’instructions) différentes. Autrement dit, ils parlent des langues différentes.

Les threads GPU diffèrent également des threads CPU, avec des coûts de création extrêmement faibles. Le CPU réduit la latence grâce à un cache multi-niveaux, tandis que le GPU réduit la latence en augmentant le débit grâce au pipeline. En raison de leurs objectifs de conception différents, le CPU a besoin d’une grande polyvalence pour traiter divers types de données, tandis que le GPU traite des données massives et uniformes sans dépendance dans un environnement de calcul pur et sans interruption.

3.3 Organisation des threads CUDA

Thread : unité de base du parallélisme

Thread Block : groupe de threads coopérants permettant de se synchroniser mutuellement via une mémoire partagée rapide, organisé en 1D, 2D ou 3D, contenant jusqu’à 1024 threads

Grid : ensemble de blocs de threads organisé en 1D, 2D ou 3D, partageant des variables globales

Kernel : programme central exécuté sur le GPU One kernel One Grid

3.4 Modèle de programmation hôte/appareil CUDA

3.4.1 Modificateurs de fonction

• _ device _: exécuté sur l’appareil et ne peut être appelé que depuis l’appareil, utilisé comme sous-fonction sur l’appareil.
• _ host _: exécuté sur l’hôte et ne peut être appelé que depuis l’hôte, identique à une fonction C normale. Ne peut pas être utilisé avec global, mais peut être utilisé avec device, auquel cas la fonction sera compilée à la fois sur l’appareil et l’hôte.
• _ global _ : fonction noyau, exécutée sur l’appareil mais appelée depuis l’hôte.

3.4.2 Limitations des fonctions noyau CUDA

1. Ne peut accéder qu’à la mémoire de l’appareil
2. Doit retourner un type void
3. Ne prend pas en charge un nombre variable de paramètres
4. Ne prend pas en charge les variables statiques
5. Comportement asynchrone affiché, ce qui signifie que l’hôte n’attendra pas que le noyau soit exécuté pour exécuter l’étape suivante

3.5 Modèle de calcul parallèle SIMT

Le bloc de threads est l’unité de lancement du programme, le warp est l’unité d’exécution du programme ;

Par exemple, lorsque nous disons qu’un bloc a une taille de 256 threads, cela signifie que la taille du bloc de threads est de 8 warps. Chaque warp contient toujours 32 threads. Ce paramètre indique que bien que le programme soit lancé avec 256 threads par bloc, cela ne signifie pas qu’ils seront exécutés immédiatement, c’est-à-dire que ces 256 threads ne seront pas tous exécutés ou terminés en même temps. Au contraire, le matériel d’exécution du GPU utilisera 8 warps pour exécuter ces threads.

SIMT appartient à la catégorie SIMD car il exécute également la même instruction sur plusieurs données. Cependant, SIMT permet à l’utilisateur d’allouer des threads, en particulier CUDA attribue un identifiant à chaque thread.

Une différence clé est que SIMD exige que tous les éléments d’un même vecteur soient exécutés ensemble dans un groupe de synchronisation unifié, tandis que SIMT permet à plusieurs threads appartenant au même warp d’être exécutés indépendamment, ces threads pouvant avoir des comportements différents. Par conséquent, SIMT permet la concurrence au niveau du thread, c’est-à-dire que les threads sous un même warp peuvent faire des choses différentes en même temps.

Trois différences :

• Chaque thread a son propre compteur d’adresse d’instruction
• Chaque thread a son propre état de registre
• Chaque thread peut avoir un chemin d’exécution indépendant

3.6 Architecture GPU

Multiprocesseur de flux SM

Un bloc de threads ne peut être planifié que sur un SM, mais un SM peut correspondre à plusieurs blocs de threads.

Lorsqu’un SM a spécifié un ou plusieurs blocs de threads à calculer, ces blocs de threads sont divisés en plusieurs warps, en attente de planification.

Les threads dans un warp exécutent la même commande sur différentes données.

Le nombre de blocs de threads qu’un SM peut contenir dépend de la mémoire partagée et des registres dans le SM ainsi que des ressources occupées par les threads.

Tous les threads dans un bloc de threads fonctionnent logiquement en parallèle, mais tous les threads ne peuvent pas être exécutés en même temps au niveau physique. (Un SM ne planifie qu’un warp à la fois, les autres warps attendent, le changement entre différents warps est sans coût car le contexte d’exécution du warp est maintenu par le SM pendant toute la durée de vie du warp)

Les capacités de calcul de la NVIDIA GeForce RTX 3090 sont de 8.6, elle contient 82 SM, chaque SM permettant un maximum de 1536 threads simultanés, calculer : combien de threads peuvent être exécutés en parallèle en même temps ? Combien de threads peuvent être exécutés en concurrence ? Elle contient 82 SM, chaque SM permettant un maximum de 1536 threads simultanés, soit un maximum de 48 warps, car les warps sont exécutés en concurrence par le planificateur de warps, 32 threads dans un warp sont exécutés en parallèle, on peut donc dire grossièrement que le nombre de threads exécutés en parallèle en même temps est de 82*32=2624, et le nombre de threads exécutés en concurrence est de 82*32*48=125952.

3.7 Modèle de mémoire

Quand les variables définies dans une fonction noyau CUDA sont-elles des variables de registre et quand sont-elles des variables locales ?

Les trois cas suivants sont des variables locales, sinon ce sont des variables de registre

• Tableau dont la taille ne peut être déterminée à la compilation
• Tableau ou structure occupant un espace important
• De nombreuses variables définies dans la fonction noyau, les registres ne peuvent pas tout contenir

Les débordements de registre vers la mémoire locale sont en fait dans la même zone de stockage que la mémoire globale

3.8 Mode d’accès à la mémoire

La mémoire globale est chargée/stocker via le cache. Tous les accès à la mémoire globale passent par le cache L2 (généralement 128 octets).

Accès aligné

La première adresse est un multiple pair de la granularité du cache (généralement 32 octets) (la ligne de cache commence à l’adresse de début des données)

Accès fusionné

Tous les threads d’un warp accèdent à un bloc de mémoire contigu. L’accès fusionné signifie qu’une requête d’accès à la mémoire globale par un warp entraîne le moins de transferts de données (fusion = 100%), sinon c’est un accès non fusionné.

5 modes d’accès, calcul de la fusion ??

Si la lecture et l’écriture ne peuvent pas être fusionnées, il est préférable de garantir une écriture fusionnée. L’accès non fusionné aux données en lecture seule peut être mis en cache avec la fonction __ldg(), ou converti en accès fusionné avec la mémoire partagée.

3.9 Mémoire partagée et conflit de banque

La mémoire partagée peut être manipulée directement par le programmeur. La mémoire partagée est divisée en de nombreuses banques.

• Tous les threads d’un warp accèdent à la même banque à la même adresse - diffusion
• Différents threads d’un warp accèdent à différentes adresses d’une banque - conflit de banque
• Plusieurs threads accèdent à la même banque à la même adresse - multidiffusion

Memory Padding Remplissage de mémoire pour résoudre le conflit de banque

Opération de remplissage : ajouter +1 à la deuxième dimension de sData, c’est-à-dire sData[BS][BS+1] La partie remplie ne peut pas être utilisée pour le stockage de données, ce qui réduit la quantité de mémoire partagée disponible.

4 Code

4.1 Inversion d’image CPU

Inversion d’image multithread, tout en maintenant manuellement le cache.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49

void * MTFlipHM(void * tid){
        struct Pixel pix; //temp swap pixel
        int row, col;
        int id = *((int *) tid);
        int start = id * ip.Vpixels/NumThreads;
        int end = start + ip.Vpixels/NumThreads;
        unsigned char buffer[16384];
        for (row = st ; row < ed; row++)
        {
            memcpy(buffer, TheImage[row], ip.Hbytes);
            col = 0;
            while (col < ip.Hpixels * 3 /2){
            pix.B = buffer[col];
            pix.G = buffer[col+1];
            pix.R = buffer[col+2];
            buffer[col]   = buffer[ip.Hpixels*3 - col -3];
            buffer[col+1] = buffer[ip.Hpixels*3 - col -2];
            buffer[col+2] = buffer[ip.Hpixels*3 - col -1];
            buffer[ip.Hpixels*3 - col -3] = pix.B;
            buffer[ip.Hpixels*3 - col -2] = pix.G;
            buffer[ip.Hpixels*3 - col -1] = pix.R;
            col += 3;
            }
        }
        memcpy(TheImage[row],buffer,ip.Hbytes);
        pthread_exit(NULL);
}

void * MTFlipVM(void * tid){
        struct Pixel pix; //temp swap pixel
        int row, col;
        int id = *((int *) tid);
        int start = id * ip.Vpixels/NumThreads;
        int end = start + ip.Vpixels/NumThreads;
        unsigned char buffer1[16384], buffer2[16384];
        for (row = start; row < end; row ++)
        {
            memcpy(buffer1,TheImage[row],ip.Hbytes);
            int  mirrow = ip.Vpixels - 1 - row;
            memcpy(buffer2,TheImage[mirrow],ip.Hbytes);
            // Copie décalée pour terminer l'échange
            memcpy(TheImage[row],buffer2,ip.Hbytes);
            memcpy(TheImage[mirrow],buffer1,ip.Hbytes);
        }
}

pthread_create(&ThHandle[i], &ThAttr, MTFlipFunc, (void *)&ThParam[i]);
for(i=0; i<NumThreads; i++)
                pthread_join(ThHandle[i], NULL);

4.2 Addition de tableaux

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28

const int a=1, b=2, c=3;
__global__ void add(double *x, double *y, double* z){
	const int n = blockIdx.x * blockDim.x + threadIdx.x;
	if (n<N)  z[n] = x[n] + y[n];
}
int main(){
	const int N = 1e9;
	const int M = sizeof(double) * N;
	double *h_x = (double*) malloc(M);
	double *h_y = (double*) malloc(M);
	double *h_z = (double*) malloc(M);
	for (int i=0; i<N; i++)
	{
		h_x[i] = a;
		h_y[i] = b;
	}
	double *d_x, *d_y, *d_z;
	cudaMalloc((void**) &d_x, M);
	cudaMalloc((void**) &d_y, M);
	cudaMalloc((void**) &d_z, M);

	const int block_size =  128;
	int grid_size = (N+block_size-1) / block_size;
	add<<<grid_size, block_size>>>(d_x, d_y, d_z);
	cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
	free(h_x),free(h_y),free(h_z);
	cudaFree(d_x),cudaFree(d_y),cudaFree(d_z);
}

4.3 Inversion d’image

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38

__global__ void Vflip(uch *ImgDst, uch *ImgSrc, ui Hpixels, ui Vpixels){
ui ThrPerBlk = blockDim.x;
ui MYbid = blockIdx.x;
ui MYtid = threadIdx.x;
ui MYgtid = ThrPerBlk * MYbid + MYtid;
ui BlkPerRow = (Hpixels + ThrPerBlk - 1) / ThrPerBlk; // ceil
ui RowBytes = (Hpixels * 3 + 3) & (~3);
ui MYrow = MYbid / BlkPerRow;
ui MYcol = MYgtid - MYrow*BlkPerRow*ThrPerBlk;
if (MYcol >= Hpixels) return;// col out of range
ui MYmirrorrow = Vpixels - 1 - MYrow;
ui MYsrcOffset = MYrow * RowBytes;
ui MYdstOffset = MYmirrorrow * RowBytes;
ui MYsrcIndex = MYsrcOffset + 3 * MYcol;
ui MYdstIndex = MYdstOffset + 3 * MYcol;
// swap pixels RGB @MYcol , @MYmirrorcol
ImgDst[MYdstIndex] = ImgSrc[MYsrcIndex];
ImgDst[MYdstIndex + 1] = ImgSrc[MYsrcIndex + 1];
ImgDst[MYdstIndex + 2] = ImgSrc[MYsrcIndex + 2];}

__global__ void Hflip(uch *ImgDst, uch *ImgSrc, ui Hpixels){
ui ThrPerBlk = blockDim.x;
ui MYbid = blockIdx.x;
ui MYtid = threadIdx.x;
ui MYgtid = ThrPerBlk * MYbid + MYtid;
ui BlkPerRow = (Hpixels + ThrPerBlk -1 ) / ThrPerBlk; // ceil
ui RowBytes = (Hpixels * 3 + 3) & (~3);
ui MYrow = MYbid / BlkPerRow;
ui MYcol = MYgtid - MYrow*BlkPerRow*ThrPerBlk;
if (MYcol >= Hpixels) return;// col out of range
ui MYmirrorcol = Hpixels - 1 - MYcol;
ui MYoffset = MYrow * RowBytes;
ui MYsrcIndex = MYoffset + 3 * MYcol;
ui MYdstIndex = MYoffset + 3 * MYmirrorcol;
// swap pixels RGB @MYcol , @MYmirrorcol
ImgDst[MYdstIndex] = ImgSrc[MYsrcIndex];
ImgDst[MYdstIndex + 1] = ImgSrc[MYsrcIndex + 1];
ImgDst[MYdstIndex + 2] = ImgSrc[MYsrcIndex + 2];}

4.4 Transposition de matrice

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15

__global__ void transpose(int a[], int b[], int N){
    // Allouer mémoire partagée
    __shared__ int S[TILE][TILE + 1];
    int bx = blockIdx.x * TILE;
    int by = blockIdx.y * TILE;
    int ix = bx + threadIdx.x;
    int iy = by + threadIdx.y;
    if (ix < N && iy < N)// Lire dans la mémoire partagée
        S[threadIdx.y][threadIdx.x] = a[iy * N + ix];
    __syncthreads();// Synchroniser, c'est indispensable
    int ix2 = bx + threadIdx.y;
    int iy2 = by + threadIdx.x;
    if (ix2 <N && iy2 < N)// Écrire
        b[ix2 * N + iy2 ] = S[threadIdx.x][threadIdx.y];
}

4.5 Multiplication de matrices carrées

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19

__shared__ float Mds[WIDTH][TILE_WIDTH];
__shared__ float Nds[TILE_WIDTH][WIDTH];
int bx=blockIdx.x ; int by = blockIdx.y;
int tx=threadIdx.x ; int ty = threadIdx.y;
int Row = by *TILE_WIDTH +ty;
int Col = bx*TILE_WIDTH + tx;
float Pvalue = 0;
for(int m=0; m<WIDTH/TILE_WIDTH; ++m)
    {
    // Chaque thread charge un élément de la sous-matrice de M
    Mds[ty][tx] = Md[Row*width+(m*TILE_WIDTH+tx)];
    // Chaque thread charge un élément de la sous-matrice de N
    Nds[ty][tx] = Nd[(m*TILE_WIDTH+ty)*width+Col];
	 __syncthreads();
	 for (int k = 0; k < TILE_WIDTH; ++k)
	    Pvalue += Mds[ty][k] * Nds[k][tx];
	 __syncthreads();
	 }
Pd[Row*WIDTH+Col] = Pvalue;// Écrire le résultat dans la matrice P

4.6 Histogramme

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25

#define SIZE (100*1024*1024)
// Générer un flux d'octets aléatoire à l'aide de la fonction utilitaire big_random_block()
unsigned char *buffer =(unsigned char*)big_random_block( SIZE );
unsigned int histo[256];
for (int i = 0; i<256; i++)
	histo[i] = 0;
for (int i = 0; i < SIZE; i++)
	histo[buffer[i]]++;
long histoCount = 0;
for (int i = 0; i<256; i++) {
	histoCount += histo[i]; }

__global__ void histo_kernel(unsigned char *buffer, long size, unsigned int *histo){
__shared__ unsigned int temp[256];
temp[threadIdx.x] = 0;
__syncthreads();
int i = threadIdx.x + blockIdx.x * blockDim.x;
int offset = blockDim.x *gridDim.x;
while (i<size){
	atomicAdd(&temp[buffer[i]], 1);
	i += offset;
}
__syncthreads();
atomicAdd(&(histo[threadIdx.x]), temp[threadIdx.x]);
}

4.7 Réduction

La réduction et le problème TOP K sont similaires, le code ci-dessous est le code officiel, pour comprendre, voir cet article

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39

__global__ void _sum_gpu(int *input, int count, int *output)
{
    __shared__ int sum_per_block[BLOCK_SIZE];

    int temp = 0;
    for (int idx = threadIdx.x + blockDim.x * blockIdx.x;
         idx < count; idx += gridDim.x * blockDim.x
	)
    {// Boucle à travers la grille, un thread additionne plusieurs données, pour faire face à des données massives
        temp += input[idx];
    }

    sum_per_block[threadIdx.x] = temp;  //la somme partielle par thread est temp!
    __syncthreads();

    //**********étape de sommation en mémoire partagée***********
    for (int length = BLOCK_SIZE / 2; length >= 1; length /= 2)
    {
        int double_kill = -1;
	if (threadIdx.x < length)
	{
	    double_kill = sum_per_block[threadIdx.x] + sum_per_block[threadIdx.x + length];
	}
	__syncthreads();  //pourquoi avons-nous besoin de deux __syncthreads() ici, et,
	
	if (threadIdx.x < length)
	{
	    sum_per_block[threadIdx.x] = double_kill;
	}
	__syncthreads();  //....ici ?
	
    } //la somme partielle par bloc est sum_per_block[0]

    if (blockDim.x * blockIdx.x < count) //au cas où nos utilisateurs seraient espiègles
    {
        //la réduction finale effectuée par atomicAdd()
        if (threadIdx.x == 0) atomicAdd(output, sum_per_block[0]);
    }
}

4.8 TOP K

Le processus de mise en œuvre spécifique est le suivant :

1. Copier les données dans la mémoire vidéo du GPU float *d_data; cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
2. Stocker les données dans un tuple

1
2
3
4
5

typedef struct {     float value;     int index; } Tuple;  
Tuple *d_tuples; 
int threadsPerBlock = 256; 
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
initializeTuples<<<blocksPerGrid, threadsPerBlock>>>(d_data, d_tuples, n);

3. Effectuer une réduction sur les tuples pour obtenir les indices des K plus grands/petits

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17

int *d_indices;
kReduceKernel<<<blocksPerGrid, threadsPerBlock>>>(d_tuples, d_indices, n, k);
__global__ void kReduceKernel(Tuple *input, int *output, int n, int k) {
    extern __shared__ Tuple shared[];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    shared[tid] = (i < n) ? input[i] : Tuple{0, 0};
    __syncthreads();
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s)
            shared[tid] = (shared[tid].value > shared[tid + s].value) ? shared[tid] : shared[tid + s];
        __syncthreads();
    }

    if (tid == 0)
        output[blockIdx.x] = shared[0].index;
}

5. Récupérer les données originales dans le CPU et trier par indice pour obtenir les K plus grands/petits

1
2
3
4
5

cudaMemcpy(h_indices, d_indices, size, cudaMemcpyDeviceToHost);  
for (int i = 0; i < k; ++i) {     
    int index = h_indices[i];     
    h_result[i] = h_data[index]; }  
std::sort(h_result, h_result + k);

5 Expériences

Expérience 1 : Trois méthodes de calcul de PI, pool de threads

Expérience 2 : Multiplication de matrices, transposition, réduction, problème TOP K

Optimisation de la mémoire globale, optimisation de la mémoire partagée, optimisation des conflits de banque

Rédiger un rapport de calcul parallèle de 3 à 4 pages, trois méthodes de calcul de PI sur CPU, modèle producteur-consommateur, programme principal GPU clair une fois, fonction noyau différente, l’accent sur la mise en œuvre de la mémoire globale et partagée, multiplication de matrices histogramme (boucle à travers la grille) réduction de la valeur maximale dans un tableau de 100w, rapport à soumettre lors de l’examen

6 Examen de 2020

6.1 Questions courtes

Loi d’Amdahl, n processeurs, 40% en série, calculer la limite du rapport d’accélération.

Image RGB de 680*480, divisée en 4 threads (pas dit comment), plage de pixels et de bytes traitée par chaque thread ;

Exemple de calcul du nombre de threads parallèles et concurrents dans un warp à partir du PPT ;

Les capacités de calcul de la NVIDIA GeForce RTX 3090 sont de 8.6, elle contient 82 SM, chaque SM permettant un maximum de 1536 threads simultanés, calculer : combien de threads peuvent être exécutés en parallèle en même temps ? Combien de threads peuvent être exécutés en concurrence ? Elle contient 82 SM, chaque SM permettant un maximum de 1536 threads simultanés, soit un maximum de 48 warps, car les warps sont exécutés en concurrence par le planificateur de warps, 32 threads dans un warp sont exécutés en parallèle, on peut donc dire grossièrement que le nombre de threads exécutés en parallèle en même temps est de 82*32=2624, et le nombre de threads exécutés en concurrence est de 82*32*48=125952.

ID global d’un élément dans le processus de transposition de matrice ;

Demander l’ID global de l’élément 3

Peut-on copier des données tout en les transposant (flux CUDA) ? ;

6.2 Analyse de programme

Première question, question originale de l’exercice en classe ; <4,4> changé en <5,5> ; expliquer le processus ;

Deuxième question, histogramme de réduction sans opération atomique, quel est le problème ? ;

6.3 Programmation CPU

Trouver les nombres premiers dans le tableau a[2,1000000], 10 threads doivent être utilisés pour diviser également ;

Pseudo-code du pool de threads : client, serveur (fonction de messagerie, fonction d’exportation, statistiques de trafic, etc.) ;

6.4 Programmation GPU

Multiplication de matrices en mémoire globale ;

Produit scalaire des vecteurs a, b, dimension = 1024000000 ; blockdim.x = blockdim.y =16 fixé ; Concevoir la grille ; Optimiser avec la mémoire partagée, résoudre le conflit de banque, retourner le résultat au CPU pour la fusion finale.

Impression : questions non codées faciles, questions codées difficiles, impossible de tout écrire.