我似乎无法弄清楚影响内核性能的潜在因素。我实现了两个简单的内核,一个加载两个图像并逐个像素地添加它们,另一个加载两个图像并对它们进行按位运算。现在,我对它们进行了模板化,以便内核可以获取 8 位和 32 位图像,以及 1、3 和 4 通道图像。
所以,最初我让两个内核都将全局内存加载为uchar3
和float3
,以及uchar4
等等。但是,由于合并,我不太确定使用三元组,所以我想我会给它一个分析运行。我想,由于操作与通道数无关,我可以像读取图像一样读取图像,就好像它是uchar
宽度为三倍的 1 通道图像,而不是uchar3
真正的图像。
事实上,uchar3
全局负载比负载慢得多uchar
。我的努力得到了证明。但是,唉,这只发生在算术内核上。按位与运算显示完全相反的结果!
现在,我知道我可以将图像数据加载为uint
s 而不是uchar
s,用于按位运算,这应该可以完美地处理合并。但是让我们假设我只是想学习和理解正在发生的事情。
让我们忘记float3
s 和float4
s 等。我的问题在于uchar
内核的版本。那么,简而言之,为什么uchar
负载有时比负载快uchar3
,有时又不是呢?
我使用的是 GTX 470,计算能力 2.0。
PS。根据 CUDA 编程指南,逻辑运算和加法运算具有相同的吞吐量。(我的内核实际上必须首先将uchar
s 转换为uint
s,但这应该在两个内核中都发生。)因此,根据我的收集,执行长度应该大致相同。
算术添加内核(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 内核也是如此。(老实说,我不确定我是否记得确切的内核......我会在明天确认)。