将原始图像为RGB的数据转存为RGBA的

这个问题和上个问题是一个项目里的,有采集到的原始的3字节一个点的,每个颜色分量为1B的彩色图像,想转存为RGBA的。我的代码如下,想问下这个代码是否需要有效率上的问题能改进
global void devGetRGBA(unsigned char* pRGB, int nRGBPitch, unsigned char* pRGBA, int nRGBAPitch)
{
int x=threadIdx.x+blockIdx.xblockDim.x;
int y=threadIdx.y+blockIdx.y
blockDim.y;

//unsigned char *ptrRGB, *ptrRGBA;
uchar4 *ptrRGB;
uchar4 *ptrRGBA;

ptrRGB= (uchar4 )(pRGB+x3+ynRGBPitch);
ptrRGBA=(uchar4 )(pRGBA+x4+y
nRGBAPitch);

(*ptrRGBA).x = (*ptrRGB).z;
(*ptrRGBA).y = (*ptrRGB).y;
(*ptrRGBA).z = (*ptrRGB).x;
(*ptrRGBA).w = 255;
}

这个算法估计没啥能改进的了。。。这不就是读取,简单赋值调换顺序/加上alpha分量, 回写到内存/显存的典型的访存密集型代码吗。

你读取3B, 回写4B, 中间却只需要几条指令简单变化而已,却有7B的访存。显然是怎么写都几乎卡死在你的卡的读写global memory的硬件极限上。我不认为将里面的神马操作(例如神马更高效的赋值之类)改进下,提高运算效率,会对你的已经达到极限的访存有好处。。

此外,如果你真的需要改进你的算法的话,例如:

你想修改你的
.x = .z;
.y = .y;
.z = .x;
.w = 255;
的操作

可以考虑fermi/kepler提供的permute操作。该操作是特别为从BGR到RGB或者RGBA这样的快速变化而优化的。

permute可以单条实现你的上面的1条固定的(.y = .y;) 2条交换的(.x = .z; .z = .x;)和1条固定的分量值赋值(.w = 255;)。 从而可能减少你的指令数(4条系列赋值=>1条permute), 从而可以在虽然可能无法提高你的程序性能(卡在访存上),但可以让你少执行的指令总数,从而给你省下了电费。

如果要用这个permute来快速变换,怎么用呢?有例子吗?