論文分享-《GPU Memory Exploitation for Fun and Profit》

Yin-SHT發表於2024-10-07
GPU

1. 研究問題

該論文對 NVIDIA GPU 上不同記憶體空間（global memory, local memory, shared memory）中存在的 buffer overflow 問題進行了深入的研究，併成功對在 GPU 上執行的 DNN 應用實現了 ROP 攻擊。

以往的研究侷限於單一記憶體空間中 buffer overflow 的影響，沒有對不同記憶體空間的跨越進行分析。另外 NVIDIA 自 2017 年釋出的 Volta architecture GPU 與以往的 GPU 架構有了很大不同，所以先前的研究不適用於當今的架構。

2. local memory 訪問方式解密

為了解析 thread 是如何訪問記憶體的，本文作者使用了非常暴力但有效的方法：利用 DMA dump 出 cuda 程式執行時的記憶體內容。在獲取整個 GPU device memory 之後，分析 memory 中的 data pattern 從而找到 thread 中 local memory 的實體記憶體位置。此外，本文作者還從 device memory 中提取了頁表的全部資訊。

2.1 local memory 訪問方式

__global__ void local_access() {
	uint32_t arr[10];
	for (int i = 0; i < 10; i ++) {
		arr[i] = 0xdead0000 + threadIdx.x;
	} 
	// 5 
	
int main() {
	local_access<<<1, 32>>>();
}

如上述程式碼所示，該 cuda 程式執行 local_access 核函式，該核函式包含一個 thread block，thread block 中包含 32 個 thread，每個執行緒會向自己的區域性私有陣列 arr 中寫入特定內容。
作者使用 cuda-gdb 使程式執行到第 5 行時暫停下來，並 dump 整個 device memory，結果如下。

如圖所示，一個執行緒塊中執行緒的記憶體空間會以四位元組為單位交錯排列。

此外，作者還發現 thread 對於 local memory 的訪問存在兩條執行路徑：

  • 當使用 LDL/STL 指令或者訪存地址字首為 0x7fff2 時，GPU 可識別出此時為 local memory 的訪問，此時的訪存地址在頁表中沒有有效的對映，GPU 會採取一條特殊的路徑來完成對記憶體的訪問，該路徑會將執行緒 ID 考慮在內。
  • 當使用頁表中的有效對映訪存時，其過程和 CPU 類似，可以訪問任意的地址空間。

相關文章