c++ - 在CUDA中的表现

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

因此，最初，我有两个内核都将全局内存作为uchar3和float3以及uchar4等加载。但是，由于合并，我不太确定要使用三元组，因此我认为应该对其进行性能分析。我认为，由于操作与通道号无关，因此我可以很好地读取图像，就好像它是宽度为原来的三倍的1通道uchar图像一样，而不是实际的uchar3图像。

实际上，uchar3全局加载要比uchar加载慢得多。我的努力被证明是正确的。但是，a，这仅在算术内核中发生。按位与运算显示完全相反的结果!

现在，我知道我可以将图像数据加载为uint而不是uchar，以进行按位操作，这应该很好地进行合并。但是，假设我只是想学习和了解正在发生的事情。

并且让我们忘记float3和float4等。我的问题是内核的uchar版本。因此，简而言之，为什么uchar加载有时比uchar3加载快，而有时却没有？

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

PS。根据CUDA编程指南，逻辑操作和添加操作具有相同的吞吐量。 (实际上，我的内核必须首先将uchar转换为uint，但这应该在两个内核中都发生。)因此，从我收集的数据来看，执行长度应大致相同。

算术添加内核(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内核类似。 (老实说，我不确定我确切地记得内核。明天我将确认)。

最佳答案

编译器将uchar3负载分为单独的负载，因为SM的指令集中没有24位负载。因此，它们永远不会合并。在某种程度上，缓存将减轻这种情况。

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

您可以比较两种版本的cuobjdump -sass可执行文件的输出，以查看编译器的静态调度是否相同。但是，运行时动态调度的工作原理基本上是未知的。

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