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

加速计算 CUDA
1

这个问题和上个问题是一个项目里的，有采集到的原始的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.yblockDim.y;

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

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

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

2

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

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

3

此外，如果你真的需要改进你的算法的话，例如：

你想修改你的
.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), 从而可以在虽然可能无法提高你的程序性能（卡在访存上），但可以让你少执行的指令总数，从而给你省下了电费。

4

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