是否可以使用 perf 事件获取 GPU 内部缓存的信息？

Question

我正在使用 perf 事件来获取性能计数或缓存信息（例如缓存访问计数、缓存未命中计数）。 现在，我想获取 GPU 的缓存信息。但是，问题是 perf 事件是否可以获取 GPU 的缓存信息。 我做了一个测试。

ioctl(fd, PERF_EVENT_IOC_RESET, 0);
ioctl(fd, PERF_EVENT_IOC_ENABLE, 0);

matrixMulCUDA<<< grid, threads >>> ( ... );

ioctl(fd, PERF_EVENT_IOC_RESET, 0);
ioctl(fd, PERF_EVENT_IOC_ENABLE, 0);

我确认数据似乎已被提取。 但我不能确定它是 GPU 内部的缓存信息。

有人知道吗？

http://man7.org/linux/man-pages/man2/perf_event_open.2.html（表演活动教程）

Answer 1

_{一个简短的版本：} No.

---

_{更长的版本：}为什么这是不可能的：

您的主机运行一个或两个 CPU 处理器_{（不涉及各自的CPU-cores）}，系统工具类似 @987654323 @ 可以在体内操作期间启动并收集一定数量的“自诊断遥测”，它在后台，在系统工作期间。

perf 语法支持所有必要的参数来指定临时哪个进程/CPU 来检查和记录数据。

虽然这可以作为 CPU 托管进程的临时观察者正常工作，但典型的 GPU 运行许多 _{（阅读 14 或更多)}Streaming Multiprocessor-s，叫SM，每个配备大约16-32-64个并行代码执行核心。

如果这是唯一的障碍，人们可以猜测主机操作系统可能有一些“备用”能力来嗅探它的少数进程，但过度订阅此能力很容易破坏稳定的操作仍然“在足迹内”其中说“备用电源”。

现在想象一个假设的情况，一个 GPU 能够并且会传送到你的操作系统 CPU-resources 几百倍大的入口这样的“性能遥测”数据，会让你闭上眼睛，依靠托管操作系统的相同“备用电源”以某种方式咀嚼它。

不，这不公平。 甚至没有尝试“吞下”来自多GPU 基础架构的数据流...

---

不过，万一呢？

虽然可能需要朝这个方向做一些疯狂的事情，但每个 SM 都有自己的一组 L1 / L2 / texL 缓存资源并假设一个设计良好的 TLP-架构将允许一些“备用电源”-in-GPU花费在SM-diagnostics 开销上，您必须以非常明智的方式自己编写代码，以限制 WARP-divergence 对 GPU 内核代码执行的副作用。虽然您可以通过GPU-code 的针孔通过DMA/RDMA/pinned-memory alchemy 将一些有限的数据移动到主机地址空间，但GPU-realm 将不允许您“读取”来自SM[13] 的遥测数据，来自SM[14] 等人。

---

一个好消息是，...

为了至少获得一些关于GPU-hardware-architecture 如何让您的特定GPU-kernel-code 与特定GPU-hardware 资源对齐的见解或原始视图， GPU-vendor 提供了一组工具，其中一个可以详细说明 SM 缓存使用/缓存溢出 _{的模拟期望 （详细说明您的主机操作系统，在您将特定的GPU-kernel-code 编译成目标GPU-hardware 汇编语言期间——这可以帮助您设计/微调/编译器优化最终的GPU-kernel - 根据您的优先级编写代码 - 最小化延迟，最大化处理速度，保持最高 TLP 和 ILP 并行度 WARP-code-executions ）}

---

如果仍有疑问，...

检查您可以从 GPU-fabric 连接到多少个片上 JTAG 诊断连接器。

无...