Postingan sebelumnya mengikuti kernel penjumlahan vektor dari nvcc hingga warp. Kali ini, kami mengikuti instruksi SASS kritis (beban global) melalui perangkat keras RTX4090. Kernel CUDA yang diselidiki memiliki dua baris: __global__ void vadd(const float* a, const float* b, float* c, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i];}. SASS yang dikompilasi menunjukkan: /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]. Instruksi-instruksi ini memuat elemen vektor b dari memori global ke register.
Satu LDG.E meminta empat byte di masing-masing dari 32 lane. Melayaninya memerlukan empat sektor 32-byte, satu baris cache, satu terjemahan alamat, satu penyeberangan crossbar, satu dari tiga puluh enam irisan L2, dan ketika miss di mana-mana, satu aktivasi dan empat pembacaan kolom di chip DRAM. Warp kami tinggal di salah satu dari empat sub-partisi SM bersama sebelas warp residan lainnya. Instruksi diterbitkan ke unit load/store (LSU), yang mengirimkan opcode, masker lane aktif, alamat yang dihitung, dan nomor register ke coalescer.
Coalescer menentukan bahwa 4 permintaan sektor berurutan diperlukan untuk 128 byte yang diminta warp. Memasuki cache L1, permintaan dikirim ke cache yang dialamatkan secara virtual, set-associative 4-arah. Pencarian membandingkan tag dalam set empat slot. Karena ini adalah beban pertama kali, ia miss dan turun lebih jauh. Hit L1 kembali dalam sekitar 15,4 ns — 40 siklus. Memori virtual memerlukan terjemahan sebelum meninggalkan SM. TLB SM menyimpan enam belas terjemahan terbaru. Miss berbiaya sekitar 4,4 ns — sebelas siklus.
Setelah terjemahan, satu permintaan per baris 128-byte melintasi crossbar menuju salah satu dari 36 irisan L2 2 MiB, dipilih oleh fungsi alamat fisiknya. Semua irisan dapat melayani secara paralel, memberikan bandwidth agregat 36 kali lipat dari satu irisan.
Sumber: Hacker News · Diringkas oleh HeadlinesBriefing