Nvidia 视觉分析器中 L1 本地存储/L1 全局读取的含义答案

【问题标题】：Meaning of L1 local store / L1 global read in Nvidia visual profilerNvidia 视觉分析器中 L1 本地存储/L1 全局读取的含义
【发布时间】：2016-06-21 13:00:38
【问题描述】：

__global__ void kern (int* gpuArr) {
    int a;

    a = gpuArr[threadIdx.x+1];
}

int main(char** argv, int argc) {
    int* gpuArr;
    cudaMalloc((void**)&gpuArr,628*sizeof(int));

    kern<<<1,32>>>(gpuArr);

    cudaDeviceSynchronize();
}

在 Nvidia 视觉分析器中分析上述代码时，我得到以下内存带宽分析。

据我所知，发生的事情是：

来自全局的内存负载 - L2 未命中

从全局复制5*32B到L2（我认为L2的缓存行大小总是32B）

从 L2 复制 2*128B 到 L1（包括 L2 的附加数据，因为 L1 缓存行是 128B？）

对每个线程执行相当于 a = L1_position[threadIdx.x] 的操作

那么每个线程中的局部变量“a”存储在哪里？根据调试器，它不在寄存器中（通常在变量选项卡中显示为@register int，但它显示为@local int）。 “Local Stores 1”这一行到底是什么意思？我们已经做了 2*128B L2 到 L1 的复制（“Global Loads”），那么“Local Store”是什么意思。为什么有 2 个“全球商店”却只有 1 个“本地商店”？

另外，我使用的是 cc 3.0，因此根据规范，L1 中的全局内存缓存是不可能的

【问题讨论】：

您是否正在使用该代码进行测量？您的代码的问题是 a = gpuArr[threadIdx.x+1]; 将被完全优化掉，因为它不会改变全局状态。
我正在使用那个确切的代码，我可以调试它，所以我认为它没有被优化掉
@Obabopisamon：打开调试时，编译器优化被禁用。但否则该内核将被编译为 NULL 存根

标签： cuda

【解决方案1】：

首先，正如已经指出的那样：

您正在运行已使用调试开关 (-G) 编译的代码。这不会为您提供最佳性能，并且不代表没有它编译的代码（无论是性能还是行为），因此分析此类代码是一个值得怀疑的活动。
由于-G 禁用编译器优化，此类代码的行为可能与您预期的不同。

那么每个线程中的局部变量“a”存储在哪里？

它存储在本地逻辑空间中。这正是分析器（和调试器）告诉你的。 “本地”逻辑空间既可以存在于寄存器中，也可以存在于物理（板载 DRAM）内存中。它不在寄存器中的原因是因为您使用-G 开关禁用了优化，并且将逻辑数据放置在寄存器中是一种优化。您将无法通过消除-G 开关直接确认这一点，因为如果您这样做了，您编写的代码将被编译器完全优化掉，因为它对任何全局状态都没有影响。

“本地商店 1”这一行究竟是什么意思？

如上所述，变量a在本地逻辑空间中，因此从全局读取和向本地写入发生在这里：

a = gpuArr[threadIdx.x+1];

当a 被“写入”时，将导致本地存储。

为什么有 2 个“全球商店”却只有 1 个“本地商店”？

存储在 DRAM 内存中的属于本地逻辑空间的变量将以这样的方式存储在内存中，即当进行 DRAM 内存事务时，warp 中的线程的连续访问将产生相邻（即“合并”）访问，读取或写入此类值。这意味着如果我为每个线程都有一个局部变量a，那么将首先存储线程0 的a，然后存储线程1 的a，然后存储线程2 的a，以此类推，这样如果每个线程读取（或写入）a，结果访问将被合并。由于您恰好有 1 个 32 个线程的 warp，每个线程写入一个 int 值 a，这会导致单个 128 字节的本地存储事务。

在您的全局内存情况下（即读取），您已经导致您的读取跨过缓存线/段边界，并将 1 添加到您的数组索引中：

a = gpuArr[threadIdx.x+1];
                       ^

所以它需要两个全局事务来收集 warp 请求的数据。如果您想确认这一点，请消除数组索引上的 +1，全局事务应该从 2 下降到 1。

需要注意的是，分析非常少量的活动可能并不总能得到您期望的结果（尽管它似乎在这种情况下有效）。原因是为某些 SM 子集捕获了一些分析器指标，然后乘以 SM 的数量以反映完整的 GPU 活动。如果非常小的数据集的结果没有意义，您可能会为更大的数据集获得更合理的结果，这些数据集更恰当地“填充”GPU 与跨 SM 的一致活动。

【讨论】：