1

我正在开发一些 CUDA 程序，我想使用常量内存来加快计算速度，但事实证明，使用常量内存会使我的代码慢约 30%。

我知道恒定内存擅长将读取广播到整个经线，我认为我的程序可以利用它。

这是常量内存代码：

__constant__ float4 constPlanes[MAX_PLANES_COUNT];

__global__ void faultsKernelConstantMem(const float3* vertices, unsigned int vertsCount, int* displacements, unsigned int planesCount) {

    unsigned int blockId = __mul24(blockIdx.y, gridDim.x) + blockIdx.x;
    unsigned int vertexIndex = __mul24(blockId, blockDim.x) + threadIdx.x;

    if (vertexIndex >= vertsCount) {
        return;
    }

    float3 v = vertices[vertexIndex];
    int displacementSteps = displacements[vertexIndex];

    //__syncthreads();

    for (unsigned int planeIndex = 0; planeIndex < planesCount; ++planeIndex) {
        float4 plane = constPlanes[planeIndex];
        if (v.x * plane.x + v.y * plane.y + v.z * plane.z + plane.w > 0) {
            ++displacementSteps;
        }
        else {
            --displacementSteps;
        }
    }

    displacements[vertexIndex] = displacementSteps;
}

全局内存代码是相同的，但它多了一个参数（带有指向平面数组的指针）并使用它而不是全局数组。

我以为那些第一个全局内存读取

float3 v = vertices[vertexIndex];
int displacementSteps = displacements[vertexIndex];

可能会导致线程的“去同步化”，然后它们不会利用持续内存读取的广播，所以我尝试调用 __syncthreads(); 在读取常量内存之前，但它没有改变任何东西。

怎么了？提前致谢！

系统：

  • CUDA 驱动程序版本：5.0
  • CUDA 能力：2.0

参数：

  • 顶点数：~250 万
  • 飞机数量：1024

结果：

  • 常量内存版本：46 毫秒
  • 全局内存版本：35 毫秒

编辑：

因此，我尝试了很多方法来使常量内存更快，例如：

1）注释掉两个全局内存读取，看看它们是否有任何影响，它们没有。全局内存仍然更快。

2) 每个线程处理更多的顶点（从 8 个到 64 个）以利用 CM 缓存。这甚至比每个线程一个顶点还要慢。

2b) 使用共享内存来存储位移和顶点 - 在开始时加载所有这些，处理并保存所有位移。同样，比显示的 CM 示例慢。

经过这次经历，我真的不明白 CM 读取广播是如何工作的，以及如何在我的代码中正确“使用”。这段代码可能无法用 CM 优化。

编辑2：

又一天的调整，我试过：

3）每个线程处理更多的顶点（8到64），内存合并（每个线程的增量等于系统中的线程总数）——这比增量等于1提供更好的结果，但仍然没有加速

4) 替换这个 if 语句

if (v.x * plane.x + v.y * plane.y + v.z * plane.z + plane.w > 0) {
    ++displacementSteps;
}
else {
    --displacementSteps;
}

这是用一点点数学来给出“不可预测”的结果，以避免使用以下代码进行分支：

float dist = v.x * plane.x + v.y * plane.y + v.z * plane.z + plane.w;
int distInt = (int)(dist * (1 << 29));  // distance is in range (0 - 2), stretch it to int range
int sign = 1 | (distInt >> (sizeof(int) * CHAR_BIT - 1));  // compute sign without using ifs
displacementSteps += sign;

不幸的是，这比使用 if so ifs 并没有我想象的那么大的邪恶要慢很多（~30%）。

编辑3：

我得出这个问题的结论是，这个问题可能无法通过使用常量内存来改善，这些是我的结果*：

全局和恒定内存性能图

*时间报告为 15 次独立测量的中值。当常量内存不足以保存所有平面（4096 和 8192）时，内核被多次调用。

4

1 回答 1

2

虽然计算能力 2.0 芯片有 64k 的常量内存，但每个多处理器只有 8k 的常量内存缓存。您的代码的每个线程都需要访问所有 16k 的常量内存，因此您会因缓存未命中而损失性能。为了有效地为平面数据使用常量内存，您需要重新构建您的实现。

于 2013-03-06T07:33:13.933 回答