Get best of the week

Komputasi GPU - Mengapa, Kapan, dan Bagaimana. Ditambah beberapa tes

Semua orang sudah lama tahu bahwa pada kartu video Anda tidak hanya bisa bermain mainan, tetapi juga melakukan hal-hal yang tidak terkait dengan game, misalnya, melatih jaringan saraf, mengingat cryptocurrency atau melakukan perhitungan ilmiah. Bagaimana ini terjadi, Anda dapat membacanya di sini , tetapi saya ingin menyentuh pada topik mengapa GPU mungkin menarik bagi programmer rata - rata (tidak terkait dengan GameDev) cara mendekati pengembangan pada GPU tanpa menghabiskan banyak waktu untuk itu, memutuskan apakah lihat ke arah ini, dan "cari tahu dengan jari Anda" berapa untung yang bisa Anda peroleh. 



Artikel ini ditulis berdasarkan presentasi saya di HighLoad ++. Ini terutama membahas teknologi yang ditawarkan oleh NVIDIA. Saya tidak punya tujuan untuk mengiklankan produk apa pun, saya hanya memberikan mereka sebagai contoh, dan pasti sesuatu yang serupa dapat ditemukan di produsen yang bersaing.

Mengapa mengandalkan GPU?


Dua prosesor dapat dibandingkan sesuai dengan kriteria yang berbeda, mungkin yang paling populer adalah frekuensi dan jumlah core, ukuran cache, dll., Tetapi pada akhirnya, kami tertarik pada berapa banyak operasi yang dapat dilakukan prosesor per unit waktu, seperti apa operasi ini, tetapi pertanyaan terpisah Metrik umum adalah jumlah operasi titik apung per detik - jepit. Dan ketika kita ingin membandingkan hangat dengan lunak, dan dalam kasus GPU kami dengan CPU, metrik ini sangat berguna.

Grafik di bawah ini menunjukkan pertumbuhan jepit yang sama ini dari waktu ke waktu untuk prosesor dan kartu video.


(Data dikumpulkan dari sumber terbuka, tidak ada data untuk 2019-20 tahun, karena tidak semuanya begitu indah di sana, tetapi GPU masih menang)

Yah, itu menggoda, bukan? Kami menggeser semua perhitungan dari CPU ke GPU dan mendapatkan delapan kali kinerja terbaik!

Tapi, tentu saja, tidak semuanya begitu sederhana. Anda tidak bisa hanya mengambil dan mentransfer semuanya ke GPU, mengapa, kami akan berbicara lebih lanjut.

Arsitektur GPU dan perbandingannya dengan CPU


Saya membawa banyak gambar yang akrab dengan arsitektur CPU dan elemen dasar:


CPU Core

Apa yang istimewa? Satu inti dan banyak blok bantu.

Sekarang mari kita lihat arsitektur GPU:


GPU Core

Kartu video memiliki banyak inti pemrosesan, biasanya beberapa ribu, tetapi digabungkan menjadi beberapa blok, untuk kartu video NVIDIA, biasanya 32 masing-masing, dan memiliki elemen umum, termasuk dan register. Arsitektur inti GPU dan elemen logis jauh lebih sederhana daripada pada CPU, yaitu, tidak ada prefetcher, prediktor brunch, dan banyak lagi.

Nah, ini adalah poin utama dari perbedaan dalam arsitektur CPU dan GPU, dan, pada kenyataannya, mereka memberlakukan batasan atau, sebaliknya, membuka kemungkinan untuk apa yang dapat kita baca secara efektif di GPU.

Saya tidak menyebutkan satu poin penting lagi, biasanya, kartu video dan prosesor tidak “mencari-cari” di antara mereka sendiri dan menulis data ke kartu video dan membaca hasilnya kembali - ini adalah operasi yang terpisah dan dapat berubah menjadi “bottleneck” di sistem Anda, grafik waktu pemompaan versus ukuran data diberikan kemudian di artikel.

Batasan dan fitur GPU


Apa keterbatasan arsitektur ini pada algoritma yang dapat dieksekusi:

  • Jika kita menghitung pada GPU, maka kita tidak dapat memilih hanya satu inti, seluruh blok inti akan dialokasikan (32 untuk NVIDIA).
  • Semua core menjalankan instruksi yang sama, tetapi dengan data yang berbeda (kami akan membicarakannya nanti), perhitungan seperti itu disebut Single-Instruction-Multiple-Data atau SIMD (walaupun NVIDIA memperkenalkan perbaikannya). 
  • Karena serangkaian blok logika dan register umum yang relatif sederhana, GPU benar-benar tidak suka bercabang, dan memang logika yang rumit dalam algoritma.

Peluang apa yang terbuka:

  • Sebenarnya, percepatan perhitungan SIMD yang sama. Contoh paling sederhana adalah penambahan matriks elemen, dan mari kita menganalisisnya.

Pengurangan algoritma klasik ke representasi SIMD


Transformasi


Kami memiliki dua larik, A dan B, dan kami ingin menambahkan elemen dari larik B ke setiap elemen larik A. Di bawah ini adalah contoh dalam C, meskipun saya harap akan jelas bagi mereka yang tidak berbicara bahasa ini:

void func(float *A, float *B, size)
{ 
   for (int i = 0; i < size; i++) 
   { 
       A[i] += B[i]
   } 
}

Loopback klasik elemen dalam loop dan runtime linier.

Sekarang mari kita lihat bagaimana kode tersebut akan mencari GPU:

void func(float *A, float *B, size) 
{ 
   int i = threadIdx.x; 
   if (i < size) 
      A[i] += B[i] 
}

Dan di sini sudah menarik, variabel threadIdx muncul, yang sepertinya tidak kami nyatakan di mana pun. Ya, sistemnya menyediakan bagi kita. Bayangkan dalam contoh sebelumnya array terdiri dari tiga elemen, dan Anda ingin menjalankannya dalam tiga utas paralel. Untuk melakukan ini, Anda perlu menambahkan parameter lain - indeks atau nomor streaming. Inilah yang dilakukan kartu video untuk kami, meskipun ia melewati indeks sebagai variabel statis dan dapat bekerja dengan beberapa dimensi sekaligus - x, y, z.

Nuansa lain, jika Anda akan memulai sejumlah besar aliran paralel sekaligus, maka aliran tersebut harus dibagi menjadi beberapa blok (fitur arsitektur kartu video). Ukuran blok maksimum tergantung pada kartu video, dan indeks elemen tempat kami melakukan perhitungan perlu diperoleh sebagai berikut:

int i = blockIdx.x * blockDim.x + threadIdx.x; // blockIdx –  , blockDim –  , threadIdx –

Akibatnya, apa yang kita miliki: banyak utas berjalan paralel yang mengeksekusi kode yang sama, tetapi dengan indeks yang berbeda, dan, dengan demikian, data, mis. SIMD yang sama.

Ini adalah contoh paling sederhana, tetapi jika Anda ingin bekerja dengan GPU, Anda harus membawa tugas Anda ke bentuk yang sama. Sayangnya, ini tidak selalu memungkinkan dan dalam beberapa kasus dapat menjadi subjek disertasi doktoral, tetapi bagaimanapun, algoritma klasik masih dapat dibawa ke formulir ini.

Pengumpulan


Sekarang mari kita lihat bagaimana agregasi yang dilemparkan ke representasi SIMD akan terlihat:
 

Kami memiliki array n elemen. Pada tahap pertama, kita mulai n / 2 utas dan setiap utas menambahkan dua elemen, yaitu dalam satu iterasi, kami menambahkan bersama setengah dari elemen dalam array. Dan kemudian dalam loop kita mengulangi hal yang sama untuk array yang baru dibuat, sampai kita menggabungkan dua elemen terakhir. Seperti yang Anda lihat, semakin kecil ukuran array, semakin sedikit thread paralel yang dapat kita mulai, mis. pada GPU, masuk akal untuk menggabungkan array dengan ukuran yang cukup besar. Algoritme seperti itu dapat digunakan untuk menghitung jumlah elemen (omong-omong, jangan lupa tentang kemungkinan melimpahnya jenis data yang Anda kerjakan), cari maksimum, minimum, atau hanya pencarian.

Penyortiran


Tetapi memilah sudah terlihat jauh lebih rumit.

Dua algoritma pemilahan paling populer pada GPU adalah:

  • Semacam bitonic
  • Radix-sort

Tetapi radix-sort masih lebih sering digunakan, dan implementasi siap produksi dapat ditemukan di beberapa perpustakaan. Saya tidak akan menganalisis secara detail bagaimana algoritma ini bekerja; mereka yang tertarik dapat menemukan deskripsi radix-sort di https://www.codeproject.com/Articles/543451/Parallel-Radix-Sort-on-the-GPU-using-Cplusplus- AMP dan https://stackoverflow.com/a/26229897

Tapi idenya adalah bahwa bahkan algoritma non-linear seperti penyortiran dapat dikurangi menjadi tampilan SIMD.

Dan sekarang, sebelum melihat bilangan real yang dapat diperoleh dari GPU, mari kita mencari tahu cara memprogram untuk keajaiban teknologi ini?

Mulai dari mana


Dua teknologi paling umum yang dapat digunakan untuk pemrograman di bawah GPU:

  • Opencl
  • Cuda

OpenCL adalah standar yang didukung oleh sebagian besar produsen kartu video, termasuk dan pada perangkat seluler, juga kode yang ditulis dalam OpenCL dapat dijalankan pada CPU.

Anda dapat menggunakan OpenCL dari C / C ++, ada binder ke bahasa lain.

Untuk OpenCL, saya paling menyukai buku OpenCL in Action . Ini juga menjelaskan berbagai algoritma pada GPU, termasuk Semacam bitonic dan semacam Radix.

CUDA adalah teknologi dan SDK milik NVIDIA. Anda dapat menulis dalam C / C ++ atau menggunakan binding ke bahasa lain.

Membandingkan OpenCL dan CUDA agak tidak benar, karena satu adalah standar, yang lainnya adalah seluruh SDK. Namun demikian, banyak orang memilih CUDA untuk pengembangan untuk kartu video, terlepas dari kenyataan bahwa teknologinya eksklusif, meskipun gratis dan hanya berfungsi pada kartu NVIDIA. Ada beberapa alasan untuk ini:

  • API
  • , GPU, (host)
  • , ..

Kekhasan termasuk fakta bahwa CUDA datang dengan kompiler sendiri, yang juga dapat mengkompilasi kode C / C ++ standar.

Buku CUDA paling komprehensif yang saya temui adalah Pemrograman CUDA C Profesional , meskipun sudah agak ketinggalan jaman, namun buku ini membahas banyak nuansa teknis pemrograman untuk kartu NVIDIA.

Tetapi bagaimana jika saya tidak ingin menghabiskan beberapa bulan membaca buku-buku ini, menulis program saya sendiri untuk kartu video, menguji dan men-debug, dan kemudian mengetahui bahwa ini bukan untuk saya? 

Seperti yang saya katakan, ada sejumlah besar perpustakaan yang menyembunyikan kompleksitas pengembangan di bawah GPU: XGBoost, cuBLAS, TensorFlow, PyTorch dan lainnya, kami akan mempertimbangkan perpustakaan dorong, karena kurang terspesialisasi dari perpustakaan lain di atas, tetapi pada saat yang sama mengimplementasikan algoritma dasar, misalnya, pengurutan, pencarian, agregasi, dan dengan probabilitas tinggi dapat diterapkan dalam tugas Anda.

Thrust adalah pustaka C ++ yang bertujuan untuk "mengganti" algoritma STL standar dengan algoritma berbasis GPU. Misalnya, mengurutkan array angka menggunakan pustaka ini pada kartu video akan terlihat seperti ini:

thrust::host_vector<DataType> h_vec(size); //    
std::generate(h_vec.begin(), h_vec.end(), rand); //   
thrust::device_vector<DataType> d_vec = h_vec; //         
thrust::sort(d_vec.begin(), d_vec.end()); //    
thrust::copy(d_vec.begin(), d_vec.end(), h_vec.begin()); //   ,

(jangan lupa bahwa contoh harus dikompilasi oleh kompiler dari NVIDIA)

Seperti yang Anda lihat, dorong :: sort sangat mirip dengan algoritma serupa dari STL. Pustaka ini menyembunyikan banyak kesulitan, khususnya pengembangan subprogram (lebih tepatnya, kernel), yang akan dieksekusi pada kartu video, tetapi pada saat yang sama menghilangkan fleksibilitas. Sebagai contoh, jika kita ingin mengurutkan beberapa gigabyte data, akan logis untuk mengirim sepotong data ke kartu untuk mulai menyortir, dan sementara penyortiran sedang berlangsung, kirim lebih banyak data ke kartu. Pendekatan ini disebut bersembunyi latensi dan memungkinkan penggunaan sumber daya peta server yang lebih efisien, tetapi, sayangnya, ketika kita menggunakan pustaka tingkat tinggi, peluang seperti itu tetap tersembunyi. Tetapi untuk membuat prototipe dan mengukur kinerja, semuanya tetap sama, terutama dengan dorongan Anda dapat mengukur overhead yang disediakan oleh transfer data.

Saya menulis patokan kecil menggunakan pustaka ini, yang menjalankan beberapa algoritma populer dengan jumlah data yang berbeda pada GPU, mari kita lihat apa hasilnya.

Hasil Algoritma GPU


Untuk menguji GPU, saya mengambil contoh di AWS dengan kartu grafis Tesla K80, ini bukan kartu server paling kuat saat ini (Tesla v100 paling kuat), tetapi yang paling terjangkau dan ada di papan:

  • 4992 kernel CUDA
  • Memori 24 GB
  • 480 Gb / s - bandwidth memori 

Dan untuk pengujian pada CPU, saya mengambil contoh dengan prosesor Intel Xeon CPU E5-2686 v4 @ 2.30GHz

Transformasi



Transformasi waktu eksekusi pada GPU dan CPU dalam ms.

Seperti yang Anda lihat, transformasi biasa elemen array kira-kira sama dalam waktu, baik pada GPU dan pada CPU. Dan mengapa? Karena overhead untuk mengirim data ke kartu dan kembali memakan seluruh peningkatan kinerja (kita akan berbicara tentang overhead secara terpisah), dan ada relatif sedikit perhitungan pada kartu. Juga, jangan lupa bahwa prosesor juga mendukung instruksi SIMD, dan kompiler dalam kasus sederhana dapat menggunakannya secara efektif. 

Sekarang mari kita lihat seberapa efisien agregasi dilakukan pada GPU.

Pengumpulan



Waktu eksekusi agregasi pada GPU dan CPU dalam ms.

Dalam contoh agregasi, kita sudah melihat peningkatan kinerja yang signifikan dengan peningkatan volume data. Penting juga untuk memperhatikan fakta bahwa kami memompa sejumlah besar data ke dalam memori kartu, dan hanya satu nilai agregat yang diambil kembali, mis. Overhead untuk mentransfer data dari kartu ke RAM minimal.

Mari kita beralih ke contoh paling menarik - pengurutan.

Penyortiran



Menyortir waktu ke GPU dan CPU dalam ms.

Terlepas dari kenyataan bahwa kami mengirim seluruh array data ke kartu video dan sebaliknya, menyortir ke GPU 800 MB data kira-kira 25 kali lebih cepat daripada pada prosesor.

Transfer data overhead


Seperti dapat dilihat dari contoh transformasi, tidak selalu jelas apakah GPU akan efektif bahkan dalam tugas-tugas yang paralel dengan baik. Alasannya adalah overhead untuk mentransfer data dari RAM komputer ke memori kartu video (di konsol game, omong-omong, memori dibagi antara CPU dan GPU, dan tidak perlu mentransfer data). Salah satu karakteristik kartu video adalah bandwidth memori atau bandwidth memori, yang menentukan bandwidth teoritis kartu. Untuk Tesla k80 itu adalah 480 GB / s, untuk Tesla v100 sudah 900 GB / s. Juga, versi PCI Express dan implementasi bagaimana Anda akan mentransfer data ke kartu akan mempengaruhi throughput, misalnya, ini dapat dilakukan dalam beberapa aliran paralel.

Mari kita lihat hasil praktis yang diperoleh untuk kartu grafis Tesla K80 di cloud Amazon:


Waktu untuk mentransfer data ke GPU, mengurutkan dan mentransfer data kembali ke RAM dalam ms

HtoD - mentransfer data ke kartu video

GPU Eksekusi - mengurutkan pada kartu video

DtoH - menyalin data dari kartu video ke RAM

Hal pertama yang perlu diperhatikan adalah membaca data dari kartu video lebih cepat daripada tuliskan di sana.

Yang kedua - ketika bekerja dengan kartu video, Anda bisa mendapatkan latensi dari 350 mikrodetik, dan ini mungkin sudah cukup untuk beberapa aplikasi latensi rendah.

Bagan di bawah ini menunjukkan overhead untuk lebih banyak data:


Saatnya mentransfer data ke GPU, mengurutkan dan mentransfer data kembali ke RAM dalam ms

Penggunaan server


Pertanyaan yang paling umum adalah bagaimana perbedaan kartu video game dari yang server? Menurut karakteristik, mereka sangat mirip, tetapi harga berbeda secara signifikan.


Perbedaan utama antara server (NVIDIA) dan kartu permainan:

  • Garansi pabrik (kartu game tidak dirancang untuk penggunaan server)
  • Kemungkinan masalah virtualisasi untuk kartu grafis konsumen
  • Ketersediaan mekanisme koreksi kesalahan pada kartu server
  • Jumlah utas paralel (bukan inti CUDA) atau dukungan untuk Hyper-Q, yang memungkinkan Anda untuk bekerja dengan kartu dari beberapa utas pada CPU, misalnya, mengunggah data ke kartu dari satu utas dan memulai perhitungan dari yang lain

Inilah, mungkin, perbedaan penting utama yang saya temukan.

Multithreading


Setelah kami menemukan cara menjalankan algoritma paling sederhana pada kartu video dan hasil apa yang bisa diharapkan, pertanyaan logis berikutnya adalah bagaimana kartu video akan berperilaku ketika memproses beberapa permintaan paralel. Sebagai jawaban, saya memiliki dua grafik komputasi pada GPU dan prosesor dengan 4 dan 32 core:


Waktu yang diperlukan untuk melakukan perhitungan matematis pada GPU dan CPU dengan matriks 1000 x 60 dalam ms

. Grafik ini melakukan perhitungan dengan matriks 1000 x 60 elemen. Perhitungan dimulai dari beberapa aliran program, aliran terpisah dibuat untuk GPU untuk setiap aliran CPU (Hyper-Q digunakan). 

Seperti yang Anda lihat, prosesor mengatasi beban ini dengan sangat baik, sedangkan latensi untuk satu permintaan per GPU meningkat secara signifikan dengan jumlah permintaan paralel.


Waktu untuk melakukan perhitungan matematis pada GPU dan CPU dengan matriks 10.000 x 60 dalam ms.

Pada grafik kedua, perhitungan yang sama, tetapi dengan matriks 10 kali lebih lama, dan GPU berperilaku jauh lebih baik di bawah beban seperti itu. Grafik ini sangat indikatif, dan kita dapat menyimpulkan: perilaku di bawah beban tergantung pada sifat beban itu sendiri. Sebuah prosesor juga dapat menangani perhitungan matriks dengan cukup efisien, tetapi sampai batas tertentu. Untuk kartu video, merupakan karakteristik bahwa untuk beban komputasi yang kecil, kinerjanya turun sekitar secara linear. Dengan peningkatan beban dan jumlah utas paralel, kartu video berupaya lebih baik. 

Sulit untuk membuat hipotesis bagaimana GPU akan berperilaku dalam berbagai situasi, tetapi seperti yang Anda lihat, dalam kondisi tertentu, kartu server dapat memproses permintaan dari beberapa aliran paralel dengan cukup efisien.

Kami akan membahas beberapa pertanyaan lagi yang mungkin Anda miliki jika Anda masih memutuskan untuk menggunakan GPU dalam proyek Anda.

Batas sumber daya


Seperti yang telah kami katakan, dua sumber utama kartu video adalah komputasi core dan memori.

Misalnya, kami memiliki beberapa proses atau wadah menggunakan kartu video, dan kami ingin dapat membagikan kartu video di antara mereka. Sayangnya, tidak ada API sederhana untuk ini. NVIDIA menawarkan teknologi vGPU , tetapi saya tidak menemukan kartu Tesla K80 dalam daftar yang didukung, dan sejauh yang saya bisa mengerti dari deskripsi, teknologi ini lebih fokus pada tampilan virtual daripada pada perhitungan. Mungkin AMD menawarkan sesuatu yang lebih cocok.

Oleh karena itu, jika Anda berencana untuk menggunakan GPU dalam proyek Anda, Anda harus bergantung pada kenyataan bahwa aplikasi akan menggunakan kartu video secara eksklusif, atau Anda akan secara terprogram mengontrol jumlah memori yang dialokasikan dan jumlah core yang digunakan untuk perhitungan.

Wadah dan GPU


Jika Anda mengetahui batas sumber daya, maka pertanyaan logis berikut: bagaimana jika ada beberapa kartu video di server?

Sekali lagi, Anda dapat memutuskan pada level aplikasi GPU mana yang akan digunakan.

Cara lain yang lebih nyaman adalah wadah Docker. Anda dapat menggunakan wadah reguler, tetapi NVIDIA menawarkan wadah NGC -nya , dengan versi yang dioptimalkan dari berbagai perangkat lunak, perpustakaan, dan driver. Untuk satu kontainer, Anda dapat membatasi jumlah GPU yang digunakan dan visibilitasnya ke kontainer. Overhead pada penggunaan kontainer adalah sekitar 3%.

Bekerja dalam sebuah cluster


Pertanyaan lain, apa yang harus dilakukan jika Anda ingin melakukan satu tugas pada beberapa GPU dalam server atau cluster yang sama?

Jika Anda memilih perpustakaan yang mirip dengan dorong atau solusi tingkat rendah, maka tugas harus diselesaikan secara manual. Kerangka kerja tingkat tinggi, misalnya, untuk pembelajaran mesin atau jaringan saraf, biasanya mendukung kemampuan untuk menggunakan banyak kartu di luar kotak.

Selain itu, saya ingin mencatat bahwa, misalnya, NVIDIA menawarkan antarmuka untuk pertukaran data langsung antar kartu - NVLINK , yang secara signifikan lebih cepat daripada PCI Express. Dan ada teknologi untuk akses langsung ke memori kartu dari perangkat PCI Express lainnya - GPUDirect RDMA , termasuk. dan jaringan .

Rekomendasi


Jika Anda mempertimbangkan untuk menggunakan GPU dalam proyek Anda, maka GPU kemungkinan besar cocok untuk Anda jika:

  • Tugas Anda dapat dikurangi menjadi tampilan SIMD
  • Dimungkinkan untuk memuat sebagian besar data di peta sebelum perhitungan (cache)
  • Tantangannya melibatkan komputasi intensif

Anda juga harus mengajukan pertanyaan terlebih dahulu:

  • Berapa banyak permintaan paralel akan 
  • Apa latensi yang Anda harapkan
  • Apakah Anda memerlukan satu kartu untuk memuat Anda? Apakah Anda memerlukan server dengan beberapa kartu atau sekelompok server GPU 

Itu saja, saya berharap materi tersebut akan berguna bagi Anda dan membantu Anda membuat keputusan yang tepat!

Referensi


Benchmark dan hasil pada github - https://github.com/tishden/gpu_benchmark/tree/master/cuda

Selain topik, rekaman laporan "Basis Data GPU - Arsitektur, Kinerja, dan Prospek untuk Penggunaan"

NVIDIA NGC Containers Webinar - http : //bit.ly/2UmVIVt atau http://bit.ly/2x4vJKF

More articles:


All Articles