High Performance Computing

Data Reorganization Patterns

Gather: crea un insieme di dati (output) leggendo da un altro insieme di dati (input), dato un insieme di indici legge i dati dalla collezione di ingresso ad ogni indice e scrive i dati nella collezione di uscita in ordine di indice;

Shifts: sposta i dati a sinistra o a destra nella memoria, gli accessi ai dati sono sfalsati di distanze fisse;

Zip: La funzione è quella di interfogliare i dati;

Unzip: Inverso dello zip, estrae i sub-array a determinati offset e passi da un array di input;

Scatter: Data una collezione di dati di input e una collezione di indici di scrittura, sparge (scatter) i dati nella collezione di output. Scatter può causare race condition, abbiamo bisogno di regole per risolvere le collisioni:

Algoritmo Pack: usato per eliminare gli elementi inutilizzati da una collezione, gli elementi conservati vengono spostati in modo che siano contigui nella memoria, in input un vettore e un array di 0 e 1;

Algoritmo di Unpack: dati

gli stessi dati su quali elementi sono stati mantenuti e quali sono stati scartati, sparge gli elementi di nuovo nelle loro posizioni originali (quando incontra uno 0 mette una posizione vuota ma non scarta l'elemento, viene messo al prossimo 1);

Split: generalizzazione del modello pack, gli elementi sono spostati nella metà superiore o inferiore della collezione di output in base a qualche stato, non perde informazioni come il pack (nel vettore di output mette prima gli elementi associati con uno 0 e poi gli elementi associati con 1);

Unsplit: crea un insieme di output basato sull'insieme originale di input (inversa della split);

Bin: è una split con più categorie (ci sono numeri maggiori di 1)

3.2.1. Array of Structures (AoS) vs Structure of Arrays (SoA)

AoS: layout più logico, difficile per (gather) letture e (scatter) scritture; difficile per la vectorization, può portare a un migliore utilizzo della cache se i dati sono acceduti in modo

casuale- SoA: separare ogni campo della struttura, mantiene l'accesso alla memoria contiguo se si accede a più istanze della struttura, spesso migliore per la vectorization e per evitare false sharing 3.2.2. Intel AVX Molte applicazioni devono riorganizzare gli elementi vettoriali per assicurare che le operazioni siano eseguite correttamente. AVX/AVX2 fornisce una serie di funzioni intrinseche per questo scopo. Due categorie principali: funzioni permute, funzioni shuffle - L'intrinseco permute accetta due argomenti, un vettore di input e un valore di controllo a 8 bit, i bit del valore di controllo determinano quale degli elementi del vettore di ingresso viene inserito nell'uscita. - L'intrinseca shuffle seleziona elementi da uno o due vettori di ingresso e li mette nel vettore di uscita in base ai bit del valore di controllo. Autovectorization and VLA Programming 3.3. Automatic vectorization I compilatori possono cercare di utilizzare automaticamente le

istruzioni SIMD, questa ottimizzazione è anche chiamata autovectorization. La vettorizzazione automatica è un caso speciale di parallelizzazione automatica dove un programma viene convertito da un'implementazione scalare.

Ci sono due strategie principali:

• LLV: vettorizzazione a livello di loop
• SLP: vettorizzazione a livello di superword

LLV vectorization

Loop vectorization è una trasformazione del compilatore che cerca di utilizzare le istruzioni vettoriali sfruttando le iterazioni (multiple) dei loop.

Loop tails è un fattore di vettorizzazione: numero di iterazioni vettorializzate in un ciclo.

Esempio: Registro AVX (256bit), loop su float (32 bit) -> fattore di vettorizzazione 8

Problema: la dimensione del loop può non essere un multiplo del fattore di vettorizzazione.

Vectorization with Non-contiguous Memory Access: quattro interi consecutivi o valori in virgola mobile possono essere caricati direttamente dalla memoria in una singola istruzione SSE:

ma se i quattro interi non sono adiacenti, devono essere caricati separatamente usando istruzioni multiple, il che è notevolmente meno efficiente. La vettorizzazione è spesso possibile, ma può essere inefficiente.

Ignoring Dependencies: il compilatore spesso assume che regioni di memoria accedute da puntatori potrebbero sovrapporsi, ciò comporta l'impedimento della vectorization. Possiamo istruire il compilatore di ignorare le dipendenze tramite #pragma ivdep

La vectorization in OpenMP si istruisce tramite #pragma omp simd [clause]

SLP Vectorization

SLP: parallelismo a livello di superword (Superword-Level Parallelism): combina istruzioni indipendenti simili in istruzioni vettoriali, vettorizzando accessi alla memoria, operazioni aritmetiche, operazioni di confronto.

Roofline Model

Modello di performance visivamente intuitivo costruito usando l'analisi dei limiti e dei bottleneck è usato per guidare i programmatori verso una comprensione intuitiva

delle prestazioni sulle moderne architetture di computer, enumerando i potenziali impedimenti delle performance utili per implementare le giuste ottimizzazioni.

Prestazioni

Le prestazioni sono determinate da:

• Architettura (NUMA, SMP) e dal kernel (intensità aritmetica dell'applicazione).

Ci sono tre componenti principali per le prestazioni:

1. Calcolo
2. La comunicazione
3. Località

Ogni architettura ha un diverso equilibrio tra queste componenti, ogni kernel ha un diverso equilibrio tra queste componenti.

Calcolo

La metrica comune è la performance in virgola mobile (Gflop/s).

Comunicazione

Si misura in larghezza di banda di DRAM (GB/s).

Località

Massimizzare la località per minimizzare la comunicazione.

Roofline Model

Obiettivo: integrare le prestazioni in-core, la larghezza di banda della memoria e la localizzazione in una singola figura di prestazioni facilmente comprensibile.

• Il modello Roofline è unico per ogni architettura.
• Le coordinate di...

un kernel sono circa uniche per ogni architettura- Costruito usando un'analisi "bound and bottleneck"- Relaziona i GB/s ai GFlop/s

Il grafico è composto da:

• Peak Single precision: limite di performance del processore;
• Peak stream bandwith: limite della banda di memoria;
• Mul/add imbalance: ci indica quando non stiamo utilizzando a pieno le unità di moltiplicazione e addizione all'interno del processore;
• w/out SIMD: risultati delle performance senza utilizzare la vectorization;
• w/out ILP: risultati delle performance quando non sfruttiamo al massimo il parallelismo al livello di istruzioni.
• w/out SW prefetching: limite della larghezza di banda quando non vengono utilizzate ed inserite delle addizionali istruzioni di prefetch;
• w/out NUMA optimization: limite della larghezza di banda in un'architettura NUMA (quando vengono usati meno controller di memoria di quelli disponibili);

Intensità aritmetica in HPC

Alcuni kernel HPC hanno

può essere acceduta solo dallo stesso workgroup e costante che contiene dati costanti che possono essere letti da tutti i threads. Le GPU sono progettate per eseguire operazioni parallele su grandi quantità di dati, quindi sono particolarmente adatte per applicazioni che richiedono calcoli intensivi come il rendering grafico, l'apprendimento automatico e la simulazione scientifica.

acceduta solo dallo stesso blocco (work-group) di threads e privata acceduta dal singolo thread.

Barrier synchronization tra thread dello stesso gruppo: in OpenCL barrier(), da usare ad esempio prima di leggere la memoria locale scritta da un altro thread.

Il problema può avere fino a 3 dimensioni 1D, 2D, 3D e ad ogni work item sarà associato un singolo punto nello spazio di iterazione.

OpenCL memory model

• Private memory: per ogni work-item
• Local memory: condivisa nello stesso work-group
• Global memory: visibile a tutti i work-group
• Host memory: sulla CPU

OpenCL Platform Model

• Contesto: ambiente in cui viene eseguito il kernel, ed include uno o più dispositivi, la loro memoria e una o più code per i comandi;
• I comandi vengono immessi attraverso la coda dei comandi (es: esecuzione kernel, sincronizzazione, trasferimenti dei dati tra le memorie);
• Ogni coda dei comandi è dedicata al singolo dispositivo;

4.1. GPU memory

La GPU per usare a pieno il bus combina

consistente solo alla fine dell'esecuzione del kernel. Per garantire la consistenza dei dati tra i work-item di un work-group, è necessario utilizzare le funzioni di sincronizzazione fornite da OpenCL, come ad esempio la funzione barrier(). Inoltre, è possibile utilizzare la memoria condivisa (shared memory) per migliorare le prestazioni, in quanto questa memoria è più veloce rispetto alla memoria globale. La memoria condivisa viene allocata nello spazio di indirizzamento locale di ogni work-group e può essere utilizzata per condividere dati tra i work-item di uno stesso work-group.

consistente da una barrier. La memoria globale è consistente nello stesso work-group grazie ad una barrier ma non è garantito tra differenti work-group.

OpenCL Atomics: operazioni che in un qualsiasi istante di tempo possono essere o completate o non ancora eseguite (garantiscono anche che i thread non interferiscano tra di loro).

Per migliorare gli accessi alla memoria si utilizza memoria locale, tiling e SoA.

4.2. GPU Reduction

Divergent Control Flow: si intende quando un gruppo di thread devono eseguire cammini diversi (es: if-else). Per come è strutturata la GPU (SIMT), è costretta a fare seguire entrambi i cammini a tutti i thread, però disattivando il gruppo di thread che dovrebbe eseguire l'altro percorso.

Occupancy: è una misura di quanto attivo stiamo tenendo attivo il processing element (singolo work-item).

Synchronization: è quando più work-item vengono portati ad un punto noto durante la loro