Jenis masalah apa yang cocok untuk komputasi GPU?

84

Jadi saya punya kepala yang layak untuk masalah apa yang saya kerjakan yang terbaik dalam serial, dan mana yang bisa dikelola secara paralel. Tapi saat ini, saya tidak punya banyak ide tentang apa yang terbaik ditangani oleh perhitungan berbasis CPU, dan apa yang harus diturunkan ke GPU.

Saya tahu ini pertanyaan mendasar, tetapi sebagian besar pencarian saya terperangkap dalam orang-orang yang dengan jelas mengadvokasi satu atau yang lain tanpa benar-benar membenarkan alasannya , atau aturan praktis yang agak kabur. Mencari respons yang lebih bermanfaat di sini.

Fomite
sumber

Jawaban:

63

Perangkat keras GPU memiliki dua kekuatan khusus: raw compute (FLOPs) dan bandwidth memori. Masalah komputasi yang paling sulit termasuk dalam salah satu dari dua kategori ini. Sebagai contoh, aljabar linier padat (A * B = C atau Memecahkan [Ax = y] atau Mendiagonalisasi [A], dll.) Berada di suatu tempat pada spektrum bandwidth komputasi / memori tergantung pada ukuran sistem. Fast Fourier transforms (FFT) juga sesuai dengan cetakan ini dengan kebutuhan bandwidth agregat tinggi. Seperti halnya transformasi lain, algoritma berbasis grid / mesh, Monte Carlo, dll. Jika Anda melihat contoh kode NVIDIA SDK , Anda dapat merasakan berbagai masalah yang paling sering ditangani.

Saya pikir jawaban yang lebih instruktif adalah untuk pertanyaan 'Masalah apa yang benar-benar buruk pada GPU?' Sebagian besar masalah yang tidak termasuk dalam kategori ini dapat dibuat untuk berjalan pada GPU, meskipun beberapa membutuhkan lebih banyak usaha daripada yang lain.

Masalah yang tidak dipetakan dengan baik umumnya terlalu kecil atau terlalu tidak terduga. Masalah yang sangat kecil tidak memiliki paralelisme yang diperlukan untuk menggunakan semua utas pada GPU dan / atau bisa masuk ke cache tingkat rendah pada CPU, secara substansial meningkatkan kinerja CPU. Masalah yang tidak dapat diprediksi memiliki banyak cabang yang bermakna, yang dapat mencegah data mengalir secara efisien dari memori GPU ke inti atau mengurangi paralelisme dengan memecah paradigma SIMD (lihat ' warps berbeda '). Contoh dari jenis masalah ini termasuk:

  • Sebagian besar algoritma grafik (terlalu sulit diprediksi, terutama di ruang memori)
  • Aljabar linier yang jarang (tetapi ini juga buruk pada CPU)
  • Masalah pemrosesan sinyal kecil (FFT lebih kecil dari 1000 poin, misalnya)
  • Cari
  • Menyortir
Max Hutchinson
sumber
3
Namun, solusi GPU untuk masalah - masalah yang "tidak dapat diprediksi" itu mungkin dan, sementara saat ini biasanya tidak layak, mungkin mendapatkan signifikansi di masa depan.
leftaroundabout
6
Saya ingin secara khusus menambahkan cabang ke daftar pemecah kinerja GPU. Anda ingin semua (ratusan) Anda menjalankan instruksi yang sama (seperti pada SIMD) untuk melakukan komputasi paralel yang benar-benar. Sebagai contoh, pada kartu AMD jika ada aliran instruksi yang bertemu cabang dan harus menyimpang - semua muka gelombang (grup paralel) menyimpang. Jika unit lain dari muka gelombang tidak boleh menyimpang - mereka harus melakukan lintasan kedua. Itulah yang dimaksud maxhutch dengan prediktabilitas, saya kira.
Violet Giraffe
2
@ VioletGiraffe, itu belum tentu benar. Dalam CUDA (yaitu pada GPU Nvidia), divergensi cabang hanya memengaruhi warp saat ini, yang paling banyak 32 utas. Lengkungan yang berbeda, meskipun mengeksekusi kode yang sama, tidak sinkron kecuali disinkronkan secara eksplisit (misalnya dengan __synchtreads()).
Pedro
1
@Pedro: Benar, tetapi bercabang secara umum merusak kinerja. Untuk kode berkinerja tinggi (kode GPU apa yang bukan?), Hampir penting untuk memperhitungkannya.
jvriesem
21

Masalah yang memiliki intensitas aritmatika tinggi dan pola akses memori biasa biasanya mudah diterapkan pada GPU, dan berkinerja baik pada mereka.

Kesulitan dasar dalam memiliki kode GPU berkinerja tinggi adalah Anda memiliki banyak core, dan Anda ingin semuanya dimanfaatkan sebaik mungkin. Masalah yang memiliki pola akses memori tidak teratur atau tidak memiliki intensitas aritmatika tinggi membuat ini sulit: apakah Anda menghabiskan waktu yang lama untuk mengkomunikasikan hasil atau Anda menghabiskan waktu yang lama untuk mengambil barang dari memori (yang lambat!), Dan tidak cukup waktu menghitung angka. Tentu saja potensi konkurensi dalam kode Anda sangat penting untuk kemampuannya untuk diterapkan dengan baik pada GPU juga.

Reid.Atcheson
sumber
Bisakah Anda menentukan apa yang Anda maksud dengan pola akses memori biasa?
Fomite
1
jawaban maxhutch lebih baik dari saya. Yang saya maksud dengan pola akses reguler adalah memori diakses secara temporal dan spasial secara lokal. Artinya: Anda tidak membuat lompatan besar di sekitar memori berulang kali. Ini juga sesuatu dari paket yang saya perhatikan. Ini juga berarti bahwa pola akses data Anda dapat ditentukan sebelumnya baik oleh kompiler atau oleh Anda programmer sehingga percabangan (pernyataan kondisional dalam kode) diminimalkan.
Reid.Atcheson
15

Ini tidak dimaksudkan sebagai jawaban sendiri tetapi lebih sebagai tambahan untuk jawaban lain oleh maxhutch dan Reid.Atcheson .

Untuk mendapatkan yang terbaik dari GPU, masalah Anda tidak hanya harus sangat (atau masif) paralel, tetapi juga algoritma inti yang akan dieksekusi pada GPU, harus sekecil mungkin. Dalam istilah OpenCL ini sebagian besar disebut sebagai kernel .

Untuk lebih tepatnya, kernel harus masuk ke dalam register setiap unit multiprosesor (atau unit komputasi ) GPU. Ukuran persis register tergantung pada GPU.

Mengingat kernelnya cukup kecil, data mentah dari masalah perlu masuk ke memori lokal GPU (baca: memori lokal (OpenCL) atau memori bersama (CUDA) dari unit komputasi). Kalau tidak, bahkan bandwidth memori GPU yang tinggi tidak cukup cepat untuk membuat elemen pemrosesan sibuk sepanjang waktu.
Biasanya memori ini sekitar 16 hingga 32 KiByte besar .

Torbjörn
sumber
Bukankah memori lokal / bersama dari setiap unit pemrosesan dibagi di antara semua lusinan (?) Thread yang berjalan dalam satu cluster inti? Dalam hal ini, bukankah Anda benar-benar harus menjaga set data kerja Anda secara signifikan lebih kecil untuk mendapatkan kinerja penuh dari GPU?
Dan Neely
Memori lokal / bersama dari unit pemrosesan hanya dapat diakses oleh unit komputasi itu sendiri dan dengan demikian hanya dibagikan oleh elemen-elemen pemrosesan unit komputasi ini. Memori global kartu grafis (biasanya 1GB) dapat diakses oleh semua unit pemrosesan. Bandwidth antara elemen pemrosesan dan memori lokal / bersama sangat cepat (> 1TB / s) tetapi bandwidth ke memori global jauh lebih lambat (~ 100GB / s) dan perlu dibagi di antara semua unit komputasi.
Torbjörn
Saya tidak bertanya tentang memori GPU utama. Saya pikir on die memory hanya dialokasikan pada cluster level inti bukan per core individu. ex untuk nVidia GF100 / 110 gpu; untuk masing-masing dari 16 cluster SM bukan 512 cuda core. Dengan setiap SM yang dirancang untuk menjalankan hingga 32 utas secara paralel memaksimalkan kinerja GPU akan membutuhkan menjaga set kerja dalam kisaran 1kb / utas.
Dan Neely
@ Tororbern Apa yang Anda inginkan adalah membuat semua pipa eksekusi GPU sibuk, GPU mencapai dua cara ini: (1) cara yang paling umum adalah meningkatkan hunian, atau mengatakan secara berbeda, dengan meningkatkan jumlah utas bersamaan (kernel kecil menggunakan lebih sedikit dari sumber daya bersama sehingga Anda dapat memiliki lebih banyak utas aktif); mungkin lebih baik, adalah (2) meningkatkan paralelisme tingkat instruksi di dalam kernel Anda, sehingga Anda dapat memiliki kernel yang lebih besar dengan tingkat hunian yang relatif rendah (sejumlah kecil thread aktif). Lihat bit.ly/Q3KdI0
fcruz
11

Mungkin tambahan yang lebih teknis untuk balasan sebelumnya: CUDA (yaitu Nvidia) GPU dapat digambarkan sebagai satu set prosesor yang bekerja secara otonom di masing-masing 32 utas. Utas di setiap prosesor berfungsi dalam langkah-kunci (pikirkan SIMD dengan vektor panjang 32).

Meskipun cara yang paling menggoda untuk bekerja dengan GPU adalah dengan berpura-pura bahwa semuanya benar-benar berjalan secara terkunci, ini tidak selalu merupakan cara yang paling efisien untuk melakukan sesuatu.

Jika kode Anda tidak memparalelkan dengan baik / otomatis ke ratusan / ribuan utas, Anda mungkin dapat memecahnya menjadi tugas-tugas asinkron individu yang melakukan paralelisasi dengan baik, dan menjalankannya dengan hanya 32 utas yang berjalan dalam langkah-kunci. CUDA menyediakan serangkaian instruksi atom yang memungkinkan untuk mengimplementasikan mutex yang pada gilirannya memungkinkan prosesor untuk melakukan sinkronisasi di antara mereka sendiri dan memproses daftar tugas dalam paradigma thread pool . Kode Anda kemudian akan bekerja banyak dengan cara yang sama seperti pada sistem multi-inti, hanya perlu diingat bahwa setiap inti kemudian memiliki 32 utas sendiri.

Berikut adalah contoh kecil, menggunakan CUDA, tentang cara kerjanya

/* Global index of the next available task, assume this has been set to
   zero before spawning the kernel. */
__device__ int next_task;

/* We will use this value as our mutex variable. Assume it has been set to
   zero before spawning the kernel. */
__device__ int tasks_mutex;

/* Mutex routines using atomic compare-and-set. */
__device__ inline void cuda_mutex_lock ( int *m ) {
    while ( atomicCAS( m , 0 , 1 ) != 0 );
    }
__device__ inline void cuda_mutex_unlock ( int *m ) {
    atomicExch( m , 0 );
    }

__device__ void task_do ( struct task *t ) {

    /* Do whatever needs to be done for the task t using the 32 threads of
       a single warp. */
    }

__global__ void main ( struct task *tasks , int nr_tasks ) {

    __shared__ task_id;

    /* Main task loop... */
    while ( next_task < nr_tasks ) {

        /* The first thread in this block is responsible for picking-up a task. */
        if ( threadIdx.x == 0 ) {

            /* Get a hold of the task mutex. */
            cuda_mutex_lock( &tasks_mutex );

            /* Store the next task in the shared task_id variable so that all
               threads in this warp can see it. */
            task_id = next_task;

            /* Increase the task counter. */
            next_tast += 1;

            /* Make sure those last two writes to local and global memory can
               be seen by everybody. */
            __threadfence();

            /* Unlock the task mutex. */
            cuda_mutex_unlock( &tasks_mutex );

            }

        /* As of here, all threads in this warp are back in sync, so if we
           got a valid task, perform it. */
        if ( task_id < nr_tasks )
            task_do( &tasks[ task_id ] );

        } /* main loop. */

    }

Anda kemudian harus memanggil kernel dengan main<<<N,32>>>(tasks,nr_tasks)untuk memastikan bahwa setiap blok hanya berisi 32 utas dan dengan demikian cocok dalam satu warp. Dalam contoh ini saya juga mengasumsikan, untuk kesederhanaan, bahwa tugas tidak memiliki dependensi (misalnya satu tugas tergantung pada hasil yang lain) atau konflik (misalnya bekerja pada memori global yang sama). Jika ini masalahnya, maka pemilihan tugas menjadi sedikit lebih rumit, tetapi struktur dasarnya sama.

Ini, tentu saja, lebih rumit daripada hanya melakukan segalanya pada satu batch besar sel, tetapi secara signifikan memperluas jenis masalah yang dapat digunakan GPU.

Pedro
sumber
2
Secara teknis ini benar, tetapi paralelisme tinggi diperlukan untuk mendapatkan bandwidth memori tinggi dan ada batasan jumlah panggilan kernel asinkron (saat ini 16). Anda juga banyak perilaku tidak berdokumen terkait dengan penjadwalan dalam rilis saat ini. Saya akan menyarankan agar tidak mengandalkan kernel asinkron untuk kinerja yang penting untuk saat ini ...
Max Hutchinson
2
Apa yang saya uraikan dapat dilakukan semuanya dalam satu panggilan kernel tunggal. Anda dapat membuat N blok yang masing-masing terdiri dari 32 utas, sehingga setiap blok cocok menjadi satu warp. Setiap blok kemudian mendapatkan tugas dari daftar tugas global (akses terkontrol menggunakan atom / mutex) dan menghitungnya menggunakan 32 utas langkah-langkah. Semua ini terjadi dalam satu panggilan kernel. Jika Anda ingin contoh kode, beri tahu saya dan saya akan mempostingnya.
Pedro
4

Satu hal yang tidak dibuat sejauh ini adalah bahwa generasi GPU saat ini tidak melakukannya dengan baik pada perhitungan floating point presisi ganda seperti halnya dengan komputasi presisi tunggal. Jika perhitungan Anda harus dilakukan dalam presisi ganda, maka Anda dapat mengharapkan waktu berjalan meningkat dengan faktor 10 atau lebih dari presisi tunggal.

Brian Borchers
sumber
Saya ingin tidak setuju. Sebagian besar (atau semua) GPU baru memiliki dukungan presisi ganda asli. Hampir setiap GPU melaporkan perhitungan presisi ganda yang berjalan sekitar setengah kecepatan presisi tunggal, kemungkinan karena penggandaan sederhana dari akses memori / bandwidth yang diperlukan.
Godric Seer
1
Walaupun memang benar bahwa kartu Nvidia Tesla terbaru dan terhebat memang menawarkan kinerja puncak ganda presisi yang setengah dari kinerja puncak tunggal presisi, rasionya adalah 8 banding 1 untuk kartu kelas konsumen arsitektur Fermi yang lebih umum.
Brian Borchers
@GodricSeer Rasio 2: 1 dari SP dan DP floating-point sangat sedikit hubungannya dengan bandwidth dan hampir semuanya berkaitan dengan berapa banyak unit perangkat keras yang ada untuk menjalankan operasi ini. Adalah umum untuk menggunakan kembali file register untuk SP dan DP, karenanya unit floating-point dapat mengeksekusi 2x ops SP sebagai ops DP. Ada banyak pengecualian untuk desain ini, misalnya IBM Blue Gene / Q (tidak memiliki logika SP dan karenanya SP beroperasi pada ~ 1,05x DP). Beberapa GPU memiliki rasio selain 2, misalnya 3 dan 5.
Jeff
Sudah empat tahun sejak saya menulis jawaban ini, dan situasi saat ini dengan NVIDIA GPU adalah untuk GeForce dan Quadro, rasio DP / SP sekarang 1/32. NVIDIA's Tesla GPUs memiliki kinerja presisi ganda yang jauh lebih kuat tetapi juga lebih mahal. Di sisi lain, AMD belum melumpuhkan kinerja presisi ganda pada GPU Radeon-nya dengan cara yang sama.
Brian Borchers
4

Dari sudut pandang metaforis, GPU dapat dilihat sebagai orang yang berbaring di atas paku. Orang yang berbaring di atas adalah data dan di dasar setiap kuku ada prosesor, sehingga kuku sebenarnya adalah panah yang menunjuk dari prosesor ke memori. Semua kuku dalam pola teratur, seperti kisi-kisi. Jika tubuh menyebar dengan baik, rasanya enak (kinerja baik), jika tubuh hanya menyentuh beberapa bintik pada kuku, maka rasa sakitnya buruk (kinerja buruk).

Ini dapat diambil sebagai jawaban pelengkap untuk jawaban yang sangat baik di atas.

labotsirc
sumber
4

Pertanyaan lama, tapi saya pikir jawaban ini dari tahun 2014 - terkait dengan metode statistik, tetapi dapat digeneralisasikan untuk siapa pun yang tahu apa itu loop - sangat ilustratif dan informatif.

GT
sumber
2

GPU memiliki latensi I / O yang lama, sehingga banyak utas perlu digunakan untuk memenuhi memori. Untuk membuat warp sibuk, dibutuhkan banyak utas. Jika jalur kode adalah 10 jam dan I / O latency 320 jam, 32 utas akan mendekati jenuh warp. Jika jalur kode adalah 5 jam, maka gandakan utasnya.

Dengan seribu inti, cari ribuan utas untuk sepenuhnya memanfaatkan GPU.

Akses memori adalah dengan jalur cache, biasanya 32 byte. Memuat satu byte memiliki biaya yang sebanding dengan 32 byte. Jadi, satukan penyimpanan untuk meningkatkan lokalitas penggunaan.

Ada banyak register dan RAM lokal untuk masing-masing warp, memungkinkan untuk berbagi tetangga.

Simulasi kedekatan perangkat besar harus dioptimalkan dengan baik.

I / O acak dan threading tunggal adalah kesenangan membunuh ...

pengguna14381
sumber
Ini adalah pertanyaan yang sangat menarik; Saya berdebat dengan diri saya sendiri apakah mungkin (atau sepadan dengan usaha) untuk 'paralel-ise' tugas yang cukup mudah (deteksi tepi dalam gambar udara) ketika setiap tugas mengambil ~ 0,06detik tetapi ada ~ 1,8 juta tugas untuk dilakukan ( per tahun, untuk data bernilai 6 tahun: tugas-tugasnya jelas dapat dipisahkan) ... dengan demikian ~ menghitung waktu komputasi 7,5 hari pada satu inti. Jika setiap perhitungan lebih cepat pada GPU, dan pekerjaan itu dapat diparalelkan 1-per-nGPU skor [n kecil], apakah mungkin waktu kerja turun menjadi ~ 1 jam? Sepertinya tidak mungkin.
GT.
0

Bayangkan sebuah masalah yang bisa diselesaikan oleh banyak kekuatan kasar, seperti Travelling Salesman. Lalu bayangkan Anda punya rak server dengan masing-masing 8 kartu video yang mencolok, dan masing-masing kartu memiliki 3000 core CUDA.

Cukup selesaikan SEMUA rute penjual yang mungkin dan kemudian urutkan berdasarkan waktu / jarak / beberapa metrik. Tentu Anda membuang hampir 100% dari pekerjaan Anda, tetapi brute force kadang-kadang merupakan solusi yang layak.

Criggie
sumber
Saya memiliki akses ke sebuah peternakan kecil yang terdiri dari 4 server selama seminggu, dan dalam lima hari saya melakukan lebih banyak blok Distribution.net daripada 10 tahun sebelumnya.
Criggie
-1

Dari mempelajari banyak ide-ide Teknik, saya akan mengatakan gpu adalah bentuk fokus tugas, manajemen memori, perhitungan berulang.

Banyak rumus mungkin sederhana untuk ditulis tetapi sulit untuk dihitung seperti dalam matematika matriks Anda tidak mendapatkan jawaban tunggal tetapi banyak nilai.

Ini penting dalam komputasi karena seberapa cepat komputer menghitung nilai dan menjalankan rumus karena beberapa rumus tidak dapat berjalan tanpa semua nilai yang dihitung (karenanya memperlambat). Komputer tidak tahu betul bagaimana menjalankan formula atau menghitung nilai yang akan digunakan dalam program ini. Ini terutama memaksa melalui kecepatan cepat dan memecah formula menjadi chuck untuk menghitung, tetapi banyak program saat ini membutuhkan chuck yang dihitung saat ini dan menunggu dalam ques (dan ques ques dan lebih banyak ques ques).

Misalnya dalam game simulasi yang harus dihitung dulu dalam tabrakan kerusakan tabrakan, posisi benda, kecepatan baru? Berapa lama waktu yang dibutuhkan? Bagaimana cara cpu menangani beban ini? Juga, sebagian besar program sangat abstrak yang membutuhkan lebih banyak waktu untuk menangani data dan tidak selalu dirancang untuk multi-threading atau tidak ada cara yang baik dalam program abstrak untuk melakukan ini secara efektif.

Seiring cpu menjadi lebih baik dan orang-orang menjadi ceroboh dalam pemrograman dan kami harus memprogram untuk berbagai jenis komputer juga. GPU dirancang untuk memaksa melalui banyak perhitungan sederhana pada saat yang sama (tidak disebutkan memori (sekunder / ram) dan pendinginan pendingin adalah leher botol utama dalam komputasi). Sebuah cpu mengelola banyak pertanyaan pada saat yang sama atau ditarik ke berbagai arah, mencari tahu apa yang tidak bisa dilakukan. (hei itu hampir manusia)

GPU adalah pekerja kasar pekerjaan yang membosankan. CPU mengelola kekacauan total dan tidak bisa menangani setiap detail.

Jadi apa yang kita pelajari? CPU melakukan detail pekerjaan yang membosankan sekaligus dan CPU adalah mesin multi-tugas yang tidak dapat fokus dengan baik dengan terlalu banyak tugas yang harus dilakukan. (Sepertinya memiliki gangguan perhatian dan autisme pada saat yang sama).

Teknik ada ide, desain, realitas, dan banyak pekerjaan kasar.

Ketika saya pergi, ingatlah untuk memulai dari yang sederhana, mulailah dengan cepat, gagal-cepat, gagal-cepat, dan tidak pernah berhenti mencoba.

Andrew G. Corbi
sumber