IF Notes

Histogram dan Operasi Atomik

6 min read

Back to IF3130 Sistem Paralel dan Terdistribusi

Pola Algoritma Paralel: Histogram dan Solusi Konflik dengan Operasi Atomik

Questions/Cues

  • Apa itu komputasi histogram?

  • Bagaimana partisi input yang efisien?

  • Apa itu Data Race?

  • Mengapa Read-Modify-Write berbahaya?

  • Apa solusi untuk Data Race?

  • Bagaimana cara kerja Operasi Atomik?

  • Apa kelemahan performa atomics?

  • Apa itu teknik Privatization?

  • Bagaimana cara kerja Histogram Terprivatisasi?

Reference Points

  • 7 - IF-3230-07-GPU-05-2023.pdf

Apa itu Komputasi Histogram?

Histogram adalah sebuah metode untuk menghitung frekuensi kemunculan setiap nilai dalam sebuah kumpulan data. Prosesnya melibatkan pendefinisian beberapa “keranjang” (bin), dan untuk setiap elemen data input, nilai keranjang yang sesuai akan diinkremen (ditambah satu).

  • Aplikasi: Sangat umum digunakan untuk ekstraksi fitur dan analisis data, seperti menghitung distribusi warna dalam gambar, mendeteksi anomali dalam transaksi keuangan, atau menganalisis frekuensi kata dalam teks.

Partisi Input yang Efisien

Saat memparalelkan pembuatan histogram, cara thread membaca data input sangat mempengaruhi efisiensi akses memori.

  1. Sectioned Partitioning: Setiap thread memproses satu blok data yang bersebelahan. Ini buruk karena thread yang bersebelahan akan mengakses lokasi memori yang jauh, menghancurkan memory coalescing.

  2. Interleaved Partitioning: Thread yang bersebelahan memproses elemen input yang bersebelahan, lalu melompat bersama ke bagian berikutnya. Ini baik karena memastikan semua akses ke Global Memory bersifat coalesced, memaksimalkan penggunaan bandwidth.

Masalah Inti: Data Race

Masalah fundamental dalam histogram paralel muncul ketika beberapa thread mencoba memperbarui bin yang sama pada saat yang bersamaan.

  • Definisi: Data Race terjadi ketika beberapa thread mengakses lokasi memori yang sama secara bersamaan, dan setidaknya salah satu dari akses tersebut adalah operasi tulis. Hasil akhir dari operasi menjadi tidak dapat diprediksi dan bergantung pada urutan eksekusi thread yang kebetulan.

Bahaya Operasi Read-Modify-Write

Inkrementasi sebuah bin (histo[bin]++) bukanlah satu operasi tunggal. Di tingkat perangkat keras, ini adalah urutan dari tiga instruksi:

  1. Read: Baca nilai lama dari histo[bin] ke dalam register thread.

  2. Modify: Tambahkan 1 ke nilai di register.

  3. Write: Tulis nilai baru dari register kembali ke histo[bin].

Skenario Bencana:

  • Thread A membaca nilai histo[5] (misalnya, 10).

  • Sebelum Thread A menulis kembali, Thread B juga membaca nilai histo[5] (masih 10).

  • Thread A menghitung 10+1=11, lalu menulis 11 ke histo[5].

  • Thread B menghitung 10+1=11, lalu menulis 11 ke histo[5].

Hasil: histo[5] bernilai 11, padahal seharusnya 12. Satu pembaruan telah hilang. Ini adalah bug yang sangat sulit dilacak karena hanya terjadi pada timing tertentu.

Solusi Perangkat Keras: Operasi Atomik

Operasi Atomik adalah instruksi perangkat keras khusus yang mengeksekusi urutan Read-Modify-Write secara tak terpisahkan (indivisible).

  • Jaminan: Ketika satu thread memulai operasi atomik pada sebuah alamat memori, perangkat keras akan “mengunci” alamat tersebut dan memastikan tidak ada thread lain yang dapat mengaksesnya sampai operasi atomik tersebut selesai sepenuhnya.

  • Serialisasi: Ini secara efektif memaksa semua thread yang ingin mengakses alamat yang sama untuk melakukannya secara berurutan (serial), sehingga menghilangkan data race dan menjamin kebenaran hasil.

  • Di CUDA: Disediakan melalui fungsi intrinsik seperti atomicAdd().

    atomicAdd(&address, value);

Kelemahan Performa Operasi Atomik

Meskipun menjamin kebenaran, operasi atomik bisa menjadi bottleneck performa yang serius, terutama jika dilakukan pada Global Memory.

  • Latensi Tinggi: Setiap operasi atomik global melibatkan siklus baca dan tulis penuh ke DRAM yang lambat.

  • Kontensi Tinggi: Jika banyak thread dari seluruh GPU (dari berbagai block) mencoba melakukan operasi atomik pada bin yang sama, mereka semua akan diserialisasi, secara efektif mengubah eksekusi paralel menjadi sekuensial pada titik tersebut dan menghancurkan throughput.

Operasi atomik pada Shared Memory jauh lebih cepat (sekitar 100x) karena latensinya yang sangat rendah.

Kinerja Tinggi: Teknik Privatization

Untuk menghindari biaya tinggi dari atomik global, kita menggunakan teknik Privatization yang cerdas. Idenya adalah membatasi konflik hanya di dalam lingkup yang cepat.

Alur Kerja:

  1. Setiap thread block mendeklarasikan histogram privat miliknya sendiri di Shared Memory.

  2. Thread-thread di dalam block memproses bagian input mereka dan memperbarui histogram privat di Shared Memory menggunakan atomicAdd() yang sangat cepat. Karena lingkupnya per-block, tidak ada konflik antar block.

  3. __syncthreads(): Menunggu semua thread dalam block selesai.

  4. Hanya beberapa thread (misalnya, thread pertama dari setiap bin) yang kemudian bertanggung jawab untuk menambahkan hasil dari histogram privat di Shared Memory ke histogram utama di Global Memory, menggunakan atomicAdd() yang lambat.

Keuntungan: Jumlah operasi atomik global yang lambat dan berpotensi konflik tinggi berkurang dari jutaan (satu per elemen input) menjadi hanya ribuan (satu per bin per block). Ini secara dramatis meningkatkan performa.

Summary

Histogram paralel menghadapi masalah kritis data race ketika banyak thread mencoba menginkremen bin yang sama melalui operasi read-modify-write. Solusi untuk menjamin kebenaran adalah Operasi Atomik, instruksi perangkat keras yang tak terpisahkan, namun bisa sangat lambat jika terjadi kontensi tinggi pada Global Memory. Teknik berkinerja tinggi yang disebut Privatization mengatasi ini dengan membuat setiap thread block membangun histogram privatnya di Shared Memory yang cepat, dan hanya pada akhirnya menggabungkan hasil-hasil privat ini ke histogram global, sehingga secara drastis mengurangi bottleneck dari operasi atomik global.

Graph View

Backlinks