内存访问使 CUDA 内核非常慢

Question

我正在尝试使用 cuda 来制作一个基本的片段着色器，我发现实际执行内核需要超过一秒钟，这对于我试图实时运行的着色器来说是不可接受的。我发现使用 synchronize 方法并通过评论一些内核，它是对输出数组的内存访问导致它如此缓慢的原因。我还没有真正尝试过任何解决问题的方法，因为我什至不知道从哪里开始。这是在 PyCUDA 中，我认为这并不重要，但这是内核代码：

__global__ void fragment_shader(int palette_lim,float *palette, float *input, float *output) {
    int fragment_idx = (3*gridDim.y*blockIdx.x)+(3*blockIdx.y);
    float min_dist = sqrtf(3);
    float color_dist;
    int best_c = 0;
    for (int c=0;c<palette_lim;c++) {
        color_dist = sqrtf(pow(input[fragment_idx]-palette[c*3],2)+pow(input[fragment_idx+1]-palette[c*3+1],2)+pow(input[fragment_idx+2]-palette[c*3+2],2));
        if (color_dist < min_dist) {
            min_dist = color_dist;
            best_c = c;
        }
    }

    //These are the lines that make it slow. If these lines get commented out, it runs in a time that would be acceptable
    output[fragment_idx] = palette[best_c*3];
    output[fragment_idx+1] = palette[best_c*3+1];
    output[fragment_idx+2] = palette[best_c*3+2];
}

编辑：在多玩了一会儿之后，我发现它也与分配给输出数组的内容有关，因为当我让它写一些常量而不是调色板中的东西时它也工作得很好，它只是当时没有做任何有用的事情。

Answer 1

首先对你的实际计算做一些评论：

• 你比较

  sqrtf(x) < sqrtf(3)

  。根很贵。比较一下

  x < 3.f

• 即使你想保留平方根以避免溢出浮动范围（可能不是问题），也不要使用

  sqrt(pow(x, 2)+...)

  ，因为那个问题不要只使用

  pow

  来进行平方。使用

  hypotf

  用于 2D 或

  norm3df

  用于 3D 矢量
• 您将最后一个低于限制的值保存在调色板中。看起来很像你想要适合最好的颜色

现在让我们分析一下你的内存访问：

片段索引

让我们看看

fragment_idx = 3*gridDim.y*blockIdx.x+3*blockIdx.y

：你没有考虑

threadIdx.x

和

threadIdx.y

。 这是你的主要问题：许多线程作用于相同的输入和输出数据。你可能想要这个：

fragment_idx = 3 * (threadIdx.y * blockDim.x + threadIdx.x)

输入

所以你加载了 3 个浮点数。对于初学者，当它不依赖于循环迭代时，为什么要在循环中重新加载它？我假设编译器将您从该访问中拯救出来，但不要养成这样做的习惯。

其次，您的访问模式未正确合并，因为 a) 这些是 3 个独立访问，b) CUDA 无法合并对

float3

向量的访问，即使您做得正确。请阅读最佳实践指南的9.2.1 合并访问全局内存部分。为了获得更好的性能，您有两种选择：

1. 您为每个

   fragment_idx

   添加1个浮点数，这样您就可以将整个东西加载为

   float4

2. 您将输入数组从 Nx3 矩阵转置为 3xN 矩阵

调色板

访问 3 个浮点数的问题相同。另外，现在每个线程都读取相同的值，因为

c

不依赖于线程索引。访问最起码要经过

__ldg

函数才能使用一级缓存。最好将调色板预取到共享内存中

输出

写访问与未合并访问有同样的问题。另外，由于

best_c

因线程而异，因此对

palette

的读取访问是随机的。您必须在循环之前加载

palette

值。只需将最佳调色板值保存在局部变量中，然后重新使用它来存储最后的输出。

方法论

两点备注：

1. 在快速运行之前尝试使您的代码有效。那会抓住

   fragment_idx

2. 如果你简化代码，比如删除输出，编译器会很乐意优化你的大部分代码。这不是您进行适当绩效评估的方式。使用分析器。 CUDA自带非常好的

最小修复

这是纠正问题的最简单的代码。它没有解决加载 vector3 变量的问题，也没有使用共享内存。这需要更多的改变

__device__ float sqr_norm(float3 a, float3 b) {
    a.x -= b.x, a.y -= b.y, a.z -= b.z;
    return a.x * a.x + a.y * a.y + a.z * a.z;
}
__global__ void fragment_shader(int palette_lim,
          const float *palette, const float *input,
          float *output) {
    int fragment_idx = 3 * (threadIdx.y * blockDim.x + threadIdx.x);
    /* TODO: Switch to float4 for better memory access patterns */
    float3 inputcolor = make_float3(
          input[fragment_idx], input[fragment_idx + 1], input[fragment_idx + 2]);
    float min_dist_sqr = 3.f;
    /* The old code always used color index 0 if there was no fit */
    float3 best_color = make_float3(
          __ldg(palette), __ldg(palette + 1), __ldg(palette + 2));
    float best_dist = sqr_norm(best_color, inputcolor);
    for(int c = 1; c < palette_lim; c++) {
        /* TODO: Prefetch into shared memory */
        float3 color = make_float3(
              __ldg(palette + c), __ldg(palette + c + 1), __ldg(palette + c + 2));
        float dist = sqr_norm(color, inputcolor);
        /* Since we always used color 0 in the old code,
         * the min_dist is somewhat pointless */
        if(dist < min_dist_sqr && dist < best_dist) {
            best_color = color;
            best_dist = dist;
        }
    }
    output[fragment_idx] = best_color.x;
    output[fragment_idx + 1] = best_color.y;
    output[fragment_idx + 2] = best_color.z;
}