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