1

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

内存中的数组布局是这样的，前 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始终是32128 字节对齐的倍数）并且没有分支分歧。

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

4

1 回答 1

1

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

于 2013-06-08T23:01:27.190 回答