Iterare sulle fasi dei tile

Moltiplicare un’intera riga per un’intera colonna è troppo per un solo tile. Si suddivide quindi la somma lunga in blocchi di larghezza TILE, uno per ogni fase.

Se le matrici hanno larghezza N e i tile hanno larghezza TILE, servono N / TILE fasi per coprire l’intera dimensione interna.

Racchiuda i passaggi di caricamento, sincronizzazione e calcolo in un ciclo su phase. A ogni iterazione, la finestra del tile avanza lungo la riga di A e la colonna di B.

La variabile locale sum si trova al di fuori del ciclo, quindi continua a crescere. Ogni fase aggiunge la propria porzione del prodotto scalare al totale progressivo.

Ogni fase sposta di phase * TILE la colonna letta da A e la riga letta da B. È così che avanza la finestra.

Come prima, chiami __syncthreads() dopo i caricamenti, in modo che il tile sia completo prima che qualcuno lo utilizzi per il calcolo.

Il ciclo interno aggiunge TILE prodotti a sum, usando solo il tile appena caricato. Contribuisce a una porzione del prodotto scalare finale.

Concluda ogni fase con un altro __syncthreads(), così nessun thread sovrascrive il tile mentre un thread più lento lo sta ancora leggendo. 🚧

Una barriera protegge le letture successive al caricamento, l’altra protegge i caricamenti successivi alle letture. Insieme mantengono ogni thread sincronizzato durante le fasi.

Al termine di tutte le fasi, la sum progressiva è il prodotto scalare completo. La memorizzi in C una sola volta, dopo aver verificato i limiti.

Quando N non è un multiplo esatto di TILE, carichi zero per gli elementi fuori intervallo, così i prodotti aggiuntivi non modificano la somma.

Rifletta su dove risiede il totale progressivo durante il ciclo sulle fasi.

Ha eseguito un ciclo sulle fasi, spostando i tile, sincronizzandosi due volte e accumulando le somme parziali in un unico totale. Ora misurerà quanto è più veloce. 📈