Pada artikel ini, kita akan mencari tahu apa untuk dua contoh.
Program pertama adalah axpy paling sederhana (analog dari hello world untuk GPGPU). Yang kedua membantu untuk memahami implementasi kondisi dan melompat pada GPU, karena semuanya berbeda di sana.
Semua bahasa Nvidia menggunakan sedikit endian encoding, jadi segera salin byte dari hex editor ke beberapa notebook (misalnya, Notepad ++) dalam urutan terbalik 8 buah (panjang instruksi konstan di sini). Kemudian, melalui kalkulator programmer (yang dari Microsoft cocok), kami menerjemahkannya ke dalam kode biner. Selanjutnya, kami mencari korek api, membuat topeng instruksi, lalu operan. Untuk memecahkan kode dan mencari topeng, hex editor dan cuobjdump disassembler digunakan, kadang-kadang assembler diperlukan, seperti pada AMDGPU (karena disassembler tidak tersedia di sana, tetapi ini adalah topik untuk artikel terpisah). Ini bekerja seperti ini: kami mencoba membalikkan semua bit yang mencurigakan dalam kalkulator, kemudian kami mendapatkan nilai heksadesimal baru untuk byte, kami menggantinya menjadi biner yang dikompilasi melalui nvcc atau assembler, jika ada, yang tidak selalu demikian.Kemudian melalui cuobjdump kita periksa.
Saya memposting kode sumber dalam format (kebanyakan dalam C, tanpa plus dan OOP untuk koneksi yang lebih dekat dengan kode mesin GPU), kemudian disasm + byte sekaligus, karena lebih nyaman dengan cara itu, mereka tidak perlu ditukar.
Salin ke axpy.cu dan kompilasi melalui cmd: nvcc axpy.cu --cubin --gpu-architecture sm_30 Bongkar
file ELF yang dihasilkan bernama axpy.cubin di tempat yang sama: cuobjdump axpy.cubin -sass
Contoh 1:
__global__ void axpy(float param_1, float* param_2, float* param_3) {
unsigned int uVar1 = threadIdx.x;
param_2[uVar1] = param_1 * param_3[uVar1];
}
Membuang
/*0000*/
/* 0x22c04282c2804307 */
/*0008*/ MOV R1, c[0x0][0x44];
/* 0x2800400110005de4 */
/*0010*/ S2R R0, SR_TID.X;
/* 0x2c00000084001c04 */
/*0018*/ MOV32I R5, 0x4;
/* 0x1800000010015de2 */
/*0020*/ ISCADD R2.CC, R0, c[0x0][0x150], 0x2;
/* 0x4001400540009c43 */
/*0030*/ LD.E R2, [R2];
/* 0x8400000000209c85 */
/*0038*/ ISCADD R4.CC, R0, c[0x0][0x148], 0x2;
/* 0x4001400520011c43 */
/*0040*/
/* 0x20000002e04283f7 */
/*0048*/ IMAD.U32.U32.HI.X R5, R0, R5, c[0x0][0x14c];
/* 0x208a800530015c43 */
/*0050*/ FMUL R0, R2, c[0x0][0x140];
/* 0x5800400500201c00 */
/*0058*/ ST.E [R4], R0;
/* 0x9400000000401c85 */
/*0060*/ EXIT;
/* 0x8000000000001de7 */
/*0068*/ BRA 0x68;
/* 0x4003ffffe0001de7 */
/*0070*/ NOP;
/* 0x4000000000001de4 */
/*0078*/ NOP;
/* 0x4000000000001de4 */
Hasil dekompilasi
void axpy(float param_1,float *param_2,float *param_3) {
uint uVar1;
uVar1 = *&threadIdx.x;
param_2[uVar1] = param_3[uVar1] * param_1;
return;
}Contoh 2:
__global__ void predicates(float* param_1, float* param_2) {
unsigned int uVar1 = threadIdx.x + blockIdx.x * blockDim.x;
if ((uVar1 > 5) & (uVar1 < 10)) param_1[uVar1] = uVar1;
else param_2[uVar1] = uVar1;
}
Membuang
/*0000*/
/* 0x2272028042823307 */
/*0008*/ MOV R1, c[0x0][0x44];
/* 0x2800400110005de4 */
/*0010*/ S2R R0, SR_TID.X;
/* 0x2c00000084001c04 */
/*0018*/ S2R R3, SR_CTAID.X;
/* 0x2c0000009400dc04 */
/*0020*/ IMAD R0, R3, c[0x0][0x28], R0;
/* 0x20004000a0301ca3 */
/*0028*/ MOV32I R3, 0x4;
/* 0x180000001000dde2 */
/*0030*/ IADD32I R2, R0, -0x6;
/* 0x0bffffffe8009c02 */
/*0038*/ I2F.F32.U32 R4, R0;
/* 0x1800000001211c04 */
/*0040*/
/* 0x22c042e04282c2c7 */
/*0048*/ ISETP.GE.U32.AND P0, PT, R2, 0x4, PT;
/* 0x1b0ec0001021dc03 */
/*0050*/ @P0 ISCADD R2.CC, R0, c[0x0][0x148], 0x2;
/* 0x4001400520008043 */
/*0058*/ @P0 IMAD.U32.U32.HI.X R3, R0, R3, c[0x0][0x14c];
/* 0x208680053000c043 */
/*0060*/ @P0 ST.E [R2], R4;
/* 0x9400000000210085 */
/*0068*/ @P0 EXIT;
/* 0x80000000000001e7 */
/*0070*/ ISCADD R2.CC, R0, c[0x0][0x140], 0x2;
/* 0x4001400500009c43 */
/*0078*/ MOV32I R3, 0x4;
/* 0x180000001000dde2 */
/*0080*/
/* 0x2000000002e04287 */
/*0088*/ IMAD.U32.U32.HI.X R3, R0, R3, c[0x0][0x144];
/* 0x208680051000dc43 */
/*0090*/ ST.E [R2], R4;
/* 0x9400000000211c85 */
/*0098*/ EXIT;
/* 0x8000000000001de7 */
/*00a0*/ BRA 0xa0;
/* 0x4003ffffe0001de7 */
/*00a8*/ NOP;
/* 0x4000000000001de4 */
/*00b0*/ NOP;
/* 0x4000000000001de4 */
/*00b8*/ NOP;
/* 0x4000000000001de4 */
Hasil dekompilasi
void predicates(float *param_1,float *param_2) {
uint uVar1;
uVar1 = *&blockIdx.x * (int)_DAT_constants_00000028 + *&threadIdx.x;
if (uVar1 - 6 < 4) {
param_1[uVar1] = (float)uVar1;
return;
}
param_2[uVar1] = (float)uVar1;
return;
}Mudah ditebak bahwa pengujian awalnya disesuaikan dengan kode mesin sehingga kompiler tidak memiliki apa pun untuk dioptimalkan. Untuk yang lainnya, Anda harus membatalkan optimasi secara manual. Dalam contoh kompleks, ini mungkin tidak mungkin sama sekali, jadi untuk kasus seperti itu Anda harus mempercayai dekompiler dan frontend.
Secara umum, aturannya adalah ini - untuk menguji frontend, kami mengambil contoh sederhana (dengan minimum kemungkinan optimasi) pertama yang cocok (mereproduksi kesalahan) contoh. Selebihnya, kode yang didekompilasi akan memiliki optimisasi (atau hanya entah bagaimana memperbaikinya melalui refactoring). Tetapi untuk sekarang, tugas utama adalah setidaknya hanya memperbaiki kode yang melakukan hal yang sama dengan kode mesin. Ini adalah Pemodelan Perangkat Lunak. "Pemodelan perangkat lunak" itu sendiri tidak menyiratkan refactoring, terjemahan C ke C ++, pemulihan kelas, dan bahkan lebih lagi hal-hal seperti identifikasi template.
Sekarang kami mencari pola untuk mnemonik, operan dan pengubah.
Untuk melakukan ini, bandingkan bit (dalam representasi biner) antara instruksi yang mencurigakan (atau string, jika lebih mudah untuk memanggilnya seperti itu). Anda juga dapat menggunakan apa yang diposkan pengguna lain dalam pertanyaan mereka tentang stackoverflow seperti "bantu saya memahami kode biner / sass / mesin", gunakan tutorial (termasuk dalam bahasa Cina) dan sumber daya lainnya. Jadi, nomor operasi utama disimpan dalam bit 58-63, tetapi ada juga bit tambahan 0-4 (mereka membedakan instruksi "I2F", "ISETP", "MOV32I"), di suatu tempat bukannya 0-2 (untuk diabaikan, 3- 4 bit dalam instruksi kosong, dalam spesifikasi mereka ditandai sebagai "UNK").
Untuk register dan angka konstan, Anda dapat bereksperimen dengan disassembler untuk menemukan semua bit yang mempengaruhi output dump, seperti yang diletakkan di bawah spoiler. Semua bidang yang berhasil saya temukan ada dalam spesifikasi pada Github, file CUDA.slaspec, bagian token.
Maka Anda harus datang dengan alamat untuk register, sekali lagi mereka ada di Github. Ini perlu karena di tingkat mikro, Sleigh mendaftar register sebagai variabel global dalam ruang dengan tipe "register_space", tetapi sejak itu Karena ruang mereka tidak ditandai sebagai "dapat disimpulkan" (dan tentu saja tidak bisa), maka mereka dalam dekompiler menjadi variabel lokal (paling sering dengan antarmuka "Var", tetapi terkadang awalan "lokal" juga suka) atau parameter (" param_ "). SP tidak pernah berguna, dibutuhkan sebagian besar secara formal untuk memastikan dekompiler berfungsi. PC (seperti IP dari x86) diperlukan untuk emulasi.
Lalu ada register predikat, seperti bendera, tetapi lebih banyak "tujuan umum" daripada untuk tujuan yang telah ditentukan, seperti melimpah, (dalam) kesetaraan ke nol, dll.
Kemudian register pemblokiran untuk mensimulasikan sekelompok instruksi ISCADD .CC dan IMAD.HI, sejak itu yang pertama dari mereka dalam implementasi saya melakukan penghitungan untuk dirinya sendiri dan untuk yang kedua, untuk menghindari mentransfer bagian dari jumlah ke atas 4 byte, karena ini akan mengacaukan dekompilasi. Tetapi kemudian Anda perlu mengunci register berikutnya sampai operasi IMAD.HI selesai. Sesuatu yang serupa, yaitu perbedaan antara dokumentasi resmi dan output yang diharapkan dari dekompiler sudah ada di SPU untuk Ghidra yang sama.
Lalu ada register khusus yang diimplementasikan melalui cpool sejauh ini. Di masa depan, saya berencana untuk menggantinya dengan simbol default untuk beberapa ruang "yang dapat disimpulkan". Ini adalah threadIdx yang sama, blockIdx.
Kemudian kita mengikat variabel ke bidang dest, par0, par1, par2, res. Lalu ada sub-tabel, dan setelah mereka - apa itu semua tentang - tabel (root) utama dengan instruksi utama.
Di sini Anda harus benar-benar mengikuti format "operan mnemonik", namun, pengecualian diberikan untuk modifier, yang, bagaimanapun, harus dilampirkan pada mnemonik atau pada bagian dengan operan. Tidak ada format lain yang diizinkan, bahkan DSP Hexagon yang sama harus disesuaikan dengan sintaks ini, yang, bagaimanapun, tidak terlalu sulit.
Langkah terakhir adalah menulis implementasi untuk instruksi dalam bahasa firmware Pcode. Satu-satunya hal yang ingin saya perhatikan dari contoh pertama adalah instruksi ISCADD .CC dan IMAD.HI, di mana yang pertama dari mereka mengambil pointer ke register dan dereferensi mereka sebagai pointer ke 8 byte, bukan 4. Ini dilakukan dengan sengaja agar lebih baik beradaptasi dengan dekompiler. dan perilakunya, terlepas dari apa yang tertulis dalam dokumentasi Nvidia tentang transfer bagian dari jumlah tersebut.
Untuk contoh kedua, lebih baik untuk memeriksa pengaturan dekompiler di seberang tulisan "Sederhanakan predikasi". Intinya adalah bahwa predikat adalah satu dan kondisi yang sama untuk instruksi yang berbeda, pada dasarnya tidak lebih dari "SIMD" yang terkenal, atau yang setara berikutnya. Itu jika bit predikat diatur, maka instruksi dieksekusi, apalagi, berturut-turut.
Anda juga perlu membiasakan diri segera menulis implementasi untuk setiap instruksi, dan bukan hanya prototipe (operand mnemonics), karena ada juga dekompiler, emulator, dan penganalisa lainnya.
Tetapi secara umum, menulis implementasi dalam Pcode adalah tugas yang bahkan lebih sederhana daripada menulis tata bahasa untuk decoder byte. Itu cepat untuk memperbaiki implementasi untuk beberapa instruksi kompleks dari x86 (dan tidak hanya), berkat bahasa perantara yang sangat nyaman, satu midland (pengoptimal), 2 backend (terutama C; sebagai alternatif - Java / C #, lebih seperti yang terakhir, dll. K. goto muncul dari waktu ke waktu, tetapi tidak berlabel break).
Dalam artikel berikut, mungkin ada juga frontend untuk bahasa yang dikelola seperti DXBC, SPIR-V, mereka akan menggunakan backend Java / C #. Namun sejauh ini hanya kode mesin yang ada dalam rencana. bytecodes memerlukan pendekatan khusus. Ghidra Bantuan
Proyek : pcode Sleigh