Performance in CUDA

Question

我似乎无法弄清楚影响内核性能的根本因素。我实现了两个简单的内核，一个是加载两个图像并逐像素相加，另一个是在逐位意义上加载两个图片并ANDS。现在，我已经对它们进行了模板化，这样内核就可以拍摄8和32位图像，以及1、3和4通道图像。

因此，最初我让两个内核将全局内存加载为uchar3和float3，以及uchar4等。然而，由于合并，我不太确定是否使用三元组，所以我想对其进行评测。我想，由于操作与通道编号无关，所以我可以像读取宽度为三倍的1通道uchar图像一样读取图像，而不是真正的uchar3图像。

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

现在，我知道我可以将图像数据加载为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];
}
}

位AND内核：

__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内核类似。(老实说，我不确定我是否记得内核……我明天会确认的)。

Answer 1

uchar3加载由编译器拆分为单独的加载，因为SM的指令集中没有24位加载。因此，它们从未结合在一起。在某种程度上，缓存将减轻这种情况。

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

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

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