
Belum lama ini saya membaca novel fiksi ilmiah
"The Three-Body Problem" oleh Liu Qixin . Di dalamnya, beberapa alien memiliki masalah - mereka tidak tahu bagaimana, dengan akurasi yang cukup bagi mereka, untuk menghitung lintasan planet asal mereka. Tidak seperti kita, mereka hidup dalam sistem bintang tiga, dan "cuaca" di planet ini sangat bergantung pada lokasi relatif mereka - dari panas pembakaran hingga dingin es. Dan saya memutuskan untuk memeriksa apakah kita dapat memecahkan masalah seperti itu.
Fisika dari fenomena tersebut
Untuk memahami masalahnya, perlu berurusan dengan fisika dari fenomena tersebut. Dalam kerangka teori klasik, gaya tarik menarik dua benda ditentukan oleh hukum
Newton :
v e c F ( v e c r 1 , v e c r 2 ) = - G m 1 m 2 f r a c v e c r 1 - v e c r 2 | v e c r 1 - v e c r 2 | 3 ,
dimana
vecr1, vecr2 - posisi benda di ruang angkasa,
m1,m2 - massa tubuh,
G - Konstanta gravitasi.
Dalam sistem
N Tubuh pada masing-masing dari mereka akan dipengaruhi oleh gaya tarik dari yang lain, yang dinyatakan oleh persamaan:
vecFn=−G sumk neqnmnmk frac vecrn− vecrk| vecrn− vecrk|3.
Menggunakan
hukum kedua Newton, kami menulis akselerasi untuk setiap partikel:
vecan= vecFn/mn=−G sumk neqnmk frac vecrn− vecrk| vecrn− vecrk|3.
Mengingat bahwa akselerasi adalah turunan kedua dari koordinat, kita memperoleh persamaan diferensial parsial orde kedua, yang harus dipecahkan untuk mendapatkan lintasan masing-masing badan:
frac partial2 vecrn partialt2=fn=−G sumk neqnmk frac vecrn− vecrk| vecrn− vecrk|3.
Penting untuk dicatat di sini bahwa kompleksitas fungsi komputasi
fn sama dengan
O(N2) dan meningkat secara signifikan dengan meningkatnya jumlah badan yang berinteraksi.
Matematika
Metode pertama dan paling sederhana untuk menyelesaikan persamaan diferensial adalah
metode Euler , yang dirancang untuk menyelesaikan persamaan bentuk:
fracdydx=f(x,y).
Setelah transisi ke wilayah diskrit, kami memperoleh:
yi=yi−1+hf(xi−1,yi−1), quadi=1,2,3, dots,m,
dimana
h Apakah langkah integrasi, dan
m - jumlah langkah integrasi. Jadi, jika kita perlu menghitung posisi tubuh pada suatu waktu
T maka kita harus melakukannya
m=T/h langkah integrasi. Di sini masalah pertama terlihat - jika
T besar, maka kita perlu mengambil sejumlah besar langkah integrasi.
Untuk menerapkan metode Euler ke masalah kita, itu harus direduksi menjadi sistem orde pertama. Untuk melakukan ini, kami memperkenalkan variabel tambahan - kecepatan partikel:
frac partial vecvn partialt=fn, frac partial vecrn partialt= vecvn.
Masalah kedua dalam memecahkan sistem persamaan diferensial adalah keakuratan solusi dan kontrolnya. Akurasi dapat ditingkatkan dengan dua cara: dengan mengurangi langkah integrasi dan memilih metode dengan tingkat akurasi yang lebih tinggi. Kedua metode mengarah pada peningkatan kompleksitas komputasi, tetapi dengan cara yang berbeda. Misalnya, Anda dapat menggunakan
metode Runge-Kutta orde empat klasik, membutuhkan empat perhitungan fungsi
fn di setiap langkah, tetapi memiliki urutan akurasi
O(h4) (Sebagai perbandingan, metode Euler memiliki urutan akurasi
O(h) dan membutuhkan satu perhitungan
fn ) Keakuratan solusi juga dapat dikontrol dalam beberapa cara: membandingkan dengan solusi analitis, menyelesaikan dengan menggunakan metode yang berbeda atau dengan langkah yang berbeda dan membandingkan hasilnya, mengontrol parameter dan batasan pihak ketiga yang harus dipenuhi oleh solusi.
Juga, masing-masing metode ini memiliki kekurangan. Keputusan analitis mungkin tidak ada, atau, bahkan dalam kebanyakan kasus, sama sekali tidak ada. Misalnya, untuk tugas kita
N solusi analitik tel hanya untuk
N=2 , tetapi bahkan ini sudah cukup untuk menguji keakuratan metode. Memecahkan masalah dengan dua metode atau dengan langkah-langkah berbeda meningkatkan waktu perhitungan, tetapi pendekatan ini dapat diterapkan untuk hampir semua tugas. Tidak setiap tugas memiliki keterbatasan, tetapi untuk kita, mereka memiliki: pada setiap langkah integrasi kita dapat mengontrol implementasi
undang-undang konservasi . Pendekatan ini juga meningkatkan kompleksitas perhitungan, tetapi ada banyak pilihan, menghitung jumlah momentum atau momen sudut dari semua partikel memiliki kompleksitas urutan.
O(N) , sementara menghitung energi total suatu sistem memiliki kompleksitas urutan
O(N2)Catatan tentang penghitungan energi totalDalam kasus kami, energi total sistem terdiri dari dua bagian - energi kinetik dan potensial. Energi kinetik terdiri dari jumlah energi kinetik semua tubuh. Untuk menghitung energi potensial, kita perlu menambahkan energi potensial dari masing-masing partikel dalam medan gravitasi partikel yang tersisa, jadi kita perlu menambahkan
O(N2) ketentuan Kesulitannya adalah bahwa semua istilah memiliki urutan yang sangat berbeda, dan bahkan dengan perhitungan presisi ganda tidak mungkin untuk menghitung nilai ini dengan akurasi yang cukup untuk perbandingan pada langkah yang berbeda. Untuk mengatasi masalah ini, perlu untuk menerapkan penjumlahan sesuai dengan
algoritma Kahan .
Gambar 1. Contoh lintasan elips.
Pertimbangkan kasus sederhana dari satelit yang bergerak dalam orbit elips di sekitar Bumi. Ketika satelit mendekati Bumi, kecepatannya meningkat, dan ketika bergerak menjauh dari Bumi, ia menurun, karenanya, kemungkinan muncul untuk mengurangi langkah integrasi tepat waktu di bagian hijau orbit, dan meningkatkannya dalam warna merah dan biru tanpa mengubah keakuratan solusi. Mari kita coba membandingkan dengan lebih detail.
Tabel 1. Metode yang diselidiki untuk menyelesaikan persamaan diferensialUntuk memilih metode terbaik untuk tugas kami, kami akan membandingkan beberapa metode yang diketahui. Untuk melakukan ini, kami mensimulasikan tabrakan dua sistem tubuh
N=512 dan mengukur perubahan relatif dalam
total energi ,
momentum dan
momennya pada akhir simulasi (waktu simulasi maksimum
Tmaks=2,5 ) Dalam hal ini, kami akan memvariasikan langkah dan parameter metode integrasi dan mengukur jumlah panggilan fungsi
fn , masing-masing, metode-metode yang dengan jumlah panggilan yang lebih kecil akan menyebabkan lebih sedikit kerugian, kami akan anggap lebih dapat diterima.

| 
|
a) | b) |
Gambar 2. Perubahan relatif dalam energi a), momentum b), pada akhir simulasi sistem N=512 tubuh dengan berbagai metode tergantung pada jumlah perhitungan fungsi fn presisi ganda |
Dari grafik Gambar 2 dapat dilihat bahwa rasio terbaik dari jumlah perhitungan fungsi
fn dan perubahan relatif dalam energi sistem tubuh dalam metode Adams dan Dorman-Prince orde kelima. Terlihat juga bahwa untuk semua metode dengan peningkatan jumlah perhitungan
fn perubahan relatif dalam momentum sistem meningkat. Untuk perubahan relatif dalam energi, ini juga terlihat, tetapi hanya untuk beberapa metode yang dapat mencapai ambang batas
dE/E0<10−12 . Efek ini tidak lagi dapat dikaitkan dengan kesalahan metode, tetapi karena kesalahan perhitungan, dan peningkatan akurasi lebih lanjut hanya dimungkinkan bersama-sama dengan peningkatan akurasi perhitungan dengan floating point.

| 
|
a) | b) |
Gambar. 3. Perubahan relatif dalam energi a), momentum b), pada akhir simulasi sistem N=512 tubuh dengan berbagai metode tergantung pada jumlah perhitungan fungsi fn dengan presisi empat kali lipat (__float128) |
Gambar 3a dan 3b menunjukkan bahwa penggunaan perhitungan dengan akurasi empat kali lipat dapat mengurangi kehilangan energi relatif hingga
10−23 , tetapi Anda perlu memahami bahwa waktu komputasi ditingkatkan dua urutan besarnya dibandingkan dengan presisi ganda. Seperti dalam kasus perhitungan presisi ganda, rasio akurasi terbaik dan jumlah perhitungan fungsi
fn memiliki metode urutan kelima Adams dan Dorman-Prince.
Metode Dorman-Prince dan Werner termasuk dalam kelas
metode bersarang dan secara bersamaan dapat menghitung dua solusi dengan urutan akurasi tinggi dan rendah (untuk metode Dorman-Prince, pesanan 5 dan 4, dan untuk metode Werner, pesanan 6 dan 5). Jika kedua solusi ini sangat berbeda, maka kita dapat memecah langkah integrasi saat ini menjadi yang lebih kecil. Itu memungkinkan kami untuk secara dinamis mengubah langkah integrasi dan menguranginya hanya di area-area di mana diperlukan.
Mari kita bandingkan metode Dorman-Prince, Werner, dan Adams urutan kelima secara lebih rinci, selama interval simulasi yang lebih lama dari sistem kami (
Tmaks=300 )

|
a) Perubahan relatif dalam energi, momentum dan momennya dalam proses pemodelan
|

|
b) Jumlah perhitungan fungsi fn pada interval simulasi DeltaT=0,3
|
Gambar 4. Ketergantungan perubahan relatif dalam parameter fisik sistem dan jumlah perhitungan fungsi fn dari waktu ke waktu pemodelan dengan metode variabel-langkah Dorman-Prince h=10−5 dots10−9
|

|
a) Perubahan relatif dalam energi, momentum dan momennya dalam proses pemodelan
|

|
b) Jumlah perhitungan fungsi fn pada interval simulasi DeltaT=0,3
|
Gambar. 5. Ketergantungan perubahan relatif dalam parameter fisik sistem dan jumlah perhitungan fungsi fn versus pemodelan waktu dengan metode Werner dengan langkah variabel h=10−5 dots10−9
|

|
Gambar. 6. Perubahan relatif dalam energi, momentum dan momennya dalam proses pemodelan dengan metode Adams-Bashfort orde lima dengan langkah h=10−6
|

|
Gambar 7. Ketergantungan jumlah perhitungan fungsi fn untuk metode kelima, metode Adams, Dorman-Prince dan Werner dari waktu simulasi
|
Dapat dilihat bahwa pada tahap awal evolusi sistem kami (
T<$2 ) ketiga metode tersebut menunjukkan karakteristik yang serupa, tetapi pada tahap selanjutnya peristiwa tertentu terjadi dalam sistem, akibatnya kesalahan dalam parameter utama sistem (energi total, momentum dan momentumnya) melonjak tajam. Tetapi metode Dorman-Prince dan Werner mengatasi perubahan ini jauh lebih baik karena kemampuan untuk mengurangi langkah integrasi di bagian "kompleks", sebagai akibatnya jumlah perhitungan fungsi meningkat
fn seperti yang terlihat pada Gambar 4b dan 5b, tetapi jumlah total perhitungan
fn metode bersarang lebih kecil dari metode Adams-Bashfort, seperti yang dapat dilihat pada Gambar 7.
Saya bertanya-tanya apa yang terjadi pada sistem pada saat-saat ini
|
Video 1. Memodelkan sistem dengan 512 badan. Metode Dorman-Prince. Langkah dinamis h=10−5 dots10−9
|
Video tersebut menunjukkan hal itu hingga ke titik waktu
T=25 gerakannya relatif tenang, dan setelah tabrakan pusat "galaksi" terjadi, yang mengarah ke perubahan tajam dalam lintasan dan peningkatan kuat dalam kecepatan beberapa partikel. Selain itu, untuk menjaga keakuratan solusi, perlu untuk mengurangi langkah integrasi. Metode bersarang dapat melakukan ini secara otomatis, grafik menunjukkan bahwa di beberapa bagian evolusi sistem, langkah integrasi berkurang hampir dua urutan besarnya dengan
10−5 sebelumnya
h=10−7 . Saat menggunakan metode Adams dan langkah seperti itu pada seluruh interval evolusi sistem, kami tidak akan menunggu solusi.
Ringkasan
Untuk solusinya, lebih baik menggunakan metode bersarang yang memungkinkan Anda mengontrol langkah integrasi secara dinamis, dan menguranginya hanya pada bagian "kompleks" dari lintasan.
Jangan mengejar metode urutan tertinggi. Bahkan ketika menggunakan tipe data 'ganda', mereka tidak mencapai kemampuan potensial mereka, menggunakan tipe data dengan akurasi yang lebih besar sangat meningkatkan waktu yang dibutuhkan untuk menyelesaikan masalah.
Implementasi CPU
Sekarang pilihan metode untuk menyelesaikan persamaan didefinisikan, mari kita coba mencari tahu perhitungan gaya interaksi untuk setiap partikel. Kami mendapatkan siklus ganda untuk semua partikel:
Kode implementasi 'sederhana'for(size_t body1 = 0; body1 < count; ++body1) { const nbvertex_t v1(rx[ body1 ], ry[ body1 ], rz[ body1 ]); nbvertex_t total_force; for(size_t body2 = 0; body2 != count; ++body2) { if(body1 == body2) { continue; } const nbvertex_t v2(rx[ body2 ], ry[ body2 ], rz[ body2 ]); const nbvertex_t force(m_data->force(v1, v2, mass[body1], mass[body2])); total_force += force; } frx[body1] = vx[body1]; fry[body1] = vy[body1]; frz[body1] = vz[body1]; fvx[body1] = total_force.x / mass[body1]; fvy[body1] = total_force.y / mass[body1]; fvz[body1] = total_force.z / mass[body1]; }
Gaya gravitasi untuk masing-masing tubuh dihitung secara independen, dan untuk menggunakan semua inti prosesor, cukup untuk menulis arahan OpenMP sebelum siklus pertama:
Sepotong kode dari implementasi 'openmp' #pragma omp parallel for for(size_t body1 = 0; body1 < count; ++body1)
Karena Karena setiap tubuh berinteraksi dengan masing-masing, untuk mengurangi jumlah interaksi prosesor dengan RAM dan meningkatkan penggunaan cache, kami memiliki kemampuan untuk memuat bagian dari data ke dalam cache dan menggunakannya berulang kali:
Kode implementasi 'openmp + block' #pragma omp parallel for for(size_t n1 = 0; n1 < count; n1 += BLOCK_SIZE) { nbcoord_t x1[BLOCK_SIZE]; nbcoord_t y1[BLOCK_SIZE]; nbcoord_t z1[BLOCK_SIZE]; nbcoord_t m1[BLOCK_SIZE]; nbvertex_t total_force[BLOCK_SIZE]; for(size_t b1 = 0; b1 != BLOCK_SIZE; ++b1) { size_t local_n1 = b1 + n1; x1[b1] = rx[local_n1]; y1[b1] = ry[local_n1]; z1[b1] = rz[local_n1]; m1[b1] = mass[local_n1]; } for(size_t n2 = 0; n2 < count; n2 += BLOCK_SIZE) { nbcoord_t x2[BLOCK_SIZE]; nbcoord_t y2[BLOCK_SIZE]; nbcoord_t z2[BLOCK_SIZE]; nbcoord_t m2[BLOCK_SIZE]; for(size_t b2 = 0; b2 != BLOCK_SIZE; ++b2) { size_t local_n2 = b2 + n2; x2[b2] = rx[local_n2]; y2[b2] = ry[local_n2]; z2[b2] = rz[local_n2]; m2[b2] = mass[n2 + b2]; } for(size_t b1 = 0; b1 != BLOCK_SIZE; ++b1) { const nbvertex_t v1(x1[ b1 ], y1[ b1 ], z1[ b1 ]); for(size_t b2 = 0; b2 != BLOCK_SIZE; ++b2) { const nbvertex_t v2(x2[ b2 ], y2[ b2 ], z2[ b2 ]); const nbvertex_t force(m_data->force(v1, v2, m1[b1], m2[b2])); total_force[b1] += force; } } } for(size_t b1 = 0; b1 != BLOCK_SIZE; ++b1) { size_t local_n1 = b1 + n1; frx[local_n1] = vx[local_n1]; fry[local_n1] = vy[local_n1]; frz[local_n1] = vz[local_n1]; fvx[local_n1] = total_force[b1].x / m1[b1]; fvy[local_n1] = total_force[b1].y / m1[b1]; fvz[local_n1] = total_force[b1].z / m1[b1]; } }
Optimalisasi lebih lanjut terdiri dari mengeluarkan fungsi menghitung gaya dalam siklus utama dan menghilangkan pembagian dan perkalian dengan massa tubuh m1 [b1]. Selain fakta bahwa kami telah sedikit mengurangi perhitungan, kompiler akan dapat menerapkan instruksi vektor prosesor SSE dan AVX pada siklus yang diperluas.
Kode implementasi 'openmp + block + optimization' #pragma omp parallel for for(size_t n1 = 0; n1 < count; n1 += BLOCK_SIZE) { nbcoord_t x1[BLOCK_SIZE]; nbcoord_t y1[BLOCK_SIZE]; nbcoord_t z1[BLOCK_SIZE]; nbcoord_t total_force_x[BLOCK_SIZE]; nbcoord_t total_force_y[BLOCK_SIZE]; nbcoord_t total_force_z[BLOCK_SIZE]; for(size_t b1 = 0; b1 != BLOCK_SIZE; ++b1) { size_t local_n1 = b1 + n1; x1[b1] = rx[local_n1]; y1[b1] = ry[local_n1]; z1[b1] = rz[local_n1]; total_force_x[b1] = 0; total_force_y[b1] = 0; total_force_z[b1] = 0; } for(size_t n2 = 0; n2 < count; n2 += BLOCK_SIZE) { nbcoord_t x2[BLOCK_SIZE]; nbcoord_t y2[BLOCK_SIZE]; nbcoord_t z2[BLOCK_SIZE]; nbcoord_t m2[BLOCK_SIZE]; for(size_t b2 = 0; b2 != BLOCK_SIZE; ++b2) { size_t local_n2 = b2 + n2; x2[b2] = rx[local_n2]; y2[b2] = ry[local_n2]; z2[b2] = rz[local_n2]; m2[b2] = mass[n2 + b2]; } for(size_t b1 = 0; b1 != BLOCK_SIZE; ++b1) { for(size_t b2 = 0; b2 != BLOCK_SIZE; ++b2) { nbcoord_t dx = x1[b1] - x2[b2]; nbcoord_t dy = y1[b1] - y2[b2]; nbcoord_t dz = z1[b1] - z2[b2]; nbcoord_t r2(dx * dx + dy * dy + dz * dz); if(r2 < NBODY_MIN_R) { r2 = NBODY_MIN_R; } nbcoord_t r = sqrt(r2); nbcoord_t coeff = (m2[b2]) / (r * r2); dx *= coeff; dy *= coeff; dz *= coeff; total_force_x[b1] -= dx; total_force_y[b1] -= dy; total_force_z[b1] -= dz; } } } for(size_t b1 = 0; b1 != BLOCK_SIZE; ++b1) { size_t local_n1 = b1 + n1; frx[local_n1] = vx[local_n1]; fry[local_n1] = vy[local_n1]; frz[local_n1] = vz[local_n1]; fvx[local_n1] = total_force_x[b1]; fvy[local_n1] = total_force_y[b1]; fvz[local_n1] = total_force_z[b1]; } }
Tabel 2. Ketergantungan waktu perhitungan (dalam detik) dari fungsi fn pada jumlah badan yang berinteraksi N untuk berbagai implementasi CPUN | 2048 | 4096 | 8192 | 16384 | 32768 |
---|
sederhana | 0,0425 | 0,1651 | 0,6594 | 2.65 | 10.52 |
openmp | 0,0078 | 0,0260 | 0,1079 | 0,417 | 1.655 |
openmp + blok + optimisasi | 0,0037 | 0,0128 | 0,0495 | 0,194 | 0,774 |
Parameter sistem:- sistem: Debian 9, Intel Core i7-5820K (6 inti)
- kompiler: gcc 6.3.0
Terlihat jelas bahwa versi dengan dukungan OpenMP dipercepat sebanyak enam kali, tepatnya dalam hal jumlah inti, dan versi yang dioptimalkan lebih cepat sedikit lebih dari dua kali. Jadi, dengan optimisasi, Anda seharusnya tidak hanya mengandalkan konkurensi. Menariknya, dalam perhitungan aliran tunggal, prosesor bekerja pada frekuensi 3,6 GHz, dalam versi paralel (openmp) menjatuhkan frekuensi menjadi 3,4 GHz, dan secara paralel dan dioptimalkan (openmp + blok + optimasi) turun ke 3,3 GHz, tetapi ini itu tidak menghentikannya untuk bekerja 13,6 kali lebih cepat. Terlihat juga bahwa peningkatan waktu komputasi dengan peningkatan ukuran masalah adalah kuadratik, dan lebih jauh
N membuat tugas tidak terpecahkan dalam waktu yang wajar.
Implementasi GPU
Tetapi ada keinginan untuk membuat perhitungan lebih cepat. Ada beberapa arah yang tersedia untuk akselerasi: Perhitungan GPU, perkiraan fungsi
fn . Pertama, untuk komputasi GPU, saya memilih teknologi OpenCL. Untuk pekerjaan yang lebih nyaman, perpustakaan
CLHPP digunakan. Keuntungan utama OpenCL adalah bahwa kode dapat dijalankan pada prosesor dan GPU, yang menyederhanakan penulisan dan debugging, serta memperluas daftar perangkat keras yang akan dijalankan. Alat
Oclgrind membantu dalam debugging, yang pada saat runtime menunjukkan kesalahan panggilan OpenCL API dan masalah akses memori.
Opencl
Untuk memulai dengan OpenCL, Anda perlu mendapatkan daftar platform yang tersedia. Platform yang paling umum adalah AMD, Intel, dan NVidia.
Kode std::vector<cl::Platform> platforms; cl::Platform::get(&platforms);
Selanjutnya, setelah memilih platform, Anda harus memilih perangkat komputasi yang diwakili platform ini:
Kode const cl::Platform& platform(platforms[platform_n]); std::vector<cl::Device> all_devices; platform.getDevices(CL_DEVICE_TYPE_ALL, &all_devices);
Dan pada akhir fase persiapan, perlu untuk membuat konteks dan antrian di mana memori akan dialokasikan dan perhitungan akan dilakukan. Misalnya, konteks yang menggabungkan semua perangkat komputasi dari platform yang dipilih dibuat sebagai berikut:
Konteks dan Kode Pembuatan Antrian cl::Context context(all_devices); std::vector<cl::CommandQueue> queues; for(cl::Device device: all_devices) queues.push_back(cl::CommandQueue(context, device));
Untuk mengunduh kode sumber ke perangkat komputer, kode itu harus dikompilasi, kelas cl :: Program ditujukan untuk ini.
Kode kompilasi kernel std::vector< std::string > source_data; cl::Program::Sources sources; for(int i = 0; i != files.size(); ++i) { source_data.push_back(load_program(files[i]));
Untuk menjelaskan parameter fungsi (kernel) yang dijalankan pada perangkat komputasi, ada templat cl: make_kernel.
Contoh kernel perhitungan kekuatan interaksi typedef cl::make_kernel< cl_int, cl_int,
Lebih lanjut, semuanya sederhana: kita mendeklarasikan variabel dengan tipe kernel, meneruskan program yang dikompilasi dan nama kernel komputasi ke dalamnya, kita dapat memulai kernel hampir seperti fungsi normal.
Kode Peluncuran Kernel ComputeBlock fcompute(prog, "ComputeBlockLocal"); cl::NDRange global_range(device_data_size); cl::NDRange local_range(block_size); cl::EnqueueArgs eargs(ctx.m_queue, global_range, local_range); fcompute(eargs, ... );
Inti komputasi untuk OpenCL sendiri sangat mirip dengan opsi 'openmp + block + optimation' untuk CPU, hanya tidak seperti versi CPU, siklus pertama dikontrol menggunakan OpenCL (kisaran siklus ditentukan oleh variabel global_range dari kode startup kernel, dan nomor iterasi saat ini tersedia dari kernel menggunakan fungsi get_global_id (0)). Pertama, sebagian data tubuh dimuat dari memori lokal, lalu diproses. Memori lokal umum untuk semua utas dalam grup, jadi unduhan terjadi sekali untuk grup, dan diproses oleh setiap utas dalam grup dan, karena memori lokal jauh lebih cepat daripada global, perhitungannya jauh lebih cepat.
Kode kernel __kernel void ComputeBlockLocal(int offset_n1, int offset_n2, __global const nbcoord_t* mass, __global const nbcoord_t* y, __global nbcoord_t* f, int yoff, int foff, int points_count, int stride) { int n1 = get_global_id(0) + offset_n1; __global const nbcoord_t* rx = y + yoff; __global const nbcoord_t* ry = rx + stride; __global const nbcoord_t* rz = rx + 2 * stride; __global const nbcoord_t* vx = rx + 3 * stride; __global const nbcoord_t* vy = rx + 4 * stride; __global const nbcoord_t* vz = rx + 5 * stride; __global nbcoord_t* frx = f + foff; __global nbcoord_t* fry = frx + stride; __global nbcoord_t* frz = frx + 2 * stride; __global nbcoord_t* fvx = frx + 3 * stride; __global nbcoord_t* fvy = frx + 4 * stride; __global nbcoord_t* fvz = frx + 5 * stride; nbcoord_t x1 = rx[n1]; nbcoord_t y1 = ry[n1]; nbcoord_t z1 = rz[n1]; nbcoord_t res_x = 0.0; nbcoord_t res_y = 0.0; nbcoord_t res_z = 0.0; __local nbcoord_t x2[NBODY_DATA_BLOCK_SIZE]; __local nbcoord_t y2[NBODY_DATA_BLOCK_SIZE]; __local nbcoord_t z2[NBODY_DATA_BLOCK_SIZE]; __local nbcoord_t m2[NBODY_DATA_BLOCK_SIZE];
Cuda
Implementasi untuk platform NVidia CUDA sedikit lebih sederhana dari OpenCL, kami tidak perlu membuat konteks perangkat dan mengelola antrian eksekusi (setidaknya sampai kami ingin membuat implementasi multi-GPU). Seperti dalam kasus OpenCL, kita perlu mengalokasikan memori pada GPU, menyalin data kita ke dalamnya, dan kemudian kita dapat memulai inti komputasi.
Anda dapat membaca lebih lanjut tentang bekerja dengan CUDA di
sini .
Kode Peluncuran Kernel CUDA dim3 grid(count / block_size); dim3 block(block_size); size_t shared_size(4 * sizeof(nbcoord_t) * block_size); kfcompute <<< grid, block, shared_size >>> (... ...);
Tidak seperti OpenCL, CUDA tidak menentukan rentang penuh iterasi (dalam implementasi OpenCL itu adalah global_range), tetapi menetapkan ukuran kisi dan ukuran blok dalam kisi, sehingga, perhitungan jumlah tubuh saat ini sedikit berubah, jika tidak, kernel sangat mirip dengan OpenCL, dengan pengecualian nama lain fungsi sinkronisasi dan penentu untuk memori bersama. Fitur pembeda lain yang berguna dari CUDA adalah bahwa kita dapat menentukan ukuran memori bersama yang diperlukan saat kernel dimulai. Seperti dalam implementasi OpenCL, pada awal setiap blok iterasi kami menyalin bagian dari data ke memori bersama dan kemudian bekerja dengan memori ini dari semua utas blok.
Kode kernel CUDA __global__ void kfcompute(int offset_n2, const nbcoord_t* y, int yoff, nbcoord_t* f, int foff, const nbcoord_t* mass, int points_count, int stride) { int n1 = blockDim.x * blockIdx.x + threadIdx.x; const nbcoord_t* rx = y + yoff; const nbcoord_t* ry = rx + stride; const nbcoord_t* rz = rx + 2 * stride; nbcoord_t x1 = rx[n1]; nbcoord_t y1 = ry[n1]; nbcoord_t z1 = rz[n1]; nbcoord_t res_x = 0.0; nbcoord_t res_y = 0.0; nbcoord_t res_z = 0.0; extern __shared__ nbcoord_t shared_xyzm_buf[]; nbcoord_t* x2 = shared_xyzm_buf; nbcoord_t* y2 = x2 + blockDim.x; nbcoord_t* z2 = y2 + blockDim.x; nbcoord_t* m2 = z2 + blockDim.x; for(int b2 = 0; b2 < points_count; b2 += blockDim.x) { int n2 = b2 + offset_n2 + threadIdx.x;
Tabel 3. Ketergantungan fungsi waktu perhitungan (dalam detik) fn pada jumlah badan yang berinteraksi N untuk berbagai implementasi GPU dan implementasi CPU yang lebih baikN | 4096 | 8192 | 16384 | 32768 | 65536 | 131072 |
---|
openmp + blok + optimisasi | 0,0128 | 0,0495 | 0,194 | 0,774 | --- | --- |
OpenCL + setengah NVidia K80 | 0,004 | 0,008 | 0,026 | 0,134 | 0,322 | 1.18 |
CUDA + setengah NVidia K80 | 0,004 | 0,008 | 0,0245 | 0,115 | 0,291 | 1.13 |
Di mana mendapatkan NVidia K80Implementasi OpenCL dan CUDA diluncurkan pada layanan
Colab gratis
Google , yang menyediakan akses ke komputer NVidia K80. Anda dapat membaca lebih lanjut tentang bekerja dengan layanan ini di
hub Secara umum, hasilnya mengecewakan, hanya 5-6 kali lebih cepat daripada implementasi CPU. Bahkan jika kita melakukan perhitungan pada seluruh K80, kita akan mendapatkan akselerasi hingga 12 kali, tetapi sejak itu Karena kompleksitas tugasnya kuadratik, maka dalam waktu yang masuk akal kita akan dapat memproses bukan 32.768 badan yang berinteraksi, tetapi 131072, yang hanya 4 kali lebih banyak.
Perkiraan fungsi fn
Jika Anda melihat dekat pada fungsi, yang diatur oleh kekuatan menarik dari dua tubuh, jelas bahwa itu berkurang secara kuadrat dengan jarak. Oleh karena itu, kita dapat secara akurat menghitung kekuatan interaksi antara benda yang dekat, dan kira-kira antara yang jauh. Satu pendekatan terkenal
adalah algoritma
treecode yang diusulkan oleh D. Barnes dan P. Hat.
Pohon octree dibangun di ruang simulasi, berisi koordinat daun dan massa badan simulasi. Simpul induk mengandung pusat massa, massa total simpul anak dan jari-jari bola yang dijelaskan di sekitar tubuh simpul anak. Akar pohon berisi pusat massa semua benda, massa totalnya, dan jari-jari bola yang dijelaskan di sekelilingnya. Saat menghitung gaya interaksi, jarak ke akar pohon pertama kali dipertimbangkan jika rasio jarak ke simpul ke jari-jarinya lebih besar dari beberapa konstanta
lambdacrit , maka root dianggap satu tubuh dengan koordinat sama dengan koordinat pusat massa tubuh yang termasuk di dalamnya, dan massa sama dengan jumlah massa simpul anak perempuan, jika rasionya kurang dari atau sama dengan
lambdacrit , lalu prosedur ini diulang secara rekursif untuk setiap simpul anak. Jika daun pohon tercapai, maka kekuatan interaksi dipertimbangkan dengan cara biasa. Jadi, jika satu benda jauh dari kelompok kompak dari badan lain, kelompok ini disajikan sebagai satu tubuh, dan gaya interaksi dihitung tidak sampai masing-masing tubuh, tetapi hanya hingga satu tubuh. Karena ini, kompleksitas algoritma menurun dengan
O(N2) sebelumnya
O(N log(N)) dengan mengorbankan beberapa kehilangan akurasi.
Dalam pendekatan saya, saya memutuskan untuk tidak menggunakan
pohon octree , tetapi
pohon kd, karena lebih mudah digunakan dan memiliki overhead penyimpanan yang lebih rendah dibandingkan dengan octree.
Mari kita kembali ke implementasi pada CPU. Node kd-tree dapat direpresentasikan dalam bentuk kelas yang berisi pointer ke turunan kiri dan kanan dan informasi tentang koordinat dan massa:
Simpul pohon Kd class node { node* m_left;
Dengan metode ini menyimpan pohon, kami memiliki dua opsi yang mungkin untuk melintasi pohon: baik menggunakan rekursi eksplisit, atau gunakan tumpukan sendiri. Saya memilih opsi kedua.
Perhitungan kekuatan interaksi dengan melintasi pohon nbvertex_t force_compute(const nbvertex_t& v1, const nbcoord_t mass1) { nbvertex_t total_force; node* stack_data[MAX_STACK_SIZE] = {}; node** stack = stack_data; node** stack_head = stack; *stack++ = m_root; while(stack != stack_head) { node* curr = *--stack; const nbcoord_t distance_sqr((v1 - curr->m_mass_center).norm()); if(distance_sqr > curr->m_radius_sqr) {
Seperti dalam kasus implementasi CPU "tepat", fungsi perhitungan gaya dipanggil untuk setiap tubuh. Siklus di semua badan dapat dengan mudah diparalelkan menggunakan arahan OpenMP.
Tetapi iterasi loop tetangga dalam kasus ini akan merujuk ke bagian pohon yang sama sekali berbeda, yang tidak memungkinkan penggunaan cache prosesor secara efisien. Untuk mengatasi masalah ini, semua benda harus dilintasi bukan dalam urutan asli, tetapi dalam urutan di mana mayat berada di daun pohon kd, maka iterasi yang berdekatan akan terjadi untuk benda yang dekat di ruang angkasa, dan akan melintasi pohon di sepanjang jalur yang hampir sama.
Lintas daun pohon template<class Visitor> void traverse(Visitor visit) { node* stack_data[MAX_STACK_SIZE] = {}; node** stack = stack_data; node** stack_head = stack; *stack++ = m_root; while(stack != stack_head) { node* curr = *--stack; if(curr->m_radius_sqr > 0) {
Implementasi ini memiliki masalah lain - tidak ada cara universal untuk memaralelkan traversal pohon semacam itu. Tetapi kita dapat sepenuhnya mengubah cara pohon disimpan dalam memori - kita dapat menyimpan semua node dalam satu array linier dan sepenuhnya meninggalkan penyimpanan pointer ke keturunan, dengan analogi dengan pembangunan
tumpukan biner . Saat memulai penomoran node dengan
1 jika simpul ada di nomor sel
i , lalu anak kirinya ada di sel
2 i , anak yang tepat di sel
2 i + 1 , dan induk di sel
i / 2 . Node kanan sesuai dengan kiri dengan nomor
saya akan memiliki nomor
i + 1 . Array itu sendiri akan memiliki panjang
2 N , dan semua simpul daun akan ditemukan di bagian terakhir
N sel, apalagi, node yang berjarak dekat akan berlokasi di sel dekat array. Fungsi traversal pohon akan berubah sedikit, tetapi kerangka kerjanya tetap sama:
Perhitungan kekuatan dengan melintasi pohon dalam array nbvertex_t force_compute(const nbvertex_t& v1, const nbcoord_t mass1) { nbvertex_t total_force; size_t stack_data[MAX_STACK_SIZE] = {}; size_t* stack = stack_data; size_t* stack_head = stack; *stack++ = NBODY_HEAP_ROOT_INDEX; while(stack != stack_head) { size_t curr = *--stack; const nbcoord_t distance_sqr((v1 - m_mass_center[curr]).norm()); if(distance_sqr > m_radius_sqr[curr]) { total_force += force(v1, m_mass_center[curr], mass1, m_mass[curr]); } else { size_t left(left_idx(curr)); size_t rght(rght_idx(curr)); if(rght < m_body_n.size()) { *stack++ = rght; } if(left < m_body_n.size()) { *stack++ = left; } } } return total_force; }
Tapi ini tidak semua kemungkinan bahwa penyimpanan node dalam array terbuka untuk kita - kita bisa menolak stack saat merangkak. Untuk melakukan ini, dalam cabang kode di mana kita pergi ke anak-anak dari simpul, kita menambahkan fungsi untuk menghitung simpul berikutnya (
n e x t u p ), dan di cabang tempat kami menghitung gaya interaksi, kami menambahkan perhitungan simpul berikutnya dengan subtree saat ini dilewati (
s k i p d o w n )
Untuk melewati subtree saat ini, kita perlu turun ke root (arah
D ), ketika kita berada di keturunan kanan, segera setelah kita mencapai kiri, kita pergi ke subtree kanan yang sesuai (arah
R ), jika kita sampai ke root, maka traversal pohon selesai.
Subtree Abaikan Kode Fungsi index_t skip_down(index_t idx) {
Gambar 8. Melewati subtree s k i p d o w n .Untuk pergi ke simpul berikutnya, jika mungkin, pergi ke anak kiri (arah
U ), dan jika tidak ada turunan, maka pergi ke simpul berikutnya 'dari bawah' menggunakan fungsi
s k i p d o w n .
Pergi ke kode fungsi simpul berikutnya index_t next_up(index_t idx, index_t tree_size) { index_t left = left_idx(idx); if(left < tree_size) { return left; } return skip_down(idx); }
Gambar 9. Transisi ke node berikutnya n e x t u p .Tampaknya kami mengganti rekursi dengan loop
s e m e n t a r a dalam fungsi
s k i p d o w n , dan penggantian seperti itu tidak memberikan apa-apa, tetapi mari kita lihat bagaimana menentukan apakah suatu simpul dengan angka
saya keturunan benar. Ini dapat dilakukan hanya dengan memeriksa nomor ganjilnya (anak yang tepat ada di dalam sel
2 i + 1 ), untuk ini cukup untuk menghitung
i \ & 1 . Yaitu kami membagi nomornya
i dua jika bit paling signifikan diatur ke satu. Tapi ini bisa dilakukan tanpa loop, di banyak prosesor ada instruksi
set pertama yang mengembalikan posisi bit set pertama, dengan menggunakannya kita mengubah loop menjadi empat instruksi prosesor.
Kode Fungsi Lewati Subtree Dioptimalkan index_t skip_down(index_t idx) { idx = idx >> (__builtin_ffs(~idx) - 1); return left2right(idx); }
Setelah itu kita bisa mengecualikan stack dari fungsi tree traversal dan menggantinya dengan pair
skipdown+nextup , setelah itu fungsinya terlihat lebih sederhana:
Perhitungan kekuatan dengan melintasi pohon dalam array tanpa menggunakan tumpukan nbvertex_t force_compute(const nbvertex_t& v1, const nbcoord_t mass1) const { nbvertex_t total_force; size_t curr = NBODY_HEAP_ROOT_INDEX; size_t tree_size = m_mass_center.size(); do { const nbcoord_t distance_sqr((v1 - m_mass_center[curr]).norm()); if(distance_sqr > m_radius_sqr[curr]) { total_force += force(v1, m_mass_center[curr], mass1, m_mass[curr]); curr = skip_down(curr); } else { curr = next_up(curr, tree_size); } } while(curr != NBODY_HEAP_ROOT_INDEX); return total_force; }
Secara total, kami mendapat enam kombinasi traversal pohon dan perhitungan gaya. Bandingkan pendekatan ini dalam hal waktu dan kualitas perhitungan. Kami mengambil sebagai ukuran kualitas perubahan relatif dalam total energi sistem setelah 100 iterasi. Sebagai sistem model, kami mengambil dua "galaksi" yang berinteraksi yang terdiri dari16384tubuh masing-masing.Tabel 4. Kombinasi metode traversal pohon dan perhitungan gayaTipe perhitungan traversal / gaya pohon | Pohon dengan tumpukan | 'Tumpukan' dengan tumpukan | 'Tumpukan' tanpa tumpukan |
---|
Iterasi berdasarkan nomor tubuh | cycle + tree | cycle + heap | cycle + heapstackless |
---|
Bypass daun | nestedtree + tree | nestedtree + heap | nestedtree + heapstackless |
---|

| 
|
a) Fungsi perhitungan ketergantungan waktu fn dari rasio jarak kritis ke simpul pohon ke jari-jarinya ( λcrit ) | b) Ketergantungan perubahan relatif energi dalam sistem pada rasio jarak kritis ke simpul pohon terhadap jari-jarinya ( λcrit ) |
Gambar 10. Hasil perhitungan fungsi fn dalam berbagai cara pada CPU (jumlah tubuh N=32768 ) |
Dapat dilihat bahwa implementasi 'nestedtree + tree' tidak ada harapan dalam kecepatan, karena tidak memiliki konkurensi. Implementasi dengan lokasi node pohon dalam array dan pengindeksan seperti dalam tumpukan biner berada di depan. Perubahan relatif dalam energi dapat diabaikan untuk semua varian denganλcrit>1 .
Juga dalam gambar. 10a menunjukkan bahwa untuk (λcrit<10 ) semua varian perhitungan fungsi fn secara signifikan menyalip dalam kecepatan versi yang paling optimal dari perhitungan yang tepat ('openmp + block + optimization'), dengan peningkatan lebih lanjut λcrit implementasi dengan pohon kehilangan versi pastinya.Traversal pohon GPU
Saya mencoba memintas pohon pada GPU menggunakan teknologi OpenCL dan CUDA. Opsi menyimpan node dalam bentuk pohon segera dibuang, dan hanya opsi yang tersisa dengan menyimpan pohon dalam array dengan pengindeksan seperti dalam tumpukan biner. Secara umum, implementasi core komputasi tidak jauh berbeda dengan versi CPU.Kernel OpenCL untuk kekuatan komputasi dengan melintasi pohon (traversal dalam urutan badan penomoran) __kernel void ComputeTreeBH(int offset_n1, int points_count, int tree_size, __global const nbcoord_t* y, __global nbcoord_t* f, __global const nbcoord_t* tree_cmx, __global const nbcoord_t* tree_cmy, __global const nbcoord_t* tree_cmz, __global const nbcoord_t* tree_mass, __global const nbcoord_t* tree_crit_r2) { int n1 = get_global_id(0) + offset_n1; int stride = points_count; __global const nbcoord_t* rx = y; __global const nbcoord_t* ry = rx + stride; __global const nbcoord_t* rz = rx + 2 * stride; __global const nbcoord_t* vx = rx + 3 * stride; __global const nbcoord_t* vy = rx + 4 * stride; __global const nbcoord_t* vz = rx + 5 * stride; __global nbcoord_t* frx = f; __global nbcoord_t* fry = frx + stride; __global nbcoord_t* frz = frx + 2 * stride; __global nbcoord_t* fvx = frx + 3 * stride; __global nbcoord_t* fvy = frx + 4 * stride; __global nbcoord_t* fvz = frx + 5 * stride; nbcoord_t x1 = rx[n1]; nbcoord_t y1 = ry[n1]; nbcoord_t z1 = rz[n1]; nbcoord_t res_x = 0.0; nbcoord_t res_y = 0.0; nbcoord_t res_z = 0.0; int stack_data[MAX_STACK_SIZE] = {}; int stack = 0; int stack_head = stack; stack_data[stack++] = NBODY_HEAP_ROOT_INDEX; while(stack != stack_head) { int curr = stack_data[--stack]; nbcoord_t dx = x1 - tree_cmx[curr]; nbcoord_t dy = y1 - tree_cmy[curr]; nbcoord_t dz = z1 - tree_cmz[curr]; nbcoord_t r2 = (dx * dx + dy * dy + dz * dz); if(r2 > tree_crit_r2[curr]) { if(r2 < NBODY_MIN_R) { r2 = NBODY_MIN_R; } nbcoord_t r = sqrt(r2); nbcoord_t coeff = tree_mass[curr] / (r * r2); dx *= coeff; dy *= coeff; dz *= coeff; res_x -= dx; res_y -= dy; res_z -= dz; } else { int left = left_idx(curr); int rght = rght_idx(curr); if(left < tree_size) { stack_data[stack++] = left; } if(rght < tree_size) { stack_data[stack++] = rght; } } } frx[n1] = vx[n1]; fry[n1] = vy[n1]; frz[n1] = vz[n1]; fvx[n1] = res_x; fvy[n1] = res_y; fvz[n1] = res_z; }
Pada versi pertama, tree traversal dimulai dengan urutan penomoran tubuh pada array asli, sehingga utas tetangga melewati bagian-bagian pohon yang sama sekali berbeda, yang secara negatif mempengaruhi kinerja cache memori GPU. Oleh karena itu, dalam perwujudan kedua, sebuah traversal diterapkan, mulai dari atas pohon, dalam hal ini, benang tetangga mulai melintasi pohon di sepanjang jalan yang sama, karena puncak pohon tetangga dekat dan di luar angkasa. Penting juga bahwa kita memilih penomoran dalam array node pohon bukan dari awal, tetapi dari satu, dalam hal ini daun pohon disimpan di paruh kedua array, dan dengan jumlah tubuh yang sama dengan kekuatan dua, kita akan memiliki akses yang sama ke memori dengan indeks tn1.Kernel OpenCL untuk menghitung kekuatan dengan melintasi pohon (melintasi simpul penomoran pohon) __kernel void ComputeHeapBH(int offset_n1, int points_count, int tree_size, __global const nbcoord_t* y, __global nbcoord_t* f, __global const nbcoord_t* tree_cmx, __global const nbcoord_t* tree_cmy, __global const nbcoord_t* tree_cmz, __global const nbcoord_t* tree_mass, __global const nbcoord_t* tree_crit_r2, __global const int* body_n) { int tree_offset = points_count - 1 + NBODY_HEAP_ROOT_INDEX; int stride = points_count; int tn1 = get_global_id(0) + offset_n1 + tree_offset; int n1 = body_n[tn1]; nbcoord_t x1 = tree_cmx[tn1]; nbcoord_t y1 = tree_cmy[tn1]; nbcoord_t z1 = tree_cmz[tn1]; nbcoord_t res_x = 0.0; nbcoord_t res_y = 0.0; nbcoord_t res_z = 0.0; int stack_data[MAX_STACK_SIZE] = {}; int stack = 0; int stack_head = stack; stack_data[stack++] = NBODY_HEAP_ROOT_INDEX; while(stack != stack_head) { int curr = stack_data[--stack]; nbcoord_t dx = x1 - tree_cmx[curr]; nbcoord_t dy = y1 - tree_cmy[curr]; nbcoord_t dz = z1 - tree_cmz[curr]; nbcoord_t r2 = (dx * dx + dy * dy + dz * dz); if(r2 > tree_crit_r2[curr]) { if(r2 < NBODY_MIN_R) { r2 = NBODY_MIN_R; } nbcoord_t r = sqrt(r2); nbcoord_t coeff = tree_mass[curr] / (r * r2); dx *= coeff; dy *= coeff; dz *= coeff; res_x -= dx; res_y -= dy; res_z -= dz; } else { int left = left_idx(curr); int rght = rght_idx(curr); if(left < tree_size) { stack_data[stack++] = left; } if(rght < tree_size) { stack_data[stack++] = rght; } } } __global const nbcoord_t* vx = y + 3 * stride; __global const nbcoord_t* vy = y + 4 * stride; __global const nbcoord_t* vz = y + 5 * stride; __global nbcoord_t* frx = f; __global nbcoord_t* fry = frx + stride; __global nbcoord_t* frz = frx + 2 * stride; __global nbcoord_t* fvx = frx + 3 * stride; __global nbcoord_t* fvy = frx + 4 * stride; __global nbcoord_t* fvz = frx + 5 * stride; frx[n1] = vx[n1]; fry[n1] = vy[n1]; frz[n1] = vz[n1]; fvx[n1] = res_x; fvy[n1] = res_y; fvz[n1] = res_z; }
Saat melintasi urutan penomoran simpul pohon, kami mendapat peningkatan kinerja. Tetapi opsi ini juga dapat ditingkatkan. Memori global tempat node pohon saat ini dioptimalkan untuk akses bersama , mis. utas satu grup harus membaca kata-kata yang terletak di satu blok memori. Dalam kasus kami, traversal pohon dimulai di sepanjang jalur yang sama, dan kami meminta data yang sama dengan semua utas grup, tetapi semakin jauh kami masuk lebih dalam ke pohon, jalur benang tetangga semakin berbeda, dan kami harus meminta data yang berbeda, yang mengurangi kinerja subsistem beberapa kali memori. Tetapi node dari masing-masing subtree milik tingkat yang sama terletak di sel-sel memori yang relatif dekat. Yaitu
ketika melintasi sisa pohon, utas tetangga dari inti komputasi tidak mengakses simpul pohon yang sama, tetapi berada dalam memori. Untuk mengoptimalkan akses memori ini, memori tekstur dapat digunakan. Tapi ada satu halangan. Saat ini, tekstur tidak mendukung kerja dengan data presisi ganda (kami ingin menghitung secara akurat). Tetapi dalam CUDA ada fungsi __hiloint2double yang mengumpulkan angka presisi ganda dari dua bilangan bulat.Kode permintaan angka presisi ganda dari integer penyimpanan tekstur template<> struct nb1Dfetch<double> { typedef double4 vec4; static __device__ double fetch(cudaTextureObject_t tex, int i) { int2 p(tex1Dfetch<int2>(tex, i)); return __hiloint2double(py, px); } static __device__ vec4 fetch4(cudaTextureObject_t tex, int i) { int ii(2 * i); int4 p1(tex1Dfetch<int4>(tex, ii)); int4 p2(tex1Dfetch<int4>(tex, ii + 1)); vec4 d4 = {__hiloint2double(p1.y, p1.x), __hiloint2double(p1.w, p1.z), __hiloint2double(p2.y, p2.x), __hiloint2double(p2.w, p2.z) }; return d4; } };
Dalam hal ini, dua implementasi dibuat, dalam satu setiap elemen pohon (x, y, z, tree_crit_r2) diminta secara independen, dan dalam implementasi kedua permintaan ini digabungkan. Permintaan massa node terjadi jauh lebih jarang, hanya jika kondisi r2> tree_crit_r2 [curr] terpenuhi , jadi tidak masuk akal untuk menggabungkan permintaan ini dengan yang lain. Fitur lain yang berguna dari kerangka kerja CUDA adalah kemampuan untuk mengontrol rasio ukuran cache L1 dan ukuran memori bersama ( cudaFuncSetCacheConfig ). Dalam hal traversal pohon, kami tidak menggunakan memori bersama, sehingga kami dapat meningkatkan cache L1 hingga merusaknya.Kernel CUDA untuk menghitung kekuatan dengan melintasi pohon (melintasi urutan penomoran simpul pohon) __global__ void kfcompute_heap_bh_tex(int offset_n1, int points_count, int tree_size, nbcoord_t* f, cudaTextureObject_t tree_xyzr, cudaTextureObject_t tree_mass, const int* body_n) { nb1Dfetch<nbcoord_t> tex; int tree_offset = points_count - 1 + NBODY_HEAP_ROOT_INDEX; int stride = points_count; int tn1 = blockDim.x * blockIdx.x + threadIdx.x + offset_n1 + tree_offset; int n1 = body_n[tn1]; nbvec4_t xyzr = tex.fetch4(tree_xyzr, tn1); nbcoord_t x1 = xyzr.x; nbcoord_t y1 = xyzr.y; nbcoord_t z1 = xyzr.z; nbcoord_t res_x = 0.0; nbcoord_t res_y = 0.0; nbcoord_t res_z = 0.0; int stack_data[MAX_STACK_SIZE] = {}; int stack = 0; int stack_head = stack; stack_data[stack++] = NBODY_HEAP_ROOT_INDEX; while(stack != stack_head) { int curr = stack_data[--stack]; nbvec4_t xyzr2 = tex.fetch4(tree_xyzr, curr); nbcoord_t dx = x1 - xyzr2.x; nbcoord_t dy = y1 - xyzr2.y; nbcoord_t dz = z1 - xyzr2.z; nbcoord_t r2 = (dx * dx + dy * dy + dz * dz); if(r2 > xyzr2.w) { if(r2 < NBODY_MIN_R) { r2 = NBODY_MIN_R; } nbcoord_t r = sqrt(r2); nbcoord_t coeff = tex.fetch(tree_mass, curr) / (r * r2); dx *= coeff; dy *= coeff; dz *= coeff; res_x -= dx; res_y -= dy; res_z -= dz; } else { int left = nbody_heap_func<int>::left_idx(curr); int rght = nbody_heap_func<int>::rght_idx(curr); if(left < tree_size) { stack_data[stack++] = left; } if(rght < tree_size) { stack_data[stack++] = rght; } } } f[n1 + 3 * stride] = res_x; f[n1 + 4 * stride] = res_y; f[n1 + 5 * stride] = res_z; }
Analisis program dalam profiler nvprof menunjukkan bahwa bahkan ketika menggunakan memori tekstur untuk menyimpan pohon, masih ada beban yang sangat tinggi pada memori global.Memang, dalam CUDA, semua memori kernel yang ditujukan ke alamat 'dihitung' disimpan dalam memori global, dan karenanya, tumpukan yang diperlukan untuk melintasi pohon terletak di memori global dan menggerogoti bagian signifikan dari bandwidth chip memori, karena tumpukan tersebut memiliki masing-masing menjalankan thread, dan ada banyak utas.Tapi, untungnya, kita sudah tahu cara melintasi pohon tanpa menggunakan tumpukan. Melengkapi inti komputasi sebelumnya dengan fungsi komputasi simpul pohon berikutnya, kami mendapatkan kernel baru, lebih dari itu, lebih ringkas.Inti CUDA untuk kekuatan komputasi dengan melintasi pohon tanpa menggunakan tumpukan __global__ void kfcompute_heap_bh_stackless(int offset_n1, int points_count, int tree_size, nbcoord_t* f, cudaTextureObject_t tree_xyzr, cudaTextureObject_t tree_mass, const int* body_n) { nb1Dfetch<nbcoord_t> tex; int tree_offset = points_count - 1 + NBODY_HEAP_ROOT_INDEX; int stride = points_count; int tn1 = blockDim.x * blockIdx.x + threadIdx.x + offset_n1 + tree_offset; int n1 = body_n[tn1]; nbvec4_t xyzr = tex.fetch4(tree_xyzr, tn1); nbcoord_t x1 = xyzr.x; nbcoord_t y1 = xyzr.y; nbcoord_t z1 = xyzr.z; nbcoord_t res_x = 0.0; nbcoord_t res_y = 0.0; nbcoord_t res_z = 0.0; int curr = NBODY_HEAP_ROOT_INDEX; do { nbvec4_t xyzr2 = tex.fetch4(tree_xyzr, curr); nbcoord_t dx = x1 - xyzr2.x; nbcoord_t dy = y1 - xyzr2.y; nbcoord_t dz = z1 - xyzr2.z; nbcoord_t r2 = (dx * dx + dy * dy + dz * dz); if(r2 > xyzr2.w) { if(r2 < NBODY_MIN_R) { r2 = NBODY_MIN_R; } nbcoord_t r = sqrt(r2); nbcoord_t coeff = tex.fetch(tree_mass, curr) / (r * r2); dx *= coeff; dy *= coeff; dz *= coeff; res_x -= dx; res_y -= dy; res_z -= dz; curr = nbody_heap_func<int>::skip_idx(curr); } else { curr = nbody_heap_func<int>::next_up(curr, tree_size); } } while(curr != NBODY_HEAP_ROOT_INDEX); f[n1 + 3 * stride] = res_x; f[n1 + 4 * stride] = res_y; f[n1 + 5 * stride] = res_z; }
Kinerja core yang berjalan pada GPU sangat tergantung pada ukuran blok yang digunakan untuk membagi tugas. Tergantung pada ukuran ini berapa banyak register, memori lokal dan sumber daya lainnya akan tersedia untuk setiap utas komputasi. Anda juga perlu mengingat bahwa sambil menunggu akses memori di satu utas, utas lain dapat melakukan perhitungan pada prosesor shader, dengan demikian, dengan jumlah yang cukup dari utas yang dijalankan secara bersamaan pada satu prosesor, waktu akses memori akan disembunyikan di balik perhitungan. Karena itu, sebelum membandingkan kinerja core kami, kami perlu menghitung ukuran blok optimal untuk masing-masing core. Mari kita bandingkan setengah yang tersedia untuk kita dari NVidia K80.Tabel 5. Ketergantungan fungsi waktu perhitungan (dalam detik) fn GPU N=131072 λcrit=10/ | 8 | 16 | 32 | 64 | 128 | 256 | 512 | 1024 |
---|
opencl+dense | 5.77 | 2.84 | 1.46 | 1.13 | 1.15 | 1.14 | 1.14 | 1.13 |
cuda+dense | 5.44 | 2.55 | 1.27 | 0.96 | 0.97 | 0.99 | 0.99 | - |
opencl+heap+cycle | 5.88 | 5.65 | 5.74 | 5.96 | 5.37 | 5.38 | 5.35 | 5.38 |
opencl+heap+nested | 4.54 | 3.68 | 3.98 | 5.25 | 5.46 | 5.41 | 5.48 | 5.31 |
cuda+heap+nested | 3.62 | 2.81 | 2.68 | 4.26 | 4.84 | 4.75 | 4.8 | 4.67 |
cuda+heap+nested+tex | 2.6 | 1.51 | 0.912 | 0.7 | 1.85 | 1.75 | 1.69 | 1.61 |
cuda+heap+nested+tex+stackless | 2.3 | 1.29 | 0.773 | 0.5 | 0.51 | 0.52 | 0.52 | 0.52 |
6. ( ) fn GPU N=1M λcrit=4/ | 8 | 16 | 32 | 64 | 128 | 256 | 512 | 1024 |
---|
opencl+dense | 366 | 179 | 89.9 | 69.3 | 70.3 | 69.1 | 68.9 | 68.0 |
cuda+dense | 346 | 162 | 79.6 | 60.8 | 60.8 | 60.7 | 59.6 | - |
opencl+heap+cycle | 16.2 | 18.2 | 20.1 | 21.2 | 21.2 | 21.3 | 21.2 | 21.1 |
opencl+heap+nested | 10.5 | 7.63 | 6.38 | 8.23 | 9.95 | 9.89 | 9.65 | 9.66 |
cuda+heap+nested | 8.67 | 6.44 | 5.39 | 5.93 | 8.65 | 8.61 | 8.41 | 8.27 |
cuda+heap+nested+tex | 6.38 | 3.57 | 2.13 | 1.44 | 3.56 | 3.46 | 3.30 | 3.29 |
cuda+heap+nested+tex+stackless | 5.78 | 3.19 | 1.83 | 1.21 | 1.11 | 1.10 | 1.11 | 1.13 |
Situasi yang sulit, tetapi, tidak seperti versi CPU dari traversal pohon, jelas bahwa setiap langkah optimasi membawa hasil nyata. Implementasi 'opencl + heap + cycle' hampir 6 kali lebih lambat dari solusi yang tepat dengan perhitungan fungsi penuhfn .
Implementasi 'opencl + heap + nested', di mana traversal pohon dimulai dari node tetangga, 1,4 kali lebih cepat dari yang sebelumnya, karena Memori cache digunakan lebih baik. Dalam implementasi 'cuda + heap + nested', L1 cache ditingkatkan hingga merusak memori bersama, yang meningkatkan kecepatannya sebesar 1,4 kali, meskipun ada kemungkinan bahwa dalam implementasi cuda kernel kompilasi lebih optimal dikompilasi (dalam versi 'opencl + padat' dan 'cuda + core padat identik, dan kinerja versi cuda ~ 1,2 kali lebih tinggi). Peningkatan kecepatan komputasi yang paling signifikan (3,8 kali) dicapai ketika pohon terletak di memori tekstur dan pertanyaan ke elemen simpul pohon digabungkan. Implementasi dengan tree traversal tanpa menggunakan tumpukan 'cuda + heap + nested + tex + stackless' juga 1,4 kali lebih cepat.di dalamnya, seluruh bandwidth bus memori hanya digunakan untuk mengakses data tentang simpul pohon dan tidak dihabiskan di stack. Demikianlah denganλcrit=10 Dimungkinkan untuk mencapai akselerasi dua kali dibandingkan dengan perhitungan penuh fungsi fn . Tapi
λcrit=10 nilai parameter yang terlalu besar, pada grafik perubahan relatif energi dalam sistem dari rasio jarak kritis ke simpul pohon hingga radiusnya untuk implementasi CPU, dapat dilihat bahwa nilai yang lebih rendah dapat digunakan λcrittanpa kehilangan akurasi. Mari kita coba bervariasiλcrit dengan ukuran blok optimal yang kami tentukan pada langkah sebelumnya.
| 
|
a) N=128K
| b) N=1M
|
Gambar 11. Ketergantungan waktu perhitungan fungsi fn dari rasio jarak kritis ke simpul pohon ke jari-jarinya ( λcrit ) untuk berbagai implementasi GPU tree walk
|
Dapat dilihat bahwa untuk kecil λcrit semua metode perhitungan fungsi fnpergi untuk menutup nilai, ditentukan oleh waktu membangun kd-tree dan menyiapkan data untuk GPU. Selain itu, waktu membangun pohon memberikan kontribusi yang signifikan terhadap total waktuλcrit≤4, maka saat ini bisa diabaikan. Sangat menarik untuk dicatat bahwa kapanN=128K kinerja meningkat lagi setelah mencapai λcrit=1024kemungkinan besar ini disebabkan oleh kenyataan bahwa semua utas GPU mengelilingi simpul pohon yang sama pada saat yang sama, yang meningkatkan penggunaan cache L1 dan sepenuhnya menghilangkan 'cabang divergensi' . Terlihat juga bahwa kinerja terbaik untuk implementasi tanpa stack (cuda + heap + nested + tex + stackless), ini mengungguli versi dengan stack sekitar1.4−1.5kali. Implementasi lain beberapa kali lebih lambat. Untuk mengkonsolidasikan hasilnya, kami akan menghitung waktu pada GPU dengan arsitektur yang lebih baru.Meluncurkan Hasil pada GeForce GTX 1080 Ti (Single Accuracy)7. ( ) fn GPU N=1M λcrit=4/ | 8 | 16 | 32 | 64 | 128 | 256 | 512 | 1024 |
---|
opencl+dense | 47.8 | 23.4 | 11.6 | 11.59 | 12.8 | 12.8 | 12.8 | 12.8 |
cuda+dense | 49.0 | 23.8 | 11.9 | 11.8 | 11.7 | 11.7 | 11.7 | 11.7 |
opencl+heap+cycle | 7.48 | 8.26 | 7.73 | 7.36 | 7.33 | 7.27 | 7.25 | 7.26 |
opencl+heap+nested | 1.33 | 1.20 | 1.41 | 2.46 | 2.53 | 2.49 | 2.44 | 2.47 |
cuda+heap+nested | 1.23 | 1.10 | 1.31 | 2.28 | 2.29 | 2.28 | 2.23 | 2.25 |
cuda+heap+nested+tex | 0.88 | 0.68 | 0.654 | 1.6 | 1.6 | 1.6 | 1.6 | 1.6 |
cuda+heap+nested+tex+stackless | 0.71 | 0.47 | 0.45 | 0.43 | 0.43 | 0.42 | 0.41 | 0.395 |

|
12. fn ( λcrit ) GPU
|
Saat menggunakan GeForce GTX 1080 Ti untuk perhitungan, perbedaan antara jalan-jalan pohon dengan tumpukan dan tanpa tumpukan mencapai dua kali, asalkan kita mengabaikan waktu yang diperlukan untuk membangun pohon. Fakta ini mendorong kami untuk memastikan bahwa di masa depan porting ke GPU dan membangun pohon. Perbandingan tabel 5-7 menunjukkan bahwa tidak ada ukuran blok optimal tunggal untuk jumlah tubuh yang berbeda, dan bahkan lebih lagi untuk arsitektur GPU yang berbeda, apalagi, perbedaan waktu perhitungan dapat mencapai beberapa kali, bahkan jika Anda tidak memperhitungkan nilai batas. Jadi, sebelum perhitungan panjang, masuk akal untuk menentukan ukuran blok optimal untuk setiap konfigurasi.Hal utama yang kami capai adalah kemampuan untuk memodelkan interaksi berpasangan dari sedikit lebih dari satu juta tubuh (220) untuk waktu yang masuk akal, bukan GPU terbaru. Pada GPU yang lebih baru (Tesla V100), tampaknya, akan mungkin untuk memproses sekitar dua juta badan yang berinteraksi dalam waktu yang wajar, karena kinerja mereka sekitar empat kali lebih tinggi daripada setengah dari Tesla K80. Meskipun jumlah ini juga tak tertandingi dengan jumlah bintang bahkan di galaksi kerdil, seperti Awan Magellan Kecil , tetapi, menurut saya, sangat mengesankan.Kesimpulan
Penggunaan metode embedded untuk menyelesaikan persamaan diferensial memungkinkan untuk menyelesaikan masalah serupa dengan tingkat akurasi tinggi dengan biaya waktu yang relatif kecil, dan penggunaan algoritma aproksimasi untuk fungsi menghitung gaya tarik berpasangan memungkinkan kita untuk memproses volume kolosal dari benda yang saling berinteraksi. Dengan demikian, tidak seperti alien dari "Tugas Tiga Tubuh", kami dapat menyelesaikan masalahN tubuh, dan untuk sejumlah kecil tubuh dan untuk seluruh gugus bintang, meskipun dengan mengorbankan sejumlah akurasi.Visualisasi
Bagi mereka yang telah membaca sampai akhir, saya akan memberikan beberapa video lagi dengan visualisasi evolusi sistem tubuh.Simulasi tabrakan dua galaksi. Jumlah mayat adalah 60 ribu.Memodelkan evolusi sebuah galaksi yang terdiri dari sejuta bintang. Di tengah adalah tubuh dengan massa 99% dari total. Partikel tunggal hampir tidak bisa dibedakan. Sudah lebih seperti gelombang dalam setetes cairan. Berwarna sesuai dengan modul kecepatan. Kecepatan rendah - biru, sedang - hijau, tinggi - merah. Terlihat bahwa di bagian tengah kecepatannya lebih tinggi, dan secara bertahap berkurang ke tepinya, dan yang terendah di bidang ekuator.Simulasi yang lebih 'dinamis' dari evolusi galaksi jutaan bintang. Di tengah adalah tubuh dengan massa 10% dari total. Tubuh pusat mempengaruhi sisanya secara signifikan lebih sedikit. Pertama, "bintang-bintang" terbang terpisah, kemudian berkumpul kembali menjadi beberapa kelompok, dan pada akhirnya kembali membentuk satu kelompok besar.Dalam proses pemodelan, sekitar 5% dari "bintang" meninggalkan wilayah awal tidak dapat ditarik kembali.Pada detik ke-10, sangat mirip dengan galaksi roda satu yang ada .Kode proyek dapat ditemukan di github .