Pembacaan Terkonsolidasi versus Berstride

Penggabungan akses bergantung pada alamat yang disentuh setiap thread. Pemetaan dari thread ke alamat menentukan apakah satu transaksi dapat melayani seluruh warp.

Ketika thread yang berdekatan membaca alamat yang berdekatan, warp mencakup satu rangkaian yang berurutan. Tata letak yang rapi ini merupakan akses terkoalesensi.

Thread 0 mengakses indeks 0, thread 1 mengakses indeks 1, dan seterusnya. Semua 32 alamat berada dalam satu baris yang selaras, sehingga GPU hanya memerlukan satu transaksi.

Langkah adalah jarak tetap antara alamat yang disentuh oleh thread-thread berturut-turut. Begitu jarak tersebut lebih besar dari satu elemen, penggabungan akses mulai terganggu.

Dengan langkah 2, thread 0 membaca 0, thread 1 membaca 2, dan thread 2 membaca 4. Kini warp mencakup memori dua kali lebih luas, sehingga memerlukan lebih banyak transaksi.

Gandakan langkahnya dan Anda kurang lebih menggandakan jumlah baris yang disentuh. Langkah yang besar dapat membuat warp memerlukan banyak transaksi, sementara hanya menggunakan sebagian kecil dari setiap baris.

Memberikan satu kolom penuh kepada setiap thread dalam matriks berurutan berdasarkan baris menciptakan langkah yang sangat besar. Kodenya tampak rapi, tetapi diam-diam menyebarkan setiap warp.

Sering kali solusinya cukup dengan menukar indeks mana yang berubah paling cepat terhadap thread. Biarkan thread yang berurutan menelusuri memori yang berurutan, maka pembacaan kembali tergabung.

Bahkan pola yang terkoalesensi dapat terganggu jika dimulai di tengah baris. Offset yang tidak selaras menggeser warp melewati batas, sehingga satu pembacaan terpecah menjadi dua.

Untuk menilai suatu pola, jangan membayangkan satu thread saja. Bayangkan semua 32 lajur sekaligus, lalu tanyakan berapa banyak baris yang dicakup oleh alamat-alamatnya secara bersama-sama. 🔍

Buat indeks yang berubah paling cepat mengikuti threadIdx.x. Kebiasaan sederhana ini membuat sebagian besar pembacaan global Anda terkoalesensi tanpa usaha tambahan.

Anda telah melihat bahwa pembacaan terkoalesensi menjaga elemen-elemen yang berdekatan tetap bersama, sedangkan langkah menyebarkannya ke berbagai baris. Biarkan threadIdx.x mengendalikan indeks yang berubah paling cepat. 🎉