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

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 类似，可以访问任意的地址空间。