L1 缓存对于矩阵乘法命中率为 0%

Question

在 GPU 上进行大型矩阵乘法运算。内核分析显示 L1 缓存命中率为 0%。矩阵步长超过 4k 字节，缓存行大小为 128 字节。 L1 缓存命中的可能原因是什么？ 谢谢

Answer 1

L1 缓存命中的可能原因是什么？

CUDA 中的加载和存储都不经过 L1 缓存。

具体来说，

• 编译器可能更喜欢使用非连贯只读缓存通过受限制的 const 指针进行加载；请参阅此处关于 SO 或 CUDA 的这个答案。这对应于

  ld.global.nc

  PTX 指令；您可能还会看到使用 CUDA C++

  __ldg()

  内在完成的操作。
• 在某些情况下，编译器会发出带有

  .cg

  修饰符的存储指令，这意味着“全局缓存”，即在 L2 级别缓存，而不是在 L1 级别缓存。请参阅PTX 存储指令缓存首选项表。

当这两种情况都发生时，L1 缓存利用率确实应该为 0。