Review parallel reduction

Pengurangan Paralel adalah salah satu jenis Paralel Mesin Random Access (PRAM) algoritma. A Parallel Random Access Machine (PRAM) adalah memori bersama mesin abstrak yang digunakan oleh algoritma paralel desainer untuk memperkirakan kawin algoritma kinerja (seperti kompleksitas waktu). jadi, penurunan paralel adalah proses PRAM di memanipulasi data yang disimpan dalam memori global register. Ada beberapa operasi secara parallel pengurangan yaitu, operasi penambahan, pengurangan operasi, dan operasi perkalian. Parallel reduksi dapat digambarkan dengan pohon biner, yang kelompok n log nilai p ditambahkan dalam langkah-langkah tambahan secara paralel.

Kahan Summation Algorithm

Dalam analisis numerik, algoritma penjumlahan Kahan (lihat algoritma 1) juga dikenal sebagai kompensasi penjumlahan secara signifikan mengurangi kesalahan numeric dalam tarif total yang diperoleh dengan menambahkan urutan presisi floating point terbatas, dibandingkan dengan pendekatan yang jelas. Hal ini dilakukan dengan menjaga kompensasi berjalan terpisah (variabel untuk mengakumulasi kesalahan kecil). Secara khusus, hanya menjumlahkan n bilangan secara berurutan memiliki kesalahan terburuk yang tumbuh proporsional untuk n, dan root mean square error yang tumbuh sebagai √n untuk input acak (kesalahan pembulatan membentuk random walk). Dengan penjumlahan kompensasi, kesalahan terburuk terikat adalah independen dari n, sehingga sejumlah besar nilai-nilai dapat disimpulkan dengan Kesalahan yang hanya tergantung pada presisi floating-point.

Interleaved Addressing with Divergent Branching by Ian Buck

Percabangan oleh Ian BuckProgram ini, untuk selanjutnya disebut sebagai kernel, melakukan pengurangan dimana setiap thread melakukan perhitungan dua elemen yang diambil dari memori bersama. Sekarang sebuah implementasi dari jumlah EREW pengurangan PRAM proses, dimana metode pengindeksan menggunakan modulo operasi untuk setiap proses iterasi.

Metode ini sederhana namun memiliki kelemahan, yaitu benang pengindeksan tidak berguna, karena ada satu hop antara benang dan lain-lain, itu akan membuat sangat warps berbeda yang mempengaruhi proses yang tidak efisien.

Solusinya adalah sesuai dengan kinerja strategi, program ini dapat direvisi dengan aritmatika instruksi dan kontrol pendekatan instruksi aliran untuk mengganti if (proses tid% untuk meminimalkan jumlah cabang divergen dan warps. Hal ini dapat diganti dengan jika (indeks <blockDim.x) aliran kontrol, di mana intindex = 2 * s * tid.

Reduction kernel with bank conflict method

Metode ini menggunakan larutan Interleaved menangani dengan divergen percabangan. Cukup ganti cara di thread pengindeksan dengan indeks langkahnya dan non cabang berbeda. Jadi, pengindeksan benang ID tidak berdasarkan lagi dengan urutan nilai elemen dalam memori bersama.

Dengan menggunakan metode ini, indeks benang telah diurutkan dengan benar tanpa hop antara satu sama lain, sehingga tidak ada gunanya penggunaan benang.

Namun masalah baru dibagi konflik bank memori. Untuk mencapai bandwidth memori yang tinggi, bersama memori dibagi menjadi modul memori berukuran sama, disebut bank, yang dapat diakses secara bersamaan. Jadi, memori dengan permintaan membaca atau menulis terbuat dari alamat n yang jatuh n memori yang berbeda bank dapat dilayani secara bersamaan, menghasilkan efektif bandwidth yang n kali setinggi bandwidth satu modul. Namun, jika dua alamat dari memori permintaan penurunan memori yang sama Bank, ada konflik perbankan dan akses harus serial. Hardware perpecahan memori permintaan dengan bank konflik menjadi sebanyak terpisah permintaan bebas konflik yang diperlukan, mengurangi bandwidth yang efektif dengan faktor sama dengan nomor permintaan memori yang terpisah. Jika jumlah permintaan memori terpisah n, memori awal permintaan dikatakan menyebabkan bank konflik n-jalan. untuk mendapatkan kinerja maksimum, karena itu penting untuk memahami bagaimana memori alamat peta ke memori bank dalam rangka untuk menjadwalkan permintaan memori sehingga untuk meminimalkan konflik perbankan.

Untuk semua thread dari warp (32 benang), mengakses memori bersama secepat mengakses register asalkan tidak ada konflik antara Bank benang. Jadi, solusinya adalah menggunakan berurutan menangani Metode untuk mengindeks benang ID.

Reduction kernel with sequential addressing

Metode pengalamatan Sequential akan menghindari konflik Bank masalah, karena dalam metode ini thread ID pengindeksan lakukan dengan metode sequential, sehingga setiap thread tidak akan mengakses alamat yang sama pada memori Bank, hanya menunjuk satu memori bersama sama alamat.

Dalam kode sumber, berpikir bahwa harus mengganti hanya mengindeks langkah dalam lingkaran batin dengan terbalik lingkaran dan benang ID pengindeksan berbasis. Tapi, cara ini masih memiliki masalah, setengah dari benang yang menganggur pada loop pertama iterasi, sehingga membuat boros dan tentu saja akan berpengaruh terhadap kinerja. Solusinya bisa mendapatkan dengan melakukan pengurangan selama proses elemen beban dari memori global untuk memori bersama.

Reduction kernel with first add during load

Untuk menghindari idle benang pada berurutan menangani Metode dengan melakukan pengurangan ketika elemen beban dari memori global untuk memori bersama atau dapat dikatakan pertama tambahkan selama metode beban.

Dalam metode ini, selama proses loop elemen pemuatan dari memori global, juga dilakukan pengurangan proses dari elemen yang telah selesai untuk memuat sebelum akhirnya menulis di memori bersama. jadi, unsur proses dalam fungsi kernel adalah hasil perhitungan jumlah juga. Intinya adalah mengganti beban tunggal dengan dua beban pada proses load dan tampil pertama kali menambahkan selama beban diproses.

Ini harus menjadi metode terbaik yang dapat digunakan untuk mencapai kinerja puncak GPU, namun pada kenyataannya jangkauan bandwidth memori masih jauh dari bandwidth yang terikat, mungkin disebabkan oleh hambatan atau instruksi instruksi latency. Latency instruksi ini adalah jumlah siklus clock yang diperlukan untuk instruksi untuk melewati pipa. Untuk menghindari kondisi ini, strategi yang dapat digunakan adalah membuka gulungan loop, karena pengurangan memiliki intensitas aritmatika rendah

Reduction kernel with unroll last warp

Kadang-kadang, compiler dapat membuka gulungan loop atau mengoptimalkan apakah atau beralih pernyataan dengan menggunakan cabang predikasi instead.In kasus ini, tidak pernah bisa warp menyimpang. Pemrogram juga dapat mengontrol lingkaran membuka gulungan.

Sebagai hasil pengurangan, jumlah aktif benang menurun. Dalam GPU arsitektur, kernel atau instruksi yang SIMD sinkron dalam warp. Warp adalah sekelompok benang dieksekusi fisik secara paralel, dan masing-masing kelompok terdiri dari warp 32 benang. Itu berarti ketika "s ¡= 32", tidak lagi perlu _syncthreads_ () dan tidak perlu lagi "If (tid ¡s)" karena tidak menyimpan pekerjaan apapun, sehingga bisa kontrol loop membuka gulungan. Mari kita membuka gulungan 6 terakhir iterasi loop batin.

Ini menghemat pekerjaan berguna dalam semua warps, bukan hanya yang terakhir. Tanpa membuka gulungan, semua warps mengeksekusi setiapiterasi untuk loop dan jika pernyataan.

Reduction kernel with completely unrolled

Jika jumlah iterasi pada waktu kompilasi adalah dikenal, itu benar-benar bisa membuka gulungan pengurangan. Untungnya, ukuran blok dibatasi oleh GPU untuk 512 benang dan juga menempel kekuatan 2 ukuran blok. Sehingga dapat dengan  udah membuka gulungan untuk ukuran blok tetap, kemudian mengganti proses berulang "untuk" dengan sepenuhnya membuka gulungan.

Reduction kernel with multiple elements per thread

Mengingat jumlah benang per grid, jumlah benang per blok, atau ekuivalen nomor blok, harus dipilih untuk memaksimalkan pemanfaatan sumber daya komputasi yang tersedia. Ini berarti bahwa harus ada setidaknya banyak blok karena ada multiprocessors dalam perangkat.

Selanjutnya, berjalan hanya satu blok per multiprosesor akan memaksa multiprosesor untuk idle selama sinkronisasi benang dan juga selama perangkat membaca memori jika ada tidak cukup benang per blok untuk menutupi beban latency. Oleh karena itu biasanya baik untuk memungkinkan dua atau lebih blok menjadi aktif pada setiap multiprosesor untuk memungkinkan tumpang tindih antara blok yang menunggu dan blok yang dapat berjalan. Agar hal ini terjadi, tidak hanya harus ada di Setidaknya dua kali lebih banyak blok karena ada multiprocessors pada perangkat, tetapi juga jumlah yang dialokasikan memori bersama per blok harus paling setengah jumlah total memori bersama yang tersedia per multiprosesor

Hasil

·         Hasil untuk 2²⁴ elemen array

Hasil menjalankan semua metode, dibandingkan dengan memori bandwidth dan memakan waktu ditunjukkan oleh Gambar 1. Itu saja dijelaskan kinerja oleh masing-masing metode untuk memproses 2²⁴ elemen array, di mana kolom "Time" adalah waktu proses yang dibutuhkan dalam milidetik, "Bandwidth" kolom tingkat di mana data dapat membaca dari atau disimpan ke dalam memori. Ingatan bandwidth biasanya dinyatakan dalam satuan byte / detik, kolom "Langkah Speedup" adalah nilai waktu kernel saat ini dibagi dengan kernel sebelumnya waktu, dan kolom "SpeedUp Kumulatif" menunjukkan nilai waktu kernel saat ini dibagi oleh kernel 1 (metode cabang divergen) waktu.

Gambar 1. Result performance by all of method

·         Hasil untuk 2¹⁷-2²⁵ elemen array

Gambar 2, menunjukkan perbandingan kinerja dengan waktu pada grafis. Pada angka itu bisa melihat bahwa dari 131,072 - 1.048.576 (2¹⁷-2²⁰) elemen array, semua metode yang masih dapat diandalkan untuk digunakan. Pada 2.097.152 - 16.777.216 (2²¹-2²⁴) elemen array, metode Divergen Percabangan, Konflik Bank, dan Sequential Mengatasi tidak dapat diandalkan untuk digunakan, karena mereka membutuhkan waktu lebih dari 1,0 detik. Selanjutnya, pada 2²⁵ menunjukkan bahwa hanya empat metode yang masih dapat diandalkan untuk digunakan. Pada akhir hasil, metode terbaik adalah Beberapa Elemen per thread karena mengkonsumsi waktu lebih kurang dari yang lain.

Gambar 2. Result Comparison by time

Gambar 3, menunjukkan perbandingan kinerja dengan bandwidth per detik pada grafis, kemampuannya bahwa metode yang terbaik adalah yang bisa mencapai tertinggi bandwidth atau lebih dekat dengan bandwidth GPU batas, pada 54,4 GB / s. Pada angka itu menunjukkan kinerja dari 131,072 - 33.554.432 (2¹⁷ to 2²⁵) elemen.

Gambar 3. Result Comparison by bandwidth

            Kesimpulan

Berdasarkan hasil penelitian, kinerja puncak GPU telah mencapai. Ini dibuktikan dengan memori akuisisi Nilai bandwith (pada 43,3 GB / s) dengan memori bandwidth yang terikat (54,4 GB / s) pada perangkat GPU (Geforce GT 240).

Untuk mengoptimalkan GPU untuk proses reduksi parallel oleh menjaga multiprocessors pada perangkat sebagai sesibuk mungkin. Perangkat yang bekerja adalah buruk seimbang di seluruh multiprocessors akan memberikan kinerja optimal. Oleh karena itu, penting untuk merancang aplikasi Anda untuk menggunakan benang dan blok dengan cara yang memaksimalkan pemanfaatan perangkat keras dan untuk membatasi praktek-praktek yang menghambat distribusi bebas kerja. Salah satu kunci untuk kinerja yang baik adalah Konsep lain yang penting adalah manajemen sumber daya sistem yang dialokasikan untuk suatu tugas tertentu.

Strategi optimasi CUDA dapat digunakan untuk memaksimalkan pemanfaatan hardware. Dalam hal ini pengurangan proses, dapat menggunakan divergen percabangan, bank konflik, memori penggabungan dengan berurutan pengalamatan dan latency bersembunyi dengan membuka gulungan lingkaran.

Sebenarnya, begitu banyak jenis pendekatan yang dapat digunakan untuk  mengoptimalkan kinerja pada GPU, namun utama untuk CUDA, ada dua strategi jenis, yaitu instruksi Throughput dan bandwidth memori.