并行处理NvidiaOpenCLCUDAWin64Visual C++ 7.1Visual C++ 8.0Windows VistaVisual C++ 7.0Windows 7Win32Windows XPC中级开发WindowsC++
学习如何使用 CUDA 进行 Alpha 混合
使用 CUDA 的爆发性能进行图像处理
引言
本文将向您介绍使用 CUDA™ 进行 Alpha 混合的方法。什么是 CUDA？简而言之，它是由 NVIDIA® 开发的并行计算架构，是 NVIDIA 图形处理单元中的计算引擎，软件开发人员可以通过编程语言访问它。
背景
对于大图像，Alpha 混合需要更多的 CPU 时间来处理。从 NVIDIA 学习，CUDA 可以加速图像处理性能。嗯，您需要试一试，看看它们之间的差异有多大。此 Alpha 混合代码在 Windows 7 和 NVIDIA GeForce G210M 上进行了测试。
传统方式的 Alpha 混合
AlphaBlending_Host() 是我用来进行 Alpha 混合的传统例程。我认为它已经是速度最快的方式，并且也具有良好的性能。
参数
PULONG pulFore [in]：前景图像缓冲区。一般来说，前景图像在背景图像之上。PULONG pulBack [in]：背景图像缓冲区。PULONG pulResult [out]：混合后的图像缓冲区。在传递之前，必须先分配此图像缓冲区。DWORD dwSize [in]：图像尺寸大小。
void AlphaBlending_Host(PULONG pulFore , PULONG pulBack, PULONG pulResult, DWORD dwSize)
{
ULONG ulResult = 0L ;
ULONG ulAlpha = 0L ;
ULONG ulVal = 0L ;
ULONG ulFore = 0L ;
ULONG ulBack = 0L ;
if ((NULL == pulFore) || (NULL == pulBack) || (NULL == pulResult) || (0L == dwSize))
{
return;
}
_asm
{
FromBeginning:
//--- get foreground pixel and move buffer forward by 1 ---
MOV EAX,pulFore
MOV ECX,DWORD PTR [EAX]
MOV ulFore,ECX
MOV EDX,pulFore
ADD EDX,4
MOV pulFore,EDX
//--- get background pixel and move buffer forward by 1 ---
MOV EAX,pulBack
MOV ECX,DWORD PTR [EAX]
MOV ulBack,ECX
MOV EDX,pulBack
ADD EDX,4
MOV pulBack,EDX
// blend foreground color (F) to a background color (B),
// using alpha channel value of F
// Result Z = (Fa * Frgb) + ((1 - Fa) * Brgb)
// EAX <- Foreground
// EDX <- Background
MOV EAX,ulFore
MOV EDX,ulBack
// Test Fa = 255 ? operation of subtraction
CMP EAX,0xFF000000 // Fa = 255 ? => Result = EAX
JNC ReturnForePixel
// Test Fa = 0 ? operation of and
TEST EAX,0xFF000000 // Fa = 0 ? => Result = EDX
JZ ReturnBackPixel
JMP EntryProcess
ReturnForePixel:
MOV ulVal,EAX
JMP ProcessFinished
ReturnBackPixel:
MOV ulVal,EDX
JMP ProcessFinished
//--- entry ---
EntryProcess:
// Get weight W = Fa * M
MOV ECX,EAX // ECX <- Fa Fr Fg Fb
SHR ECX,24 // ECX <- 00 00 00 Fa
PUSH EBX
// P = W * F
MOV EBX,EAX // EBX <- Fa Fr Fg Fb
AND EAX,0x00FF00FF // EAX <- 00 Fr 00 Fb
AND EBX,0xFF00FF00 // EBX <- Fa 00 Fg 00
IMUL EAX,ECX // EAX <- Pr ** Pb **
SHR EBX,8 // EBX <- 00 Fa 00 Fg
IMUL EBX,ECX // EBX <- Pa ** Pg **
ADD EAX,0x00800080
AND EAX,0xFF00FF00 // EAX <- Pr 00 Pb 00
SHR EAX,8 // EAX <- 00 Pr ** Pb
ADD EBX,0x00800080
AND EBX,0xFF00FF00 // EBX <- Pa 00 Pg 00
OR EAX,EBX // EAX <- Pa Pr Pg Pb
// W = (1 - W) ; Q = W * B
XOR ECX,0x000000FF // ECX <- 1 - ECX
MOV EBX,EDX // EBX <- Ba Br Bg Bb
AND EDX,0x00FF00FF // EDX <- 00 Br 00 Bb
AND EBX,0xFF00FF00 // EBX <- Ba 00 Bg 00
IMUL EDX,ECX // EDX <- Qr ** Qb **
SHR EBX,8 // EBX <- 00 Ba 00 Bg
IMUL EBX,ECX // EBX <- Qa ** Qg **
ADD EDX,0x00800080
AND EDX,0xFF00FF00 // EDX <- Qr 00 Qb 00
SHR EDX,8 // EDX <- 00 Qr ** Qb
ADD EBX,0x00800080
AND EBX,0xFF00FF00 // EBX <- Qa 00 Qg 00
OR EBX,EDX // EBX <- Qa Qr Qg Qb
// Z = P + Q (assuming no overflow at each byte)
ADD EAX,EBX // EAX <- Za Zr Zg Zb
POP EBX
MOV ulVal,EAX // new blended RGB color
ProcessFinished:
//--- saved to result buffer and move buffer forward by 1 ---
MOV EAX,pulResult
MOV ECX,ulVal
MOV DWORD PTR [EAX],ECX
MOV EDX,pulResult
ADD EDX,4
MOV pulResult,EDX
//--- check next pixel until the final one ---
DEC dwSize
JNZ FromBeginning
}
}
在这种情况下，在 CPU 上的处理时间为 0.006027 (ms)。
使用 CUDA 进行 Alpha 混合
DilutePixel() 用于使用特定的 alpha 通道值混合像素，然后返回混合后的像素。请注意，此函数有一个 __device__ 关键字作为前缀。它由 CUDA __global__ 函数调用。
参数
unsigned long ulPixel [in]：源设备像素。unsigned long ulAlpha [in]：用于混合像素的 alpha 通道值。
该函数返回混合后的像素。
__device__
unsigned long DilutePixel(unsigned long ulPixel, unsigned long ulAlpha)
{
unsigned long nResult = 0;
nResult = ulPixel;
ulPixel &= 0x00ff00ff;
nResult &= 0xff00ff00;
ulPixel *= ulAlpha;
nResult >>= 8;
nResult *= ulAlpha;
ulPixel += 0x00800080;
ulPixel &= 0xff00ff00;
ulPixel >>= 8;
nResult += 0x00800080;
nResult &= 0xff00ff00;
nResult |= ulPixel;
return( nResult );
}
AlphaBlending_Texture() 是执行 Alpha 混合的线程例程，它处理前景和背景图像缓冲区的每个像素。请注意，此函数有一个 __global__ 关键字作为前缀。
参数
unsigned long* pResult [out]：混合后的设备图像缓冲区。unsigned nSize [in]：图像尺寸大小。
__global__
void AlphaBlending_Texture(unsigned long* pResult, unsigned nSize)
{
unsigned nIndex = (__umul24(blockIdx.x, blockDim.x) + threadIdx.x);
unsigned long ulPixelF = 0L;
unsigned long ulPixelB = 0L;
unsigned long ulAlphaF = 0L;
if (nIndex >= nSize)
{
return;
}
ulPixelF = tex1Dfetch(texForegnd, nIndex);
ulPixelB = tex1Dfetch(texBackgnd, nIndex);
ulAlphaF = (ulPixelF >> 24L);
if (ulAlphaF == 0xffL)
{
*(pResult + nIndex) = ulPixelF;
}
else if (ulAlphaF == 0L)
{
*(pResult + nIndex) = ulPixelB;
}
else
{
ulPixelF = DilutePixel(ulPixelF, ulAlphaF);
ulPixelB = DilutePixel(ulPixelB, (0xffL ^ ulAlphaF));
*(pResult + nIndex) = (ulPixelF + ulPixelB);
}
}
AlphaBlending_Device() 是处理两个图像的 Alpha 混合的入口进程。
参数
unsigned long* pMemA [in]：主机前景图像缓冲区。unsigned long* pMemB [in]：主机背景图像缓冲区。unsigned long* pResult [out]：主机混合图像缓冲区。unsigned nWidth [in]：图像宽度。unsigned nHeight [in]：图像高度。
extern "C"
void AlphaBlending_Device(unsigned long* pMemA, unsigned long* pMemB,
unsigned long* pMemResult,
unsigned nWidth, unsigned nHeight)
{
unsigned nDimen = (nWidth * nHeight);
unsigned nSize = (nDimen << 2);
unsigned char *pDevA = NULL, *pDevB = NULL, *pDevResult = NULL;
cudaMalloc((void**)&pDevA, nSize);
cudaMalloc((void**)&pDevB, nSize);
if (false == gm_bMapHostMemory)
{
cudaMalloc((void**)&pDevResult, nSize);
}
else
{
cudaHostGetDevicePointer((void**)&pDevResult, (void*)pMemResult, 0);
}
cudaMemcpy(pDevA, pMemA, nSize, cudaMemcpyHostToDevice);
cudaMemcpy(pDevB, pMemB, nSize, cudaMemcpyHostToDevice);
cudaBindTexture(0, texForegnd, pDevA);
cudaBindTexture(0, texBackgnd, pDevB);
AlphaBlending_Texture<<<::ceil((float)nDimen /
(float)BLOCK_DIM), BLOCK_DIM>>>((unsigned long*)pDevResult, nDimen);
if (false == gm_bMapHostMemory)
{
cudaMemcpy(pMemResult, pDevResult, nSize, cudaMemcpyDeviceToHost);
}
cudaUnbindTexture( texForegnd );
cudaUnbindTexture( texBackgnd );
cudaFree( pDevA );
cudaFree( pDevB );
if (false == gm_bMapHostMemory)
{
cudaFree( pDevResult );
}
}
在这种情况下，在 GPU 上的处理时间仅为 0.000067 (ms)。
关注点
结果非常棒。使用 CUDA 的处理时间比传统方式快近 100 倍。我想尝试另一种方法；OpenCL，它也是一种并行计算语言，看看它们之间的差距有多大。
修订
- 版本：1.0 初始发布