Memahami dimensi grid CUDA, dimensi blok dan organisasi utas (penjelasan sederhana) [ditutup]

161

Bagaimana thread diatur untuk dijalankan oleh GPU?

cuda nvidia warga negara1
sumber
Panduan Pemrograman CUDA harus menjadi tempat yang baik untuk memulai ini. Saya juga merekomendasikan untuk memeriksa pengenalan CUDA dari sini .
Tom

Jawaban:

287

Perangkat keras

Jika perangkat GPU memiliki, misalnya, 4 unit multiprosesing, dan mereka dapat menjalankan masing-masing 768 utas: maka pada saat tertentu tidak lebih dari 4 * 768 utas akan benar-benar berjalan secara paralel (jika Anda merencanakan lebih banyak utas, mereka akan menunggu giliran mereka).

Perangkat lunak

utas disusun dalam blok. Blok dijalankan oleh unit multiprosesing. Utas blok dapat diidentifikasi (diindeks) menggunakan 1Dimensi (x), 2Dimensi (x, y) atau indeks 3Dim (x, y, z) tetapi dalam kasus apa pun x y z <= 768 untuk contoh kita (pembatasan lain berlaku untuk x, y, z, lihat panduan dan kemampuan perangkat Anda).

Jelas, jika Anda membutuhkan lebih dari 4 * 768 utas itu, Anda membutuhkan lebih dari 4 blok. Blok juga dapat diindeks 1D, 2D atau 3D. Ada antrian blok yang menunggu untuk memasuki GPU (karena, dalam contoh kami, GPU memiliki 4 multiprosesor dan hanya 4 blok yang dieksekusi secara bersamaan).

Sekarang kasus sederhana: memproses gambar 512x512

Misalkan kita ingin satu utas untuk memproses satu piksel (i, j).

Kita dapat menggunakan blok masing-masing 64 utas. Maka kita membutuhkan 512 * 512/64 = 4096 blok (jadi untuk memiliki 512x512 utas = 4096 * 64)

Sudah umum untuk mengatur (untuk mempermudah pengindeksan gambar) utas dalam blok 2D yang memiliki blockDim = 8 x 8 (64 utas per blok). Saya lebih suka menyebutnya threadsPerBlock.

dim3 threadsPerBlock(8, 8);  // 64 threads

dan 2D gridDim = 64 x 64 blok (4096 blok diperlukan). Saya lebih suka menyebutnya numBlocks.

dim3 numBlocks(imageWidth/threadsPerBlock.x,  /* for instance 512/8 = 64*/
              imageHeight/threadsPerBlock.y);

Kernel diluncurkan seperti ini:

myKernel <<<numBlocks,threadsPerBlock>>>( /* params for the kernel function */ );

Akhirnya: akan ada sesuatu seperti "antrian 4096 blok", di mana satu blok sedang menunggu untuk ditugaskan salah satu dari multiprosesor GPU untuk menjalankan 64 utasnya.

Dalam kernel piksel (i, j) yang akan diproses oleh utas dihitung dengan cara ini:

uint i = (blockIdx.x * blockDim.x) + threadIdx.x;
uint j = (blockIdx.y * blockDim.y) + threadIdx.y;
warga negara1
sumber
11
Jika setiap blok dapat menjalankan 768 utas, mengapa hanya menggunakan 64? Jika Anda menggunakan batas maksimal 768, Anda akan memiliki blok lebih sedikit dan kinerja jadi lebih baik.
Aliza
10
@Aliza: blok logis , batas 768 utas adalah untuk setiap unit pemrosesan fisik . Anda menggunakan blok sesuai dengan spesifikasi masalah Anda untuk mendistribusikan pekerjaan ke utas. Kemungkinan Anda tidak selalu dapat menggunakan blok 768 utas untuk setiap masalah yang Anda miliki. Bayangkan Anda harus memproses gambar 64x64 (4096 piksel). 4096/768 = 5.333333 blok?
cibercitizen1
1
blok logis, tetapi setiap blok ditugaskan ke inti. jika ada lebih banyak blok dari inti, blok-blok tersebut di-antri sampai inti menjadi bebas. Dalam contoh Anda, Anda dapat menggunakan 6 blok dan membuat utas tambahan tidak melakukan apa-apa (2/3 dari utas pada blok ke-6).
Aliza
3
@ cibercitizen1 - Saya pikir poin Aliza bagus: jika memungkinkan, seseorang ingin menggunakan sebanyak mungkin utas per blok. Jika ada kendala yang membutuhkan lebih sedikit utas, lebih baik untuk menjelaskan mengapa itu mungkin terjadi dalam contoh kedua (tetapi masih menjelaskan kasus yang lebih sederhana dan lebih diinginkan, pertama).
6
@kamu Ya, mungkin. Tetapi kasusnya adalah bahwa jumlah memori yang dibutuhkan oleh setiap utas tergantung pada aplikasi. Misalnya, dalam program terakhir saya, setiap utas memanggil fungsi optimisasi kuadrat-terkecil, yang membutuhkan "banyak" memori. Begitu banyak, blok itu tidak bisa lebih besar dari utas 4x4. Meski begitu, speedup yang diperoleh sangat dramatis, vs versi sekuensial.
cibercitizen1
9

Misalkan GPU 9800GT:

  • ini memiliki 14 multiprosesor (SM)
  • setiap SM memiliki 8 thread-prosesor (stream prosesor AKA, SP atau core)
  • memungkinkan hingga 512 utas per blok
  • warpsize adalah 32 (yang berarti masing-masing dari 14x8 = 112 thread-prosesor dapat menjadwalkan hingga 32 thread)

https://www.tutorialspoint.com/cuda/cuda_threads.htm

Blok tidak dapat memiliki utas lebih aktif dari 512 sehingga __syncthreadshanya dapat menyinkronkan utas dalam jumlah terbatas. yaitu Jika Anda menjalankan yang berikut dengan 600 utas:

func1();
__syncthreads();
func2();
__syncthreads();

maka kernel harus dijalankan dua kali dan urutan eksekusi adalah:

  1. func1 dijalankan untuk 512 utas pertama
  2. func2 dijalankan untuk 512 utas pertama
  3. func1 dieksekusi untuk utas yang tersisa
  4. func2 dijalankan untuk utas yang tersisa

catatan:

Titik utama __syncthreadsadalah operasi blok-lebar dan tidak menyinkronkan semua utas.

Saya tidak yakin tentang jumlah persis utas yang __syncthreadsdapat disinkronkan, karena Anda dapat membuat blok dengan lebih dari 512 utas dan membiarkan warp menangani penjadwalan. Untuk pemahaman saya itu lebih akurat untuk mengatakan: func1 dijalankan setidaknya untuk 512 utas pertama.

Sebelum saya mengedit jawaban ini (kembali pada tahun 2010) saya mengukur 14x8x32 utas yang disinkronkan menggunakan __syncthreads.

Saya akan sangat menghargai jika seseorang menguji ini lagi untuk informasi yang lebih akurat.

Bizhan
sumber
Apa yang terjadi jika func2 () tergantung pada hasil dari func1 (). Saya pikir ini salah
Chris
@ Chris saya menulis ini tujuh tahun lalu, tetapi jika saya ingat dengan benar saya melakukan tes ini dan mendapatkan kesimpulan ini bahwa kernel dengan lebih banyak thread daripada gpu berperilaku seperti ini. Jika Anda menguji kasus ini dan mencapai hasil yang berbeda maka saya harus menghapus posting ini.
Bizhan
Maaf saya pikir ini salah, juga, GPU hanya bisa menjalankan 112 utas secara bersamaan.
Steven Lu
@ SevenLu, apakah Anda sudah mencobanya? juga saya tidak berpikir 112 thread bersamaan masuk akal untuk GPU. 112 adalah jumlah pemroses arus. Saya hampir tidak dapat mengingat CUDA sekarang :)
Bizhan
1
@ SevenLu jumlah maksimum utas bukan masalah di sini, __syncthreadsadalah operasi blok-lebar dan fakta bahwa sebenarnya tidak menyinkronkan semua utas adalah gangguan bagi pelajar CUDA. Jadi saya memperbarui jawaban saya berdasarkan informasi yang Anda berikan kepada saya. Saya sangat menghargai itu.
Bizhan