From 59f3ea02848204cdd0abd47dba553ff56d9fdbc3 Mon Sep 17 00:00:00 2001 From: LucaCorra02 Date: Wed, 28 Jan 2026 17:12:37 +0100 Subject: [PATCH 1/3] add raw note --- chapters/Lezione5.typ | 94 +++++++++++++++++++++++++++++++++++++++++++ main.typ | 5 ++- 2 files changed, 98 insertions(+), 1 deletion(-) create mode 100644 chapters/Lezione5.typ diff --git a/chapters/Lezione5.typ b/chapters/Lezione5.typ new file mode 100644 index 0000000..81b31f5 --- /dev/null +++ b/chapters/Lezione5.typ @@ -0,0 +1,94 @@ +#import "../template.typ": * + += Lezione 5 + +== Memory +Le cache sulla GPU possono essere : +- L1 (una per ogni SM) +- L2 (condivisa tra tutti gli SM) + +Le cache a volte sono fisicamente in un area della device memory. + +``` cudaFuncSetAttribute``` è una primitiva run time o compile time che permette di dividere in modo arbitrario la L1 cache, tra L1 e shared memory. + +Il numero massimo di blocchi ospitati all'interno di un SM può scendere se all'interno di un blocchi richiedo molte risorse (registre e shared memory), in quanto esse sono condivise nell'SM. Il numero massimo di blocchi non è mai fisso ma dipende dall'occupancy (tasso di warp attivi simltuanamente nell'SM). A compile time si viene a sapere il nuemro di blocchi operativi simultaneamente. + +== Shared memory + +Dobbiamo trasportare dei dati all'interno dei core, perchè ci metto di mezzo una memoria shared tra global memory (fuori dal chip) e registri? Ha senso usare la shared memory quanto il tasso di accesso del singolo dato è molto elevato ad esempio quando i thread di un certo blocco cooperano sul blocco di dati in shared memory. il beneficio è che la bandwidth della shared memory è molto altra rispetto a quella fuori dal chip. + +La shared memory viene caricata e usata intensamente dai thread di un blocco, i thread non vendono la shared memory degli altri blocchi (a meno di clustering). + +Barriera di sincronizzazione. Il trasferimento di dati dalla global memory alla shared memory è distrubuiti tra i thread del blocco, ogni thread trasporta un certo set di dati.\ +Tuttavia se il trasferimento è parallelo srve sincronizzazione per garantire che per iniziare la computazione tutti i dati sono stati trasferiti, abbiamo così dei dati afffidabili (no race condition). + +=== Access via banking + +La memoria shared è organizzata in 32 bank fisici (per la dimensione del warp). + +#nota()[ + La shared memory è tagliata sulla dimensione del warp. +] + +I 32 thread arrivano e prevelano i dati purche siano distribuiti in 32 banchi distinti. La dimensione dei banchi dipende dalla dimensione dei dati, solitamente $4$ byte(int), ovvero una word. I banchi solitamente implementato sequenze di word. + + +Pattern di accesso. Possiamo avere: +- Acesso parallelo, ogni thread accede ad un blocco differente +- Accesso sequenziale, si possono creare dei conflitti se diversi thred cercano di accedere allo stesso blocco. + +=== SMEM + +Se la dimensione non è nota a compile time è possibile dichiarare una variabile adimensionale con la keyword ``` extern```: +- dinamicamente si possono allocare solo array 1D +- può essere sia interna al kernl sia esterna +``` extern __shared__ int array[]```. Alloca un array nelal shared memory. + +Per allocare la shared memory dinamicamente la variabile occorre indicare un terzo argomento all'interno della chiamate del kernel: +```c kernel<>> ```. + +#informalmente()[ + Ho dei dati di grandi dimensioni che vengono partizionati in blocchi. Se ragiono a livello di blocco devo pensare che un certo blocco di thread usa dei dati, i dati vengono trasferiti a chunk quando un certo blocco viene scelto per l'esecuzione dall'SM. +] + +#esempio()[ + Somma partiziale su blocchi + + Al livello della reduction + //Inserire esempio + +] + +=== TILING + +Cerco di rendre efficiente la tail 1024 alla volta (32*32). Devo suddividere in tail indipendenti tra di loro, essi andranno combinati tra di loro per ottenere il risultato finale. + +#esempio()[ + Prodotto matriciale. I parametei fondamentali sono quelli del risultato finale ovvero righe e colonne. + + Se ogni thread si occupa di una certa cella della matrice risultante (lavora in modo asincrono) se metto in gioco un numero di thread pari alla dimensione della matrice ottengo un thread per ogni cella risultante. + + La shared memory ha senso. Ogni entry (riga) della matrice di input viene usata per molti prodotti. + + #attenzione()[ + Leggere la matrice riga per riga è diversa da leggerla per colonna, in quanto è salavata row major order. Due righe sono contigue ma non è detto che due colonne adiacienti lo siano in global memory. + + la shared memory risolve anche questi problemi di accesso alla memori globale (come sono salvati i dati). + ] + + La SMEM Ha senso improntala 2D in questo caso. Tappezziamo le due matrici di input (le dividiamo in matrici più piccole), otteniamo dei blocchi disgiunti che coprono le due matrici. Successivamente all'interno di ogni blocco mi concentrerò su una determinata entry. + + #nota()[ + Se dividessi la matrice in tanti blocchi (massimo 1024 elementi) in casi di matrice di input molto grande creerei un numero troppo elevato di blocchi. + ] + + //aggiungere immagine + Strategia. In un blocco lavorano tutti i thread assieme. Il blocco di output conviene usi la shared memory, ma per avere il risultato finale di tutti i thread che appartengono, devo spazoolzare più blocchi delle matrici di input (in base alla lenght e alla height delle matrici di input). + + Il risultato viene calcolato quadrato per quadrato (block size) simultaneamente. Devo avere due SMEM. Per arrivare al risultato dentro il kernel mi serve caricare -> computare -> sincronizzare +] + +=== TIling Numba +//aggungere + +Convoluzione. Le memorie shard dovrebbero essere anche arricchite con l'alone. diff --git a/main.typ b/main.typ index b88ff8a..ebc780d 100644 --- a/main.typ +++ b/main.typ @@ -34,4 +34,7 @@ // #include "chapters/Lezione1-Introduzione.typ" #include "chapters/Lezione2-Gerarchia-Thread.typ" -#include "chapters/Lezione3-Warp.typ" \ No newline at end of file +#include "chapters/Lezione3-Warp.typ" + + +#include "chapters/Lezione5.typ" \ No newline at end of file From 244da9db27489da6ffac4a3227bc7dfe42b54e31 Mon Sep 17 00:00:00 2001 From: LucaCorra02 Date: Mon, 2 Feb 2026 12:12:10 +0100 Subject: [PATCH 2/3] update --- chapters/Lezione5.typ | 40 +++++++++++++++++++++++----------------- 1 file changed, 23 insertions(+), 17 deletions(-) diff --git a/chapters/Lezione5.typ b/chapters/Lezione5.typ index 81b31f5..8313969 100644 --- a/chapters/Lezione5.typ +++ b/chapters/Lezione5.typ @@ -1,37 +1,43 @@ #import "../template.typ": * -= Lezione 5 +== Cache su GPU -== Memory -Le cache sulla GPU possono essere : -- L1 (una per ogni SM) -- L2 (condivisa tra tutti gli SM) +Su una GPU esistono in genere $4$ tipi di cache: +- L1: Una per ogni SM +- L2: Condivisa tra tutti gli SM della GPU +- Read only constant +- Read only texture -Le cache a volte sono fisicamente in un area della device memory. +Le ultime due cache sono presenti su ogni SM e vengono utilizzate per migliorare le prestazioni in lettura dai rispettivi in lettura dai rispettivi spazi di memoria sul device. -``` cudaFuncSetAttribute``` è una primitiva run time o compile time che permette di dividere in modo arbitrario la L1 cache, tra L1 e shared memory. - -Il numero massimo di blocchi ospitati all'interno di un SM può scendere se all'interno di un blocchi richiedo molte risorse (registre e shared memory), in quanto esse sono condivise nell'SM. Il numero massimo di blocchi non è mai fisso ma dipende dall'occupancy (tasso di warp attivi simltuanamente nell'SM). A compile time si viene a sapere il nuemro di blocchi operativi simultaneamente. +Solitamente le cache L1 e L2 vengono usate per memorizzare dati in *memoria locale* e *globale*, incluso lo spilling dei registri (eccessi nell'uso di local memory).\ +Inoltre, tramite la primitiva ` cudaFuncSetAttribute` è possibile a run time o compile time dividere in modo arbitrario la L1 cache, tra L1 e shared memory. == Shared memory -Dobbiamo trasportare dei dati all'interno dei core, perchè ci metto di mezzo una memoria shared tra global memory (fuori dal chip) e registri? Ha senso usare la shared memory quanto il tasso di accesso del singolo dato è molto elevato ad esempio quando i thread di un certo blocco cooperano sul blocco di dati in shared memory. il beneficio è che la bandwidth della shared memory è molto altra rispetto a quella fuori dal chip. +A differenza della global memory (DRAM) situata al di fuori dei chip, la *shared memory* si tratta di una memoria che risiede sul chip, offrendo una bandwidth molto più alta e una minore latenza rispetto alla DRAM globale, giustificando il costo dell'overhead aggiunto. + +La shared memory ha una *visibilità a livello di blocco*. I thread all'interno dello stesso blocco comunicano attraverso shared memory, per questo motivo, la shared memory ha una vita correlata con la durata di esecuzione del blocco. -La shared memory viene caricata e usata intensamente dai thread di un blocco, i thread non vendono la shared memory degli altri blocchi (a meno di clustering). +#nota()[ + Ha senso usare la shared memory quanto il *tasso di accesso al singolo dato è molto elevato*, ad esempio quando i thread di un certo blocco cooperano su un blocco di dati in shared memory. +] -Barriera di sincronizzazione. Il trasferimento di dati dalla global memory alla shared memory è distrubuiti tra i thread del blocco, ogni thread trasporta un certo set di dati.\ -Tuttavia se il trasferimento è parallelo srve sincronizzazione per garantire che per iniziare la computazione tutti i dati sono stati trasferiti, abbiamo così dei dati afffidabili (no race condition). +Siccome il *trasferimento di dati* dalla global memory alla shared memory è *distribuito* tra i thread del blocco (ogni thread trasporta un certo set di dati), è *necessaria sincronizzazione*. In quanto per iniziare la computazione è necessario che tutti i dati siano stati trasferiti, abbiamo così dei dati afffidabili (no race condition). -=== Access via banking +=== Parallel access via banking -La memoria shared è organizzata in 32 bank fisici (per la dimensione del warp). +La memoria shared viene divisa in $32$ moduli della stessa ampiezza chiamati *bank*. La dimensione dei banchi dipende dalla dimensione dei dati, solitamente $4$ byte(int), ovvero una word. I banchi solitamente implementato sequenze di word. #nota()[ - La shared memory è tagliata sulla dimensione del warp. + La shared memory è tagliata sulla dimensione del warp, in modo tale da aumentare l'efficienza. ] -I 32 thread arrivano e prevelano i dati purche siano distribuiti in 32 banchi distinti. La dimensione dei banchi dipende dalla dimensione dei dati, solitamente $4$ byte(int), ovvero una word. I banchi solitamente implementato sequenze di word. +Ogni richiesta di accesso fatta di $N$ indirizzi che riguardino $N$ distinti bank sono serviti simultaneamente. +#nota()[ + Ogni SM ha una quantità limitata di shared memory che viene ripartita tra i blocchi di thread. Il numero massimo di blocchi ospitati all'interno di un SM può scendere se all'interno di un blocco vengono richieste molte risorse (registri e shared memory), in quanto esse sono condivise nell'SM. Il numero massimo di blocchi non è mai fisso ma dipende dall'*occupancy* (tasso di warp attivi simltuanamente nell'SM). +] Pattern di accesso. Possiamo avere: - Acesso parallelo, ogni thread accede ad un blocco differente From 82ed827abed532fe378395d670859ebf72de515c Mon Sep 17 00:00:00 2001 From: LucaCorra02 Date: Mon, 2 Feb 2026 13:15:12 +0100 Subject: [PATCH 3/3] update acess --- chapters/Lezione5.typ | 235 ++++++++++++++++++++++++++++++++++++------ 1 file changed, 205 insertions(+), 30 deletions(-) diff --git a/chapters/Lezione5.typ b/chapters/Lezione5.typ index 8313969..37627e1 100644 --- a/chapters/Lezione5.typ +++ b/chapters/Lezione5.typ @@ -2,22 +2,22 @@ == Cache su GPU -Su una GPU esistono in genere $4$ tipi di cache: +Su una GPU esistono in genere $4$ tipi di cache: - L1: Una per ogni SM - L2: Condivisa tra tutti gli SM della GPU - Read only constant -- Read only texture +- Read only texture -Le ultime due cache sono presenti su ogni SM e vengono utilizzate per migliorare le prestazioni in lettura dai rispettivi in lettura dai rispettivi spazi di memoria sul device. +Le ultime due cache sono presenti su ogni SM e vengono utilizzate per migliorare le prestazioni in lettura dai rispettivi in lettura dai rispettivi spazi di memoria sul device. Solitamente le cache L1 e L2 vengono usate per memorizzare dati in *memoria locale* e *globale*, incluso lo spilling dei registri (eccessi nell'uso di local memory).\ Inoltre, tramite la primitiva ` cudaFuncSetAttribute` è possibile a run time o compile time dividere in modo arbitrario la L1 cache, tra L1 e shared memory. -== Shared memory +== Shared memory -A differenza della global memory (DRAM) situata al di fuori dei chip, la *shared memory* si tratta di una memoria che risiede sul chip, offrendo una bandwidth molto più alta e una minore latenza rispetto alla DRAM globale, giustificando il costo dell'overhead aggiunto. +A differenza della global memory (DRAM) situata al di fuori dei chip, la *shared memory* si tratta di una memoria che risiede sul chip, offrendo una bandwidth molto più alta e una minore latenza rispetto alla DRAM globale, giustificando il costo dell'overhead aggiunto. -La shared memory ha una *visibilità a livello di blocco*. I thread all'interno dello stesso blocco comunicano attraverso shared memory, per questo motivo, la shared memory ha una vita correlata con la durata di esecuzione del blocco. +La shared memory ha una *visibilità a livello di blocco*. I thread all'interno dello stesso blocco comunicano attraverso shared memory, per questo motivo, la shared memory ha una vita correlata con la durata di esecuzione del blocco. #nota()[ Ha senso usare la shared memory quanto il *tasso di accesso al singolo dato è molto elevato*, ad esempio quando i thread di un certo blocco cooperano su un blocco di dati in shared memory. @@ -25,6 +25,13 @@ La shared memory ha una *visibilità a livello di blocco*. I thread all'interno Siccome il *trasferimento di dati* dalla global memory alla shared memory è *distribuito* tra i thread del blocco (ogni thread trasporta un certo set di dati), è *necessaria sincronizzazione*. In quanto per iniziare la computazione è necessario che tutti i dati siano stati trasferiti, abbiamo così dei dati afffidabili (no race condition). +Maggiore è la shared memory richiesta da un kernel, minore è il numero di blocchi attivi concorrenti. + +#nota()[ + Ogni SM ha una quantità limitata di shared memory che viene ripartita tra i blocchi di thread. Il numero massimo di blocchi ospitati all'interno di un SM può scendere se all'interno di un blocco vengono richieste molte risorse (registri e shared memory), in quanto esse sono condivise nell'SM.\ + Il numero massimo di blocchi non è mai fisso ma dipende dall'*occupancy* (tasso di warp attivi simltuanamente nell'SM). +] + === Parallel access via banking La memoria shared viene divisa in $32$ moduli della stessa ampiezza chiamati *bank*. La dimensione dei banchi dipende dalla dimensione dei dati, solitamente $4$ byte(int), ovvero una word. I banchi solitamente implementato sequenze di word. @@ -33,25 +40,193 @@ La memoria shared viene divisa in $32$ moduli della stessa ampiezza chiamati *ba La shared memory è tagliata sulla dimensione del warp, in modo tale da aumentare l'efficienza. ] -Ogni richiesta di accesso fatta di $N$ indirizzi che riguardino $N$ distinti bank sono serviti simultaneamente. +L'accesso avviene per warp. Ogni richiesta di accesso fatta di $N$ indirizzi che riguardino $N$ distinti bank sono serviti simultaneamente. + +Possiamo avere due pattern di accesso alla shared memory: +- *Acesso parallelo*: se abbiamo un warp dove ogni thread richiede al più un accesso per bank, possiamo effettuare in una sola transizione il trasferimento dati. + +- *Accesso sequenziale*: si possono creare dei conflitti se diversi thread cercano di accedere allo stesso blocco. In questo caso l'hardware effetua tante transizioni quante ne sono necessarie per eliminare i conflitti + +#figure( + stack( + dir: ttb, + spacing: 2em, + [ + // Accesso parallelo - nessun conflitto + #align(center)[ + #text(size: 10pt, weight: "bold", fill: rgb(100, 180, 50))[Accesso Parallelo (senza conflitti)] + ] + #import cetz.draw: * + #cetz.canvas({ + let thread_w = 0.35 + let bank_w = 0.35 + let thread_h = 0.4 + let bank_h = 0.4 + + // Thread row + for i in range(32) { + let x = i * thread_w + rect((x, 2.5), (x + thread_w, 2.5 + thread_h), + fill: rgb(200, 220, 200), + stroke: black) + content((x + thread_w/2, 2.7), + text(size: 5pt, weight: "bold")[T#i]) + } + + // Arrows - ogni thread accede a un bank diverso + for i in range(32) { + let x = i * thread_w + thread_w/2 + line((x, 2.5), (x, 1.5), + mark: (end: "stealth"), + stroke: (paint: rgb(100, 180, 50), thickness: 0.8pt)) + } + + // Bank row + for i in range(32) { + let x = i * bank_w + rect((x, 1), (x + bank_w, 1 + bank_h), + fill: rgb(150, 200, 255), + stroke: black) + content((x + bank_w/2, 1.2), + text(size: 5pt, weight: "bold")[B#i]) + } + + // Label + content((5.6, 0.5), text(size: 8pt, fill: rgb(100, 180, 50))[ + *1 transazione* - Tutti i thread accedono simultaneamente + ]) + }) + ], + [ + // Accesso sequenziale - conflitti + #align(center)[ + #text(size: 10pt, weight: "bold", fill: rgb(200, 50, 50))[Accesso Sequenziale (con conflitti)] + ] + #import cetz.draw: * + #cetz.canvas({ + let thread_w = 0.35 + let bank_w = 0.35 + let thread_h = 0.4 + let bank_h = 0.4 + + // Thread row + for i in range(32) { + let x = i * thread_w + rect((x, 2.5), (x + thread_w, 2.5 + thread_h), + fill: rgb(220, 200, 200), + stroke: black) + content((x + thread_w/2, 2.7), + text(size: 5pt, weight: "bold")[T#i]) + } + + // Arrows - più thread accedono allo stesso bank (esempio: bank conflict) + // Threads 0-3 tutti su Bank 0 (rosso) + for i in range(4) { + let x_thread = i * thread_w + thread_w/2 + let x_bank = 0 + bank_w/2 + line((x_thread, 2.5), (x_bank, 1.4), + mark: (end: "stealth"), + stroke: (paint: rgb(200, 50, 50), thickness: 0.8pt)) + } + + // Threads 4-7 tutti su Bank 4 (rosso) + for i in range(4, 8) { + let x_thread = i * thread_w + thread_w/2 + let x_bank = 4 * bank_w + bank_w/2 + line((x_thread, 2.5), (x_bank, 1.4), + mark: (end: "stealth"), + stroke: (paint: rgb(200, 50, 50), thickness: 0.8pt)) + } + + // Threads 8-11 su Bank 8 (arancione) + for i in range(8, 12) { + let x_thread = i * thread_w + thread_w/2 + let x_bank = 8 * bank_w + bank_w/2 + line((x_thread, 2.5), (x_bank, 1.4), + mark: (end: "stealth"), + stroke: (paint: rgb(255, 150, 50), thickness: 0.8pt)) + } + + // Altri thread - pattern simile + for i in range(12, 16) { + let x_thread = i * thread_w + thread_w/2 + let x_bank = 12 * bank_w + bank_w/2 + line((x_thread, 2.5), (x_bank, 1.4), + mark: (end: "stealth"), + stroke: (paint: rgb(255, 150, 50), thickness: 0.8pt)) + } + + // Bank row + for i in range(32) { + let x = i * bank_w + let color = if i == 0 or i == 4 { + rgb(255, 150, 150) + } else if i == 8 or i == 12 { + rgb(255, 200, 150) + } else { + rgb(150, 200, 255) + } + + rect((x, 1), (x + bank_w, 1 + bank_h), + fill: color, + stroke: black) + content((x + bank_w/2, 1.2), + text(size: 5pt, weight: "bold")[B#i]) + } + + // Label + content((5.6, 0.5), text(size: 8pt, fill: rgb(200, 50, 50))[ + *4 transazioni* - Conflitti serializzano gli accessi + ]) + }) + ] + ), + caption: [ + Differenza tra accesso parallelo e sequenziale alla shared memory. + ] +) + + #nota()[ - Ogni SM ha una quantità limitata di shared memory che viene ripartita tra i blocchi di thread. Il numero massimo di blocchi ospitati all'interno di un SM può scendere se all'interno di un blocco vengono richieste molte risorse (registri e shared memory), in quanto esse sono condivise nell'SM. Il numero massimo di blocchi non è mai fisso ma dipende dall'*occupancy* (tasso di warp attivi simltuanamente nell'SM). + La bandwidth effettiva si riduce di un fattore pari al numero di transizioni separate necessarie. ] -Pattern di accesso. Possiamo avere: -- Acesso parallelo, ogni thread accede ad un blocco differente -- Accesso sequenziale, si possono creare dei conflitti se diversi thred cercano di accedere allo stesso blocco. +Inoltre possono essere effetuati i seguenti pattern di accesso: +- Caso *broadcat*: un singolo valore letto da tutti i thread in un warp da un singolo bank (tutti i thread vogliono un certo dato) +``` + float f = smem[20]; +``` -=== SMEM +- Caso *parallelo*: un singolo valore letto da un singolo bank. Tutti i thread leggono un dato in posizione diversa +``` + float f = smem[threadIdx.x] +``` + +- Caso *conflitto doppio*: acceso agli elementi saltandone uno ogni volta (passo 2). Pattern per gestire strutture di $8$ byte (come `double` o coppie di `float` complesse) +``` + float f = smem[(threadIdx.x * 2)%32] +``` + La GPU deve serializzare l'accesso: prima serve il primo gruppo, poi il secondo. Questo dimezza la velocità effettiva della memoria. + +- *Nessun conflitto*: accesso agli elementi con un passo di $3$, ad esempio per leggere le coordinate $X, Y, Z$ di un punto: +``` + float f = smem[(threadIdx.x * 3)%32] +``` + +#nota()[ + In questo caso *non ci sono conflitti*, in quanto matematicamente, il numero $3$ e il numero $32$ sono coprimi (non hanno divisori comuni). Quando il passo (stride) è dispari e non multiplo di $32$, i thread si distribuiscono perfettamente sui banchi senza conflitti. Tutti i $32$ thread accedono in parallelo, massima velocità. +] + +=== SMEM Se la dimensione non è nota a compile time è possibile dichiarare una variabile adimensionale con la keyword ``` extern```: -- dinamicamente si possono allocare solo array 1D +- dinamicamente si possono allocare solo array 1D - può essere sia interna al kernl sia esterna ``` extern __shared__ int array[]```. Alloca un array nelal shared memory. -Per allocare la shared memory dinamicamente la variabile occorre indicare un terzo argomento all'interno della chiamate del kernel: -```c kernel<>> ```. +Per allocare la shared memory dinamicamente la variabile occorre indicare un terzo argomento all'interno della chiamate del kernel: +```c kernel<>> ```. #informalmente()[ Ho dei dati di grandi dimensioni che vengono partizionati in blocchi. Se ragiono a livello di blocco devo pensare che un certo blocco di thread usa dei dati, i dati vengono trasferiti a chunk quando un certo blocco viene scelto per l'esecuzione dall'SM. @@ -60,41 +235,41 @@ Per allocare la shared memory dinamicamente la variabile occorre indicare un ter #esempio()[ Somma partiziale su blocchi - Al livello della reduction + Al livello della reduction //Inserire esempio ] -=== TILING +=== TILING -Cerco di rendre efficiente la tail 1024 alla volta (32*32). Devo suddividere in tail indipendenti tra di loro, essi andranno combinati tra di loro per ottenere il risultato finale. +Cerco di rendre efficiente la tail 1024 alla volta (32*32). Devo suddividere in tail indipendenti tra di loro, essi andranno combinati tra di loro per ottenere il risultato finale. #esempio()[ - Prodotto matriciale. I parametei fondamentali sono quelli del risultato finale ovvero righe e colonne. + Prodotto matriciale. I parametei fondamentali sono quelli del risultato finale ovvero righe e colonne. - Se ogni thread si occupa di una certa cella della matrice risultante (lavora in modo asincrono) se metto in gioco un numero di thread pari alla dimensione della matrice ottengo un thread per ogni cella risultante. + Se ogni thread si occupa di una certa cella della matrice risultante (lavora in modo asincrono) se metto in gioco un numero di thread pari alla dimensione della matrice ottengo un thread per ogni cella risultante. - La shared memory ha senso. Ogni entry (riga) della matrice di input viene usata per molti prodotti. + La shared memory ha senso. Ogni entry (riga) della matrice di input viene usata per molti prodotti. #attenzione()[ - Leggere la matrice riga per riga è diversa da leggerla per colonna, in quanto è salavata row major order. Due righe sono contigue ma non è detto che due colonne adiacienti lo siano in global memory. + Leggere la matrice riga per riga è diversa da leggerla per colonna, in quanto è salavata row major order. Due righe sono contigue ma non è detto che due colonne adiacienti lo siano in global memory. la shared memory risolve anche questi problemi di accesso alla memori globale (come sono salvati i dati). ] - La SMEM Ha senso improntala 2D in questo caso. Tappezziamo le due matrici di input (le dividiamo in matrici più piccole), otteniamo dei blocchi disgiunti che coprono le due matrici. Successivamente all'interno di ogni blocco mi concentrerò su una determinata entry. + La SMEM Ha senso improntala 2D in questo caso. Tappezziamo le due matrici di input (le dividiamo in matrici più piccole), otteniamo dei blocchi disgiunti che coprono le due matrici. Successivamente all'interno di ogni blocco mi concentrerò su una determinata entry. - #nota()[ - Se dividessi la matrice in tanti blocchi (massimo 1024 elementi) in casi di matrice di input molto grande creerei un numero troppo elevato di blocchi. - ] + #nota()[ + Se dividessi la matrice in tanti blocchi (massimo 1024 elementi) in casi di matrice di input molto grande creerei un numero troppo elevato di blocchi. + ] - //aggiungere immagine - Strategia. In un blocco lavorano tutti i thread assieme. Il blocco di output conviene usi la shared memory, ma per avere il risultato finale di tutti i thread che appartengono, devo spazoolzare più blocchi delle matrici di input (in base alla lenght e alla height delle matrici di input). + //aggiungere immagine + Strategia. In un blocco lavorano tutti i thread assieme. Il blocco di output conviene usi la shared memory, ma per avere il risultato finale di tutti i thread che appartengono, devo spazoolzare più blocchi delle matrici di input (in base alla lenght e alla height delle matrici di input). - Il risultato viene calcolato quadrato per quadrato (block size) simultaneamente. Devo avere due SMEM. Per arrivare al risultato dentro il kernel mi serve caricare -> computare -> sincronizzare + Il risultato viene calcolato quadrato per quadrato (block size) simultaneamente. Devo avere due SMEM. Per arrivare al risultato dentro il kernel mi serve caricare -> computare -> sincronizzare ] === TIling Numba //aggungere -Convoluzione. Le memorie shard dovrebbero essere anche arricchite con l'alone. +Convoluzione. Le memorie shard dovrebbero essere anche arricchite con l'alone.