高全局内存指令开销 - 不知道它来自哪里答案

【问题标题】：High global memory instruction overhead - no idea where it comes from高全局内存指令开销 - 不知道它来自哪里
【发布时间】：2013-04-06 06:47:09
【问题描述】：

我编写了一个内核，用于计算给定 D 维向量 q（存储在常量内存中）和 N 个向量（也是 D 维）的数组 pts 之间的欧几里得距离。

内存中的数组布局是这样的，前 N 个元素是所有 N 个向量的第一个坐标，然后是 N 个第二个坐标的序列，依此类推。

这是内核：

__constant__ float q[20];

__global__ void compute_dists(float *pt, float *dst,
        int n, int d) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; 
             i += gridDim.x * blockDim.x) {
        float ld = 0;
        for (int j = 0; j < d; ++j) {
            float tmp = (q[j] - pts[j * n + i]); 
            ld += tmp * tmp;
        }
        dst[i] = ld;
    }
}r

调用如下：

const int N = 1000000, D = 20;
compute_dists<<<32, 512>>>(vecs, dists, vec, N, D);

现在，在 Quadro K1000M 上使用 NVIDIA Visual Profiler 分析此内核会导致警告

高指令重放开销 (31.2%) 和
高全局内存指令开销 (31.2%)。

这对我来说非常令人惊讶，因为据我所知，内存访问是合并的（因为j * n + i 始终是线程中第一个扭曲的32 的倍数，这为我们提供了 128 字节对齐）并且没有分支分歧..

是否还有其他一些因素会影响指令重播开销指标，或者我是否遗漏了其他内容？

【问题讨论】：

请发布完整的可复制的。除了地址分歧之外，还有几个原因会导致内存指令重放。剖析器不提供有关其他原因的信息。 Nsight VSE CUDA Profiler Memory Transaction 实验可以在源码视图中显示每条内存指令的事务。
我已经很久没有这样做了。所以我可能完全关闭了，但为什么最里面的循环是尺寸而不是线程索引。这不会导致你无缘无故地切换内存流吗？
您发布的代码中有足够多的语法错误表明这不是您实际运行的代码。您能否发布您所询问的实际代码？

标签： memory cuda gpu

【解决方案1】：

我认为您有来自“pts[j * n + i]”的高 TLB（翻译后备缓冲区）未命中率的问题。由于 n 很大，连续的第 j 个元素很可能不会出现在加载的内存页面中。 TLB 硬件在加载给定内存位置的页面所在的信息时具有很高的延迟。这会导致内存加载指令重播。如果缓存中不存在数据或页面未加载到 TLB，则重新发出每个内存加载指令。尽管我不完全确定后者，但可能就是这种情况。希望能帮助到你。我有同样的问题，但更严重，97% 重播。 My question 也可能会回答你的问题。

【讨论】：