Algoritma paralel (GPU) untuk automata seluler asinkron

12

Saya memiliki koleksi model komputasi yang dapat digambarkan sebagai automata seluler asinkron. Model-model ini menyerupai model Ising, tetapi sedikit lebih rumit. Sepertinya model seperti itu akan mendapat manfaat dari dijalankan pada GPU daripada CPU. Sayangnya itu tidak cukup mudah untuk memparalelkan model seperti itu, dan sama sekali tidak jelas bagi saya bagaimana cara melakukannya. Saya sadar bahwa ada literatur tentang masalah ini, tetapi semuanya tampaknya ditujukan pada ilmuwan komputer hardcore yang tertarik pada detail kompleksitas algoritmik, daripada seseorang seperti saya yang hanya menginginkan deskripsi tentang sesuatu yang dapat saya terapkan, dan akibatnya saya merasa agak tidak bisa ditembus.

Untuk kejelasan, saya tidak mencari algoritma yang optimal sebanyak sesuatu yang saya dapat dengan cepat mengimplementasikan di CUDA yang cenderung memberikan peningkatan yang signifikan atas implementasi CPU saya. Waktu pemrogram jauh lebih merupakan faktor pembatas daripada waktu komputer dalam proyek ini.

Saya juga harus mengklarifikasi bahwa otomat seluler asinkron adalah hal yang agak berbeda dari yang sinkron, dan teknik untuk paralelisasi CA sinkron (seperti kehidupan Conway) tidak dapat dengan mudah disesuaikan dengan masalah ini. Perbedaannya adalah bahwa CA sinkron memperbarui setiap sel secara bersamaan pada setiap langkah waktu, sedangkan yang asinkron memperbarui wilayah lokal yang dipilih secara acak pada setiap langkah waktu seperti diuraikan di bawah ini.

Model yang ingin saya paralelkan diimplementasikan pada kisi (biasanya heksagonal) yang terdiri dari ~ 100000 sel (meskipun saya ingin menggunakan lebih banyak), dan algoritma non-paralel untuk menjalankannya terlihat seperti ini:

  1. Pilih pasangan sel tetangga secara acak

  2. Hitung fungsi "energi" berdasarkan lingkungan lokal di sekitar sel-sel iniΔE

  3. Dengan probabilitas yang bergantung pada (dengan parameter β a), dapat menukar status kedua sel atau tidak melakukan apa pun.eβΔEβ

  4. Ulangi langkah-langkah di atas tanpa batas.

Ada juga beberapa komplikasi yang berkaitan dengan kondisi batas, tetapi saya membayangkan ini tidak akan menimbulkan banyak kesulitan untuk paralisis.

Perlu disebutkan bahwa saya tertarik pada dinamika transien dari sistem-sistem ini daripada hanya pada kondisi keseimbangan, jadi saya memerlukan sesuatu yang memiliki dinamika setara dengan yang di atas, daripada hanya sesuatu yang akan mendekati distribusi kesetimbangan yang sama. (Jadi variasi dari algoritma papan catur bukan yang saya cari.)

Kesulitan utama dalam memparalelkan algoritma di atas adalah tumbukan. Karena semua perhitungan hanya bergantung pada wilayah lokal kisi, ada kemungkinan banyak situs kisi diperbarui secara paralel, selama lingkungannya tidak tumpang tindih. Pertanyaannya adalah bagaimana menghindari tumpang tindih tersebut. Saya dapat memikirkan beberapa cara, tetapi saya tidak tahu mana yang terbaik untuk diterapkan. Ini adalah sebagai berikut:

  • Gunakan CPU untuk menghasilkan daftar situs grid acak dan periksa tabrakan. Ketika jumlah situs grid sama dengan jumlah prosesor GPU, atau jika tabrakan terdeteksi, kirim setiap set koordinat ke unit GPU untuk memperbarui situs grid yang sesuai. Ini akan mudah diimplementasikan tetapi mungkin tidak akan mempercepat, karena memeriksa tabrakan pada CPU mungkin tidak akan jauh lebih murah daripada melakukan seluruh pembaruan pada CPU.

  • Bagilah kisi-kisi menjadi beberapa wilayah (satu per unit GPU), dan minta satu unit GPU yang bertanggung jawab untuk secara acak memilih dan memperbarui sel-sel kisi di dalam wilayahnya. Tetapi ada banyak masalah dengan gagasan ini yang saya tidak tahu bagaimana menyelesaikannya, yang paling jelas adalah apa yang seharusnya terjadi ketika sebuah unit memilih lingkungan yang tumpang tindih dengan tepi wilayahnya.

  • Perkirakan sistem sebagai berikut: biarkan waktu melanjutkan dalam langkah-langkah tersendiri. Bagilah kisi menjadi berbedamengatur wilayah pada setiap langkah waktu sesuai dengan beberapa skema yang telah ditentukan sebelumnya, dan meminta setiap unit GPU secara acak memilih dan memperbarui sepasang sel kisi yang lingkungannya tidak tumpang tindih dengan batas wilayah. Karena batas berubah setiap langkah, kendala ini mungkin tidak terlalu mempengaruhi dinamika, asalkan daerah tersebut relatif besar. Ini tampaknya mudah diimplementasikan dan cenderung cepat, tetapi saya tidak tahu seberapa baik perkiraannya terhadap dinamika, atau apa skema terbaik untuk memilih batas wilayah pada setiap langkah waktu. Saya menemukan beberapa referensi untuk "automata seluler blok-sinkron", yang mungkin atau mungkin tidak sama dengan ide ini. (Saya tidak tahu karena sepertinya semua deskripsi metode ini dalam bahasa Rusia atau dalam sumber yang tidak dapat saya akses.)

Pertanyaan spesifik saya adalah sebagai berikut:

  • Adakah salah satu dari algoritma di atas merupakan cara yang masuk akal untuk mendekati parallelisation GPU dari model CA asinkron?

  • Apakah ada cara yang lebih baik?

  • Apakah ada kode perpustakaan yang ada untuk jenis masalah ini?

  • Di mana saya dapat menemukan deskripsi bahasa Inggris yang jelas tentang metode "blok-sinkron"?

Kemajuan

Saya yakin saya telah menemukan cara untuk memparalelkan CA asinkron yang mungkin cocok. Algoritme yang diuraikan di bawah ini adalah untuk CA asinkron normal yang memperbarui hanya satu sel pada satu waktu, daripada sepasang sel tetangga seperti saya. Ada beberapa masalah dengan menggeneralisasikannya ke kasus spesifik saya, tetapi saya rasa saya memiliki ide bagaimana menyelesaikannya. Namun, saya tidak yakin berapa banyak manfaat kecepatan itu akan memberikan, karena alasan yang dibahas di bawah ini.

Idenya adalah untuk mengganti CA asinkron (selanjutnya ACA) dengan CA sinkron stokastik (SCA) yang berperilaku setara. Untuk melakukan ini pertama-tama kita bayangkan bahwa ACA adalah proses Poisson. Yaitu, waktu berlangsung terus menerus, dan setiap sel sebagai probabilitas konstan per unit waktu untuk melakukan fungsi pembaruannya, terlepas dari sel lainnya.

Xijtijtij(0)Exp(λ)λ adalah parameter yang nilainya dapat dipilih secara sewenang-wenang.)

Pada setiap langkah waktu logis, sel-sel SCA diperbarui sebagai berikut:

  • k,li,jtkl<tij

  • XijXklΔtExp(λ)tijtij+Δt

Saya percaya ini menjamin bahwa sel-sel akan diperbarui dalam urutan yang dapat "diterjemahkan" agar sesuai dengan ACA asli, sambil menghindari tabrakan dan memungkinkan beberapa sel diperbarui secara paralel. Namun, karena poin pertama di atas, itu berarti bahwa sebagian besar prosesor GPU sebagian besar akan menganggur pada setiap langkah waktu SCA, yang kurang dari ideal.

Saya perlu lebih memikirkan apakah kinerja algoritma ini dapat ditingkatkan, dan bagaimana memperluas algoritma ini untuk menangani kasus di mana beberapa sel diperbarui secara bersamaan dalam ACA. Namun, itu terlihat menjanjikan sehingga saya pikir saya akan menggambarkannya di sini kalau-kalau ada yang (a) tahu apa pun yang serupa dalam literatur, atau (b) dapat menawarkan wawasan apa pun tentang masalah yang tersisa ini.

Nathaniel
sumber
Mungkin Anda dapat merumuskan masalah Anda dalam pendekatan berbasis stensil. Banyak perangkat lunak yang ada untuk masalah berbasis stensil. Anda dapat melihat di: libgeodecomp.org/gallery.html , Conway's Game of Life. Ini mungkin memiliki beberapa kesamaan.
vanCompute
@vanCompute yang terlihat seperti alat yang fantastis, tetapi dari penyelidikan awal saya (agak sepintas), sepertinya paradigma kode stensil secara inheren sinkron, jadi mungkin tidak cocok dengan apa yang saya coba lakukan. Namun, saya akan memeriksanya lebih jauh.
Nathaniel
Bisakah Anda memberikan beberapa detail lebih lanjut tentang bagaimana Anda akan memparalelkan ini menggunakan SIMT? Apakah Anda menggunakan satu utas per pasangan? Atau dapatkah pekerjaan yang terlibat dengan memperbarui satu pasangan tersebar di lebih dari 32 utas?
Pedro
@Pedro pekerjaan yang terlibat dalam memperbarui satu pasangan cukup kecil (pada dasarnya hanya menjumlahkan lingkungan, ditambah satu iterasi dari generator nomor acak dan satu exp()) jadi saya tidak akan berpikir itu masuk akal untuk menyebarkannya ke beberapa utas. Saya pikir lebih baik (dan lebih mudah bagi saya) untuk mencoba dan memperbarui beberapa pasangan secara paralel, dengan satu pasang per utas.
Nathaniel
Oke, dan bagaimana Anda mendefinisikan tumpang tindih antara untuk memasangkan pembaruan? Jika pasangan itu tumpang tindih, atau jika lingkungan mereka tumpang tindih?
Pedro

Jawaban:

4

Saya akan menggunakan opsi pertama dan akan menggunakan AC sinkron sebelum (menggunakan GPU), untuk mendeteksi tabrakan, menjalankan langkah AC heksagonal yang aturannya adalah nilai sel pusat = Jumlah (tetangga), CA ini harus memiliki tujuh status harus dimulai dengan sel yang dipilih secara acak, dan statusnya diverifikasi sebelum menjalankan aturan pembaruan untuk setiap GPU.

Sampel 1. Nilai sel tetangga dibagikan

0 0 0 0 0 0 0

  0 0 1 0 0 0

0 0 0 0 0 0 0

  0 0 0 1 0 0

0 0 0 0 0 0 0

langkah CA yang aturannya adalah sel pusat heksagonal = Jumlah (tetangga)

0 0 1 1 0 0 0

  0 1 1 1 0 0

0 0 1 2 1 0 0

  0 0 1 1 1 0

0 0 0 1 1 0 0

Contoh 2. Nilai sel untuk memperbarui diperhitungkan sebagai tetangga di sisi lain

0 0 0 0 0 0 0

  0 0 1 0 0 0

0 0 0 1 0 0 0

  0 0 0 0 0 0

0 0 0 0 0 0 0

Setelah iterasi

0 0 1 1 0 0 0

  0 1 2 2 0 0

0 0 2 2 1 0 0

  0 0 1 1 0 0

0 0 0 0 0 0 0

Contoh 3. Tidak ada hubungan

  0 0 0 0 0 0

0 0 1 0 0 0 0

  0 0 0 0 0 0

0 0 0 0 0 0 0

  0 0 0 1 0 0

0 0 0 0 0 0 0

Setelah iterasi

  0 1 1 0 0 0

0 1 1 1 0 0 0

  0 1 1 0 0 0

0 0 0 1 1 0 0

  0 0 1 1 1 0

0 0 0 1 1 0 0

jlopez1967
sumber
O(n)n
Saya pikir ada banyak yang bisa diparalelkan. Pemrosesan tabrakan yang sepenuhnya dipengaruhi GPU adalah langkah dalam AC sinkron seperti yang ditunjukkan pada tautan yang diposting di atas. untuk verifikasi akan menggunakan aturan lokal jika Sum (tetangga) = 8 TANPA tabrakan, Jumlah (tetangga)> 8 Tabrakan, itu akan diverifikasi sebelum menjalankan perubahan aturan pembaruan Anda jika tidak ada status sel tabrakan, karena keduanya harus diletakkan di dekat poin yang akan dievaluasi jika tidak dekat adalah milik sel lain.
jlopez1967
Saya mengerti itu, tetapi masalahnya adalah, apa yang Anda lakukan ketika Anda mendeteksi tabrakan? Seperti yang saya jelaskan di atas, algoritma CA Anda hanyalah langkah pertama dalam mendeteksi tabrakan. Langkah kedua adalah mencari sel untuk kotak dengan keadaan> = 2, dan ini tidak sepele.
Nathaniel
mis. Bayangkan bahwa kita ingin mendeteksi sel tumbukan (5.7), pada automata seluler dan jumlah yang dieksekusi (tetangga sel (5,7)) dan jika nilainya 8 dan jika tidak ada tumbukan lebih besar dari 8 tumbukan ini harus dalam fungsi yang mengevaluasi setiap sel untuk menentukan keadaan sel selanjutnya dalam automata seluler asinkron. Deteksi tabrakan untuk setiap sel adalah aturan lokal yang hanya melibatkan sel
jlopez1967
Ya, tetapi pertanyaan yang kami harus dapat jawab untuk memparalelkan CA asinkron bukan "apakah ada tabrakan dalam sel (5,7)" tetapi "apakah ada tabrakan di suatu tempat di grid, dan jika demikian di mana ada Itu?" Itu tidak bisa dijawab tanpa iterasi di grid.
Nathaniel
1

Mengikuti jawaban Anda atas pertanyaan saya di komentar di atas, saya sarankan Anda mencoba pendekatan berbasis kunci di mana setiap utas mencoba mengunci lingkungan yang akan diperbarui sebelum menghitung pembaruan aktual.

Anda dapat melakukan ini menggunakan operasi atom yang diatur dalam CUDA, dan array intberisi kunci untuk setiap sel, misalnya lock. Setiap utas kemudian melakukan hal berikut:

ci, cj = choose a pair at random.

int locked = 0;

/* Try to lock the cell ci. */
if ( atomicCAS( &lock[ci] , 0 , 1 ) == 0 ) {

    /* Try to lock the cell cj. */
    if ( atomicCAS( &lock[cj] , 0 , 1 ) == 0 ) {

        /* Now try to lock all the neigbourhood cells. */
        for ( cn = indices of all neighbours )
            if ( atomicCAS( &lock[cn] , 0 , 1 ) != 0 )
                break;

        /* If we hit a break above, we have to unroll all the locks. */
        if ( cn < number of neighbours ) {
            lock[ci] = 0;
            lock[cj] = 0;
            for ( int i = 0 ; i < cn ; i++ )
                lock[i] = 0;
            }

        /* Otherwise, we've successfully locked-down the neighbourhood. */
        else
            locked = 1;

        }

    /* Otherwise, back off. */
    else
        lock[ci] = 0;
    }

/* If we got everything locked-down... */
if ( locked ) {

    do whatever needs to be done...

    /* Release all the locks. */
    lock[ci] = 0;
    lock[cj] = 0;
    for ( int i = 0 ; i < cn ; i++ )
        lock[i] = 0;

    }

Perhatikan bahwa pendekatan ini mungkin bukan yang paling optimal, tetapi dapat memberikan titik awal yang menarik. Jika ada banyak tabrakan antar thread, yaitu satu atau lebih per 32 thread (seperti dalam satu collision per warp), maka akan ada sedikit pengalihan cabang. Juga, operasi atom bisa agak lambat, tetapi karena Anda hanya melakukan operasi perbandingan dan pertukaran, seharusnya skala ok.

Mengunci overhead mungkin terlihat menakutkan, tetapi sebenarnya hanya beberapa tugas dan cabang, tidak lebih.

Perhatikan juga bahwa saya menjadi cepat dan longgar dengan notasi di loop iatas tetangga.

Tambahan: Saya cukup angkuh untuk berasumsi bahwa Anda bisa mundur ketika berpasangan. Jika ini bukan masalahnya, maka Anda dapat membungkus semuanya pada baris kedua di while-loop dan menambahkannya breakdi akhir pernyataan-akhir if.

Semua utas harus menunggu sampai yang terakhir selesai, tetapi jika tabrakan jarang terjadi, Anda harus bisa lolos.

Addendum 2: Do tidak tergoda untuk menambahkan panggilan ke __syncthreads()mana saja di kode ini, terutama itu versi perulangan dijelaskan dalam addendum sebelumnya! Asynchronicity sangat penting dalam menghindari tabrakan berulang dalam kasus yang terakhir.

Pedro
sumber
Terima kasih, ini terlihat cukup bagus. Mungkin lebih baik daripada ide rumit yang saya pertimbangkan, dan jauh lebih mudah untuk diterapkan. Saya bisa membuat tabrakan langka dengan menggunakan kotak yang cukup besar, yang mungkin baik-baik saja. Jika metode just-back-off ternyata secara signifikan lebih cepat saya dapat menggunakannya untuk menyelidiki parameter secara informal, dan beralih ke metode menunggu-untuk-semua-orang-untuk-menyelesaikan ketika saya perlu menghasilkan hasil resmi. Saya akan mencobanya sebentar lagi.
Nathaniel
1

Saya adalah pengembang utama LibGeoDecomp. Meskipun saya setuju dengan vanCompute bahwa Anda dapat meniru ACA Anda dengan CA, Anda benar bahwa ini tidak akan sangat efisien, karena hanya beberapa sel dalam setiap langkah yang dimaksudkan untuk diperbarui. Ini memang aplikasi yang sangat menarik - dan menyenangkan untuk bermain-main!

Saya menyarankan Anda untuk menggabungkan solusi yang diusulkan oleh jlopez1967 dan Pedro: Algoritma Pedro menangkap paralelisme dengan baik, tetapi kunci-kunci atom itu sangat lambat. Solusi jlopez1967 adalah elegan ketika datang untuk mendeteksi tabrakan, tetapi memeriksa semua nsel, ketika hanya sebagian kecil (saya akan mulai sekarang berasumsi bahwa ada beberapa parameter kyang menunjukkan jumlah sel yang akan diperbarui secara bersamaan) aktif, jelas merupakan penghalang.

__global__ void markPoints(Cell *grid, int gridWidth, int *posX, int *posY)
{
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    int x, y;
    generateRandomCoord(&x, &y);
    posX[id] = x;
    posY[id] = y;
    grid[y * gridWidth + x].flag = 1;
}

__global__ void checkPoints(Cell *grid, int gridWidth, int *posX, int *posY, bool *active)
{
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    int x = posX[id];
    int y = posY[id];
    int markedNeighbors = 
        grid[(y - 1) * gridWidth + x + 0].flag +
        grid[(y - 1) * gridWidth + x + 1].flag +
        grid[(y + 0) * gridWidth + x - 1].flag +
        grid[(y + 0) * gridWidth + x + 1].flag +
        grid[(y + 1) * gridWidth + x + 0].flag +
        grid[(y + 1) * gridWidth + x + 1].flag;
    active[id] = (markedNeighbors > 0);
}


__global__ void update(Cell *grid, int gridWidth, int *posX, int *posY, bool *active)
{
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    int x = posX[id];
    int y = posY[id];
    grid[y * gridWidth + x].flag = 0;
    if (active[id]) {
        // do your fancy stuff here
    }
}

int main() 
{
  // alloc grid here, update up to k cells simultaneously
  int n = 1024 * 1024;
  int k = 1234;
  for (;;) {
      markPoints<<<gridDim,blockDim>>>(grid, gridWidth, posX, posY);
      checkPoints<<<gridDim,blockDim>>>(grid, gridWidth, posX, posY, active);
      update<<<gridDim,blockDim>>>(grid, gridWidth, posX, posY, active);
  }
}

Dengan tidak adanya sinkronisasi global yang baik pada GPU, Anda perlu memanggil beberapa kernel untuk fase yang berbeda. Di Kepler Nvidia, Anda bahkan dapat memindahkan loop utama ke GPU, tapi saya tidak berharap itu akan mendapatkan banyak.

Algoritma mencapai tingkat paralelisme (dapat dikonfigurasi). Saya kira, pertanyaan menarik adalah apakah tumbukan akan mempengaruhi distribusi acak Anda ketika Anda meningkat k.

Gentryx
sumber
0

Saya sarankan kepada Anda agar Anda melihat tautan ini http://www.wolfram.com/training/courses/hpc021.html tentu saja sekitar 14:15 menit ke dalam video, tentu saja, pelatihan Mathematica di mana mereka membuat implementasi automata seluler menggunakan CUDA , dari sana dan Anda dapat memodifikasinya.

Juan Carlos Lopez
sumber
Sayangnya itu adalah CA sinkron, yang merupakan jenis binatang yang agak berbeda dari yang asinkron yang saya hadapi. Dalam CA sinkron, setiap sel diperbarui secara bersamaan, dan ini mudah untuk diparalelkan pada GPU, tetapi dalam CA asinkron satu sel yang dipilih secara acak diperbarui setiap langkah waktu (sebenarnya dalam kasus saya ini adalah dua sel tetangga), dan ini membuat Paralelisasinya jauh lebih sulit. Masalah yang diuraikan dalam pertanyaan saya khusus untuk memerlukan fungsi pembaruan asinkron.
Nathaniel