CUDA 中的性能答案

【问题标题】：Performance in CUDACUDA 中的性能
【发布时间】：2012-11-29 09:59:37
【问题描述】：

我似乎无法弄清楚影响内核性能的潜在因素。我实现了两个简单的内核，一个加载两个图像并逐像素添加它们，另一个加载两个图像并ANDS它们，在按位意义上。现在，我已经对它们进行了模板化，以便内核可以获取 8 位和 32 位图像，以及 1、3 和 4 通道图像。

所以，最初我让两个内核将全局内存加载为 uchar3 和 float3，以及 uchar4 等。但是，由于合并，我不太确定使用三元组，所以我想我会给它一个分析运行。我认为由于操作与通道号无关，因此我可以像读取图像一样读取图像，就像它是宽度为三倍的 1 通道 uchar 图像，而不是真正的 uchar3 图像。

事实上，uchar3 全局加载比uchar 加载慢得多，多。我的努力得到了证明。但是，唉，这只发生在算术内核上。按位与运算显示完全相反的结果！

现在，我知道我可以将图像数据加载为uints 而不是uchars，用于按位运算，这应该可以完美地处理合并。但是让我们假设我只是想学习和理解正在发生的事情。

让我们忘记float3s 和float4s 等。我的问题在于内核的uchar 版本。那么，简而言之，为什么 uchar 加载有时比 uchar3 加载快，而有时则不然？

我使用的是 GTX 470，计算能力 2.0。

PS。根据 CUDA 编程指南，逻辑运算和加法运算具有相同的吞吐量。（我的内核实际上必须首先将uchars 转换为uints，但这应该在两个内核中都发生。）所以根据我的收集，执行长度应该大致相同。

算术加法内核（uchar版本）：

__global__ void add_8uc1(uchar* inputOne, uchar* inputTwo, uchar* output, unsigned int width, unsigned int height, unsigned int widthStep)
{
    const int xCoordinateBase = blockIdx.x * IMAGE_X * IMAGE_MULTIPLIER + threadIdx.x;
    const int yCoordinate = blockIdx.y * IMAGE_Y + threadIdx.y;

    if (yCoordinate >= height)
        return;

#pragma unroll IMAGE_MULTIPLIER
    for (int i = 0; i < IMAGE_MULTIPLIER && xCoordinateBase + i * IMAGE_X < width; ++i)
    {
        //  Load memory.
        uchar* inputElementOne = (inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x));
        uchar* inputElementTwo = (inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x));

        //  Write output.
        *(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)) = inputElementOne[0] + inputElementTwo[0];
    }
}

按位与内核：

__global__ void and_8uc1(uchar* inputOne, uchar* inputTwo, uchar* output, unsigned int width, unsigned int height, unsigned int widthStep)
{
    const int xCoordinateBase = blockIdx.x * IMAGE_X * IMAGE_MULTIPLIER + threadIdx.x;
    const int yCoordinate = blockIdx.y * IMAGE_Y + threadIdx.y;

    if (yCoordinate >= height)
        return;

#pragma unroll IMAGE_MULTIPLIER
    for (int i = 0; i < IMAGE_MULTIPLIER && xCoordinateBase + i * IMAGE_X < width; ++i)
    {
        //  Load memory.
        uchar* inputElementOne = (inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x));
        uchar* inputElementTwo = (inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x));

        //  Write output.
        *(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)) = inputElementOne[0] & inputElementTwo[0];
    }
}

uchar3 版本相同，只是加载/存储行现在如下：

        //  Load memory.
    uchar3 inputElementOne = *reinterpret_cast<uchar3*>(inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3);
    uchar3 inputElementTwo = *reinterpret_cast<uchar3*>(inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3);

    //  Write output.
    *reinterpret_cast<uchar3*>(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3) 
        = make_uchar3(inputElementOne.x + inputElementTwo.x, inputElementOne.y + inputElementTwo.y, inputElementOne.z + inputElementTwo.z);

对于 AND 内核也是如此。（老实说，我不确定我是否记得确切的内核......我会在明天确认）。

【问题讨论】：

你能给我们看看你的内核吗？对未知代码进行推理有点困难。或许可以告诉我们执行时间的比较（即内核在一个版本或另一个版本上花费的时间是否相似，更快，...）。
你能提出一个简洁的问题吗？您想知道为什么加载一堆uchar 会比加载一堆uchar3 更快吗？如果未合并的加载/存储是一个问题，CUDA 5 中的分析器会发出通知，即使对于最基本的分析运行类型也是如此。它对这 2 个案例的百分比表示了什么？
负载是线性的并且完全合并，至少在 uchar 的情况下是这样。我无法粘贴内核，因为我现在不在工作。我想知道为什么它有时更快，有时却没有。

标签： c++ performance cuda

【解决方案1】：

uchar3 加载由编译器拆分为单独的加载，因为 SM 的指令集中没有 24 位加载。因此，它们永远不会合并。在一定程度上，缓存会缓解这种情况。

但是，根据具体的执行配置，每个线程可能只有大约 10.7 字节的缓存（您的示例可能会接近该值，因为内核很简单，因此许多线程可以在一个 SM 上同时运行） .由于缓存不是完全关联的，因此在发生抖动之前每个线程的可用字节数可能要小得多。具体何时发生取决于许多因素，包括指令的确切调度，即使对于具有相同记录吞吐量的指令，这也可能有所不同。

您可以比较两个版本的cuobjdump -sassexecutable 的输出，看看编译器的静态调度是否相同。然而，运行时的动态调度如何工作基本上是不可观察的。

您已经注意到，图像的所有通道都以相同的方式处理，因此您如何在线程之间分配它们并不重要。最好的选择是使用uchar4 而不是uchar3 或uchar，这（假设图像的适当对齐）将为您提供独立于缓存的合并访问。这应该会导致执行时间更短且更一致。

【讨论】：