Elakkan Konflik Bank

Memori dikongsi dibahagikan kepada 32 bank, satu bagi setiap utas dalam kumpulan warp. Menyebarkan akses merentasi bank tersebut membolehkan kesemua 32 utas membaca pada kelajuan penuh.

Perkataan 4 bait berturutan ditempatkan dalam bank berturutan, berulang semula selepas 32. Jadi perkataan 0 berada dalam bank 0 dan perkataan 32 berada dalam bank 0 sekali lagi.

Konflik bank berlaku apabila dua utas dalam kumpulan warp mengakses bank yang sama tetapi perkataan yang berbeza. Perkakasan mesti melaksanakan akses tersebut secara bersiri.

Konflik N hala menukarkan satu akses pantas kepada N akses perlahan. Konflik 32 hala boleh menjadikan memori dikongsi sama perlahan seperti akses satu demi satu. 🐢

Apabila setiap utas membaca tile[threadIdx.x], setiap utas mengakses bank yang berbeza. Tiada konflik berlaku dan ia berjalan pada jalur lebar penuh.

Pembacaan dengan langkah 32 menghantar setiap utas ke bank yang sama. Itulah keadaan terburuk, iaitu konflik 32 hala penuh.

Sebarang langkah genap yang berkongsi faktor dengan 32 menyebabkan konflik separa. Langkah sebanyak 2, contohnya, menghasilkan konflik 2-hala merentasi warp.

Berita baik: jika semua thread membaca alamat yang sama, perkakasan menyiarkannya sekali gus. Perkataan yang sama tidak menjadi masalah; hanya bank yang sama dengan perkataan berbeza menyebabkan konflik.

Untuk jubin 2D, tambahkan satu lajur tambahan dengan padding +1. Ini mengalihkan setiap baris supaya capaian lajur tidak lagi semuanya berada dalam satu bank.

Lajur tambahan menjadikan panjang baris relatif perdana dengan 32. Kini, bergerak ke bawah lajur akan mengunjungi bank yang berbeza setiap kali, lalu menghapuskan konflik.

Konflik bank tidak kelihatan dalam kod sumber. Gunakan Nsight Compute untuk mengukur konflik memori dikongsi sebelum meluangkan usaha untuk membaikinya.

Anda telah mempelajari bahawa memori dikongsi mempunyai 32 bank, bahawa capaian perkataan berbeza dalam bank yang sama akan disiri, dan bahawa padding +1 membaiki konflik lajur. Seterusnya: memori dikongsi dinamik. 🎯