0

这是 CUDA SDK (2.3) matrixMultiply 内核的一部分：

for (int a = aBegin, b = bBegin;
         a <= aEnd;
         a += aStep, b += bStep) {

    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

    int XI=wA * ty + tx;
    int XII=wB * ty + tx;
    ////////////////////
    // PREFETCH BLOCK //
    ////////////////////
    AS(ty, tx) = A[a + XI];
    BS(ty, tx) = B[b + XII];

    __syncthreads();

    for (int k = 0; k < BLOCK_SIZE; ++k)
        Csub += AS(ty, k) * BS(k, tx);

    __syncthreads();
}

此版本的矩阵乘法将切片带入共享内存并在共享内存带宽上执行计算。我想通过将下一次迭代的数据预取到 L1 缓存中来提高性能。我按照这里的建议使用预取内在函数，并将以下命令插入到PREFETCH BLOCK上面：

    long long int k,kk;
    k=((long long int)A+aStep); if(k<=aEnd) prefetch_l1(k+XI);
    kk=((long long int)B+bStep); if(kk<=aEnd) prefetch_l1(kk+XII);

经过测试，两个版本（有或没有预取）执行非常相似（平均 3 次运行）：

没有预取： 6434.866211（毫秒）

预取： 6480.041016（毫秒）

问题：

我希望看到预取的一些加速，但我对结果感到困惑。任何机构都有理由说明为什么这两个实现非常接近？也许我正在执行错误的预取。

先感谢您。

更多信息：

显卡：特斯拉 C2050

CUDA 版本：4.0

inline __device__ void prefetch_l1 (unsigned int addr)
{

  asm(" prefetch.global.L1 [ %1 ];": "=r"(addr) : "r"(addr));
}
4

1 回答 1

1

预取（在任何架构上）仅在以下情况下才有用：

  • 你有空闲的内存带宽和
  • 您可以在正确的时间启动预取，即在实际需要数据之前足够远的时间以及在空闲内存带宽可用的时间。

如果您不能满足这些标准，那么预取将无济于事，甚至可能弊大于利。

于 2012-12-02T22:05:10.727 回答