BerandaComputers and TechnologyCatatan tentang Bercabang Dalam Shader

Catatan tentang Bercabang Dalam Shader

An keabadian yang lalu, saya menerbitkan entri blog tentang shader yang saya tulis. Di postingan itu, saya dengan santai mengulangi sepotong cerita rakyat gpu yang saya ambil sebagai penggemar wee opengl hampir satu dekade yang lalu,

“Menggunakan pernyataan if di dalam shader akan menyebabkan penurunan kinerja.”

Orang-orang di internet tampak agak skeptis tentang pernyataan ini, dan saya menyadari bahwa saran saya mungkin sudah ketinggalan zaman. Jadi posting ini adalah tindak lanjut singkat tentang biaya percabangan pada GPU yang lebih modern.

Sebelum kita berbicara tentang percabangan, apa itu GPU? GPU atau unit pemrosesan grafis, adalah perangkat keras khusus yang awalnya dikembangkan untuk menghitung angka untuk kalkulasi grafis. Ini menyelesaikan tujuan ini dengan cepat melalui sejumlah besar paralelisasi.

GPU dapat menjalankan banyak utas sekaligus secara paralel. Utas ini umumnya dieksekusi dalam grup yang disebut warps (CUDA), invocations (Vulkan), dan gelombang (saya akan menggunakan istilah warp, tetapi semuanya dapat dipertukarkan). Pada perangkat keras Nvidia baru-baru ini (Ampere / Volta / Pascal) lungsin ini berisi 32 utas. Pascal misalnya secara teoritis dapat menjalankan hingga empat instruksi independen per warp lebih dari 56 lungsin dari 32 utas yang keluar ke pikiran meniup 7168 instruksi per siklus.

Setiap warp hanya dapat menjalankan satu instruksi dalam satu waktu. Namun instruksi itu dapat dijalankan untuk setiap utas di warp. Ini berarti GPU dapat mengeksekusi 32 salinan dari instruksi yang sama secara paralel untuk setiap utas di warp sekaligus.

Namun, mengambil kedua sisi cabang akan menyebabkan utas di dalam warp menjadi “menyimpang”. Ini berarti beberapa utas perlu mengeksekusi satu sisi cabang dan beberapa perlu mengeksekusi sisi lainnya. Sayangnya, kedua instruksi tersebut tidak dapat dijalankan secara bersamaan untuk utas di warp yang sama. Jadi instruksi yang berbeda ini harus dijalankan secara berurutan.

Gambar di atas diambil dari whitepaper Volta.

Pertimbangkan shader fragmen berikut :
void mainImage (keluar vec4 fragColor, di vec2 fragCoord) {

// SESUAIKAN JUMLAH INI DENGAN KEKUATAN GPU ANDA
int workAmount=2000;
float incr=1. / float (workAmount);
float outColor=0,0;

// GUNAKAN VARIABEL INI UNTUK MENGGUNAKAN CABANG
bool branch=true;

if (mod (fragCoord.x, 2.)

untuk (int i=0; i

outColor +=incr;

}
fragColor=vec4 (0.0, outColor, 0.0,1.0);

} lain {

untuk (int i=0; i

outColor +=incr;

}
fragColor=vec4 (outColor, 0.0,0.0,1.0);

}

}

Dalam shader fragmen seperti yang di atas, setiap lengkungan terdiri dari sekumpulan piksel yang dekat secara spasial. Setiap warp di shader ini harus divergen karena setiap piksel yang berdekatan secara horizontal mengambil sisi cabang yang berbeda. Misalkan biaya setiap loop dalam adalah n siklus. Karena setiap sisi cabang dijalankan secara seri, setiap warp harus mengambil setidaknya 2n siklus, meskipun setiap utas hanya menggunakan satu sisi cabang. Efek ini berskala dengan jumlah utas di jalur berbeda yang diambil (hingga 32x di perangkat keras Nvidia). Inilah shader lain yang seharusnya memiliki sekitar empat cabang per warp. Shader ini akan mengambil setidaknya 4n siklus.

Namun, mari kita lihat shader kedua :

void mainImage (keluar vec4 fragColor, di vec2 fragCoord) {

// GUNAKAN VARIABEL INI UNTUK MENYALA CABANG
int workAmount=2000;
float incr=1. / float (workAmount);
float outColor=0,0;

// GUNAKAN VARIABEL INI UNTUK MENGGUNAKAN CABANG
bool branch=true;

if (fragCoord.x

untuk (int i=0; i

outColor +=incr;

}
fragColor=vec4 (0.0, outColor, 0.0,1.0);

} lain {

untuk (int i=0; i

outColor +=incr;

}
fragColor=vec4 (outColor, 0.0,0.0,1.0);

}

}

Shader ini harus berjalan kira-kira dua kali lebih cepat dari shader lainnya meskipun jumlah piksel yang kira-kira sama untuk setiap cabang. Hal ini karena sebagian besar lengkungan tidak menyimpang (dengan pengecualian beberapa di sekitar setengah layar).

Spesifik Nvidia
Sejauh yang saya tahu, kira-kira ada dua cara berbeda dari kartu grafis nvidia menangani divergensi dengan dalam sebuah warp.

Thread Masking (Sebelum Volta)
Pada kartu pra volta Nvidia, divergensi ditangani oleh sesuatu yang disebut thread masking.

Pada dasarnya ketika kondisional tercapai, thread mask dibuat. Kita dapat menganggap mask ini sebagai array 32 entri. Setiap entri menentukan apakah utas terkait mengambil sisi pertama dari cabang atau tidak. Sisi pertama kemudian dieksekusi untuk setiap utas tersebut. Setelah sisi pertama dari cabang dieksekusi, mask dibalik dan sisi lain dari cabang dieksekusi.

Strategi topeng utas ini ada sebagian karena hanya ada satu penghitung program per warp di pra volta gpus. Jadi penting setiap utas di warp harus berada di tempat yang sama dalam program. Satu-satunya cara kita mengetahui utas mana yang akan dieksekusi, adalah dengan menggunakan topeng.

Gambar di atas diambil dari whitepaper Volta.

Perlu dicatat bahwa hanya utas yang perlu mengeksekusi cabang yang benar-benar ditetapkan ke inti komputasi. Jadi jika tidak ada utas yang mengeksekusi satu sisi cabang, tidak ada pekerjaan yang dilakukan.

Penjadwalan Untaian Independen (Volta and Beyond)
Satu perubahan besar dalam arsitektur Volta dan seterusnya adalah pengenalan “Penjadwalan Benang Independen”.

Dalam penjadwalan utas independen, setiap utas memiliki penghitung programnya sendiri. Hal ini memungkinkan setiap warp untuk melacak eksekusi setiap utas pada tingkat butiran halus. Namun, kami masih hanya dapat menjalankan satu instruksi dalam satu waktu. Jadi runtime sebenarnya dari sebuah cabang tidak berubah.

Jadi jika Penjadwalan Utas Independen tidak mempercepat percabangan, apa fungsinya? Penjadwalan Benang Independen memungkinkan kedua sisi cabang untuk dieksekusi secara bersamaan tetapi tidak secara paralel.

Gambar di atas diambil dari whitepaper Volta.

Alasan pergeseran ini karena sebelum Volta, cabang adalah tempat di mana kebuntuan yang menipu bisa terjadi.
id_pemimpin=0
Jika (threadIdx.x==leader_id) {

// Lakukan pekerjaan awal

} lain {

// Tunggu sampai pemimpin menyelesaikan pekerjaannya
// lalu lakukan pekerjaanku

}

Dalam kode pseudo shader komputasi di atas, kita memiliki 31 utas “pengikut” yang menunggu “utas pemimpin” selesai (menggunakan memori lokal bersama in-warp) sebelum menjalankan instruksinya sendiri.

Pada GPU Pascal atau di bawahnya, jika sisi lain dari cabang dijalankan lebih dulu, akan terjadi kebuntuan. Pemimpin tidak akan pernah mendapat kesempatan untuk memulai pekerjaannya karena pengikutnya akan menunggu.

Penjadwalan Benang Independen mengatasi masalah penguncian ini dengan menghubungkan dua jalur divergen.

Artikel TLDR
Cabang divergensi di warp membuat semua bagian dari cabang dieksekusi secara berurutan yang memperlambat shader. Divergensi percabangan antara warps tidak mempengaruhi runtime.

Bacaan lebih lanjut
Eksekusi cabang pada gpu sangat dalam, dan ada banyak kasus tepi yang tidak saya tangani dan saya juga tidak menyentuh vendor gpu selain Nvidia. Berikut beberapa sumber yang bagus jika Anda ingin mempelajari lebih lanjut:

The Volta Whitepaper (Sebagian besar posting ini dapat ditemukan antara angka 20 dan 22)
Memahami Pascal GPU Instruction Pipeline
Ada apa dengan Cabang Saya di GPU (posting ini mencakup beberapa kasus tepi setan, termasuk beberapa di sekitar pengambilan sampel tekstur)

Punya pertanyaan / komentar / koreksi?
Hubungi: pstefek.dev@gmail.com

Read More

RELATED ARTICLES

LEAVE A REPLY

Please enter your comment!
Please enter your name here

Most Popular

Recent Comments