前回の記事では、ベクトル加算カーネルをnvccからワープまで追跡しました。今回は、重要なSASS命令(グローバルロード)をRTX4090のハードウェアを通じて追跡します。調査対象のCUDAカーネルは2行です:__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は次を示します:/*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]。これらの命令は、ベクトルbの要素をグローバルメモリからレジスタにロードします。
1つのLDG.Eは、32のレーンのそれぞれで4バイトを要求します。それに応答するには、4つの32バイトセクター、1つのキャッシュライン、1回のアドレス変換、1回のクロスバークロッシング、36個のL2スライスのうちの1つ、そしてすべてでミスした場合は、DRAMチップでの1回のアクティベートと4回のカラムリードが必要です。私たちのワープは、SMの4つのサブパーティションのうちの1つに、他の11個の常駐ワープと共に存在します。命令はロード/ストアユニット(LSU)に発行され、LSUはオペコード、アクティブレーンマスク、計算されたアドレス、レジスタ番号をコアレッサーに送ります。
コアレッサーは、ワープが要求した128バイトに対して4つの連続したセクター要求が必要であると判断します。L1キャッシュに入ると、要求は仮想アドレス指定の4ウェイセットアソシアティブキャッシュに送られます。ルックアップは4スロットのセットでタグを比較します。これが初回のロードであるため、ミスし、さらに下へ降ります。L1ヒットは約15.4ナノ秒で返ります — 40サイクル。仮想メモリは、SMを離れる前に変換を必要とします。SMのTLBは最近の16個の変換を保持します。ミスのコストは約4.4ナノ秒 — 11サイクルです。
変換後、128バイトラインごとに1つの要求がクロスバーを横断して、物理アドレスの関数によって選択された36個の2 MiB L2スライスのうちの1つへ進みます。すべてのスライスは並列に応答できるため、合計帯域幅は単一スライスの36倍になります。
出典: Hacker News · 要約：HeadlinesBriefing