使用 CPU、GPU 和 C++ AMP 进行图像滤镜处理

• 下载源代码 - 包含 C++ AMP，CUDA 6.5，VS 2013 解决方案 - 693.5 KB
• 下载源代码 - 不包含 C++ AMP，VS 2010 解决方案 - 151.2 KB

• 项目 Doxygen 文档 - 链接
• GitHub 上的最新代码

引言

在本文中，我们将研究一个简单的通用框架，用于处理单输入、单输出的图像滤镜。滤镜可以以纯 C++ 在 CPU 上运行，使用 nVidia CUDA 在 GPU 上运行，使用 C++ AMP (至少需要 VS 2012 RC) 或这些技术的组合来实现。

我们将使用 OpenCV 捕获网络摄像头输入，并使用实现的图像滤镜处理捕获的帧。

要求

• 补丁版适用于 Windows 的 OpenCV 2.2。
• CUDA 6.5 SDK，可与 Visual Studio 2013 编译器配合使用。
• Visual Studio 2013
• 一款支持 CUDA 的 nVidia 显卡，已安装最新驱动程序。

背景

本文并非介绍 OpenCV、CUDA 或 C++ AMP。只提供简要说明。

什么是图像处理？

图像处理滤镜表示应用于图像的数学运算，例如颜色反转、任意旋转、模糊等。

什么是 OpenCV？

OpenCV 是一个开源的计算机视觉库，它除了其他功能之外，还可以帮助我们为本文捕获网络摄像头。

什么是 CUDA？

CUDA 是并行编程领域的颠覆者，它是一种范式转变，它不是多核计算，而是众核计算。为什么是众核？核心数量如此之多，以至于您无需确切知道有多少个核心，而缺乏此知识可以实现透明的可扩展性。

使用 CUDA，您可以编写在 GPU 上运行的 C/C++ 代码，利用硬件优化的并行处理。

CUDA 是 nVidia 的专有技术，要使用它，您需要一块 nVidia 显卡和支持 CUDA 的最新显卡驱动程序。

什么是 C++ AMP？

C++ Accelerated Massive Parallelism 是一个库，它在幕后使用 DirectX 11 在 GPU 上进行计算，并回退到 CPU。该库由 Microsoft 提供，可用于 VS 2012 及更高版本。更多信息 (MSDN)

Using the Code

OpenCV 帧捕获精炼

1. 打开网络摄像头设备。
2. 捕获第一帧。
3. 如果第一帧不 OK，则退出。
4. 进入主循环
   1. 捕获下一帧。
   2. 如有必要，调整帧大小。
   3. 调用图像滤镜处理捕获的帧。
   4. 在屏幕上显示处理后的图像。

接口、类、继承等

我们正在开发的点在于，滤镜的操作独立于图像的创建方式。我们可以将过滤方法写在 main.cpp 中作为全局函数，或者我们可以使用一个接口来在自定义滤镜类的实例上调用我们的过滤方法。

最简单的形式是，可以看到过滤图像的接口至少需要一个过滤图像的方法，一个传递图像尺寸的方法，以及一个管理实现者所持资源的释放方法。假设我们可以在 filter 方法中传递图像尺寸并在返回前清理内存；但我们需要一种方法来分配和释放 GPU 上的内存。GPU 上的内存操作成本很高，这就是为什么我们使用 InitFilter 和 ReleaseFilter 方法与接口 ISingleImageFilter 中的 FilterImage 方法。

最简单的滤镜：IdentityFilter

IdentityFilter 按原样返回图像，实现了 ISingleImageFilter 接口。

class IdentityFilter : public SingleImageFilter
{
public:
    IdentityFilter()
    {
    }

    /** Görüntüde degisiklik yapmadan çikar. */
    virtual void FilterImage(char* imageData)
    {
        return; // imaji degistirmeden dön.
    }
};

要使用任何 ISingleImageFilter，您需要执行以下操作。首先，必须调用 InitFilter 并传入图像尺寸。在帧捕获循环中，调用 FilterImage 并传入图像数据。FilterImage 方法执行就地处理，这意味着您的引用将被结果覆盖。当您完成滤镜使用并且不再需要它时，调用 ReleaseFilter 方法。

请注意，我们仅在 IdentityFilter 中重写了 FilterImage 方法。图像尺寸保存在基类 SingleImageFilter 中，它是一个 abstract 类，只将 FilterImage 方法留为一个纯虚方法。

CPU 上的颜色反转滤镜

在本文中，图像被捕获为 BGR 三通道图像，每个颜色通道有 8 位。颜色反转意味着每个像素的每个通道都被反转，因此

value + inverted value = max value

在 8 位情况下，最大值为 255，因此这是颜色反转滤镜的代码

/** Görüntünün RGB kanallarinin tersini alir. */
virtual void FilterImage(char* imageData)
{
    for(int i=0; i<3*width*height; i++)
    {
        *( imageData + i ) = ( unsigned char ) ( 255 - *( imageData + i ) ); // her pikselin her kanalinin negatifini al.
    }
}

在滤镜中使用 CUDA

要编译 CUDA 示例，您需要安装 CUDA SDK 6.5。

将在 GPU 上运行的代码称为内核。要将方法标记为内核，您需要应用方法装饰器 __global__。

__global__
void gpuInvert(
    float* image,
    int width,
    int height
    )
{
// ...
}

要启动内核，我们使用配置参数，这些参数告诉 GPU 我们希望内核在多少块和线程上运行。

// launch kernel
dim3 dimBlock( BLOCK_SIZE, BLOCK_SIZE );
dim3 dimGrid( width / dimBlock.x, height / dimBlock.y );

gpuInvert<<< dimGrid, dimBlock >>>( d_Image, width, height );

图像处理是 CUDA 的一个广阔领域，因为许多图像处理算法非常适合内存访问模式和数学复杂性的并行处理。

我们已经提到 CUDA 程序本身就在 GPU 上运行，那么数据应该放在哪里？我们在每一帧中都将图像数据传输到显卡内存中。我们在 InitFilter 中分配设备内存，在 ReleaseFilter 中释放它。

支持 CUDA 的滤镜的基类

在 FilterImage 方法中，我们需要将图像数据复制到 GPU 内存，处理数据，并将结果取回我们原来的 RAM。由于我们使用的是单输入、单输出的图像处理滤镜，我们可以在 GPU 内核中使用相同的函数签名。

我们将大量使用内核启动函数指针，我们将为它们定义一个类型。

typedef void (*ptKernelLauncher)(float*, int, int); // define

ptKernelLauncher kernelLauncher; // declare

kernelLauncher = deviceInvertLaunch; // assign

kernelLauncher( d_Image, width, height ); // invoke

从 SingleCudaFilter 类开始，我们不会直接启动内核。包含 CUDA 内核的 C/C++ 代码需要通过 **nvcc** nVidia 编译器驱动程序。我们通过使用所谓的内核启动器方法来抽象内核代码，并在我们的滤镜类中调用这些启动器方法。

在 SingleCudaFilter 类中，我们不会有任何内核或启动器依赖。我们将在该类的构造函数中传递启动器函数指针。当在图像处理循环中调用 FilterImage 方法时，SingleCudaFilter 将调用内核启动器，而内核启动器又会在 GPU 上启动内核。

在此代码片段中，为 SingleCudaFilter 类包含了 InitFilter、FilterImage 和 ReleaseFilter 方法。

class SingleCudaFilter : public SingleImageFilter
{
protected:
    float* h_Image; /**< Normalize edilmis görüntünün CPU bellek adresi. */
    float* d_Image; /**< Normalize edilmis görüntünün GPU bellek adresi. */

    ptKernelLauncher kernelLauncher;

public:

    /**
        kernelLauncher metod isaretçisini alan yaratici.

        \param kernelLauncher \ref kernelLauncher tipinde metod isaretçisi alan yaratici.
    */
    explicit SingleCudaFilter( ptKernelLauncher kernelLauncher )
        : kernelLauncher(kernelLauncher)
    {
    }

    /**
        CPU ve GPU üzerinde normalize edilmis görüntüler için bellek ayirir.
    */
    virtual void InitFilter(int width, int height)
    {
        SingleImageFilter::InitFilter(width, height);

        /*
        allocate device memory
        */

        cudaMalloc( (void**) &d_Image, 3 * sizeof(float) * width * height );
        checkCUDAError("malloc device image");


        /*
        allocate host memory
        */

        cudaMallocHost( (void**) &h_Image, 3 * sizeof(float) * width * height );
        checkCUDAError("malloc host image");
    }

    /**
        Yaraticida alinan kerneli çagirir.

        \param imageData Görüntünün BGR kanal sirali bellekteki adresi.


        Görüntüyü normalize ederek kernelLauncher isaretçisinin gösterdigi 
        metodu çagirir ardindan görüntüyü denormalize eder( [0, 255] araligina ).
        Kernelde islenen görüntüden sonuç olarak [0, 1] araligi disinda bir 
        deger dönerse o kanalin degeri [0, 255] araligindan disarida olabilir. 
        Bu durumda deger yakin oldugu sinira indirgenir.
    */
    virtual void FilterImage(char* imageData)
    {
        // imageData degiskenindeki görüntü verisi normalize edilerek h_Image degiskenine aktarilir.
        for(int i=0; i<3*width*height; i++)
        {
            *(h_Image + i) = (unsigned char)*(imageData + i) / 255.0f; // normalize and copy image
        }

        /*
            Görüntü GPU bellegine kopyalanir.
        */

        cudaMemcpy( d_Image, h_Image, 3 * sizeof(float) * width * height, cudaMemcpyHostToDevice );
        checkCUDAError("FilterImage: memcpy");

        /*
            Constructorda verilen kernel çalistirilir.
        */
        kernelLauncher( d_Image, width, height );
    
        /*
            Sonuçlar CPU bellegine kopyalanir.
        */
        cudaMemcpy( h_Image, d_Image, 3 * sizeof(float) * width * height, cudaMemcpyDeviceToHost);
        checkCUDAError("FilterImage: memcpy2");

        /*
            h_Image degiskenindeki normalize edilmis görüntü verisi [0, 255] araligina çekilir.
        */
        for(int i=0; i<3*width*height; i++)
        {
            *(imageData + i) = satchar(*(h_Image + i) * 255);
        }
    }

    /**
        CPU ve GPU üzerinde normalize edilmis görüntüler için ayrilmis bellegi serbest birakir.
    */
    virtual void ReleaseFilter()
    {
        SingleImageFilter::ReleaseFilter();

        cudaFree( d_Image );
        checkCUDAError("free device image");
    
        cudaFreeHost( h_Image );
        checkCUDAError("free host image");
    }
};

GPU 上的颜色反转滤镜

CudaInvertFilter 类不过是 boilerplate 代码，用于将内核启动器传递给 SingleCudaFilter 类。

class CudaInvertFilter : public SingleCudaFilter
{
public:

    /**
        \ref deviceInvertLaunch metod isaretçisi parametresi ile SingleCudaFilter yaraticisini çagirir.
    */
    CudaInvertFilter()
        : SingleCudaFilter(deviceInvertLaunch)
    {
    }
    
};

而且内核中也没有什么特别之处。唯一需要考虑的是，我们通过从 1 减去通道值来反转通道，而不是从 255。我们在 FilterImage 方法中将归一化的图像传递给了内核启动器。

/**
    Görüntünün tersini alan kernel.

    \param image [0, 1] araligina normalize edilmis, BGR kanal sirali görüntünün GPU bellegindeki adresi.
    \param width Görüntünün piksel olarak genisligi
    \param height Görüntünün piksel olarak yüksekligi

    Metod GPU üzerinde çalisir, çiktisini image parametresinin üzerine yazar.

    */
__global__
void gpuInvert(
    float* image,
    int width,
    int height
    )
{
    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;

    int cIdx = ( row * width + col ) * 3; // 3 ile çarpim RGB için, linearIndex.

    // normalize edilmis pikselleri 1'den çikarttigimizda görüntünün negatifini almis oluruz.
    *( image + cIdx     ) = 1 - *( image + cIdx     ); // Blue kanali
    *( image + cIdx + 1 ) = 1 - *( image + cIdx + 1 ); // Green kanali
    *( image + cIdx + 2 ) = 1 - *( image + cIdx + 2 ); // Red kanali
}

我们使用了以下安排来启动内核

/**
    \ref ptKernelLauncher tipinde metod.

    \param d_Image [0, 1] araligina normalize edilmis, BGR kanal sirali görüntünün GPU bellegindeki adresi.
    \param width Görüntünün piksel olarak genisligi
    \param height Görüntünün piksel olarak yüksekligi

    \ref gpuInvert kernelini Grid ve Block boyutlarini ayarlayarak çagiran metod.
*/
void deviceInvertLaunch(
    float *d_Image,
    int width,
    int height
    )
{
    // launch kernel
    dim3 dimBlock( BLOCK_SIZE, BLOCK_SIZE );
    dim3 dimGrid( width / dimBlock.x, height / dimBlock.y );

#if ENABLE_TIMING_CODE

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);

#endif

    gpuInvert<<< dimGrid, dimBlock >>>( d_Image, width, height);

#if ENABLE_TIMING_CODE
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    float elapsedTime;
    cudaEventElapsedTime(&elapsedTime, start, stop);

    // block until the device has completed
    cudaThreadSynchronize();
    
    printf("gpuInvert kernel time: %.3f ms\n", elapsedTime);
#endif

    cudaThreadSynchronize();

    // check if kernel execution generated an error
    // Check for any CUDA errors
    checkCUDAError("kernel invocation");
}

在 CUDA 中使用纹理

使用前面讨论的方法实现的滤镜将使用显卡的全局内存。全局内存是可用内存中最慢的类型。我们可以使用 **纹理** 内存，它更快，具有空间缓存，并且是只读的。只读部分与我们无关，我们可以使用纹理内存。

在启动纹理支持并运行后，出现了一个问题。nvcc 编译器驱动程序通过代码文件为纹理引用分配作用域。这意味着您不能在头文件中声明一个纹理并在不同文件中使用它。如果您从另一个代码文件中引用纹理，您的代码将编译成功，但会发生错误！您将得到一个零（全黑）纹理等待着您。

为了从基类加载纹理并将当前帧提供给纹理中的内核，我们需要通过 CUDA 驱动程序 API 获取的引用来访问纹理。

我们的纹理问题尚未结束。可能是我做得不对，但在内核中使用纹理并从另一个文件加载它不应该如此困难。当您使用以下代码通过 CUDA 获取纹理引用时，您会面临挑战：

驱动程序 API 期望一个 const 指针。但如果您将指针声明为 const，您如何配置它指向的结构中的纹理参数？const_cast<> 拯救了我们。

const textureReference* constTexRefPtr;
textureReference* texRefPtr;

...

// cudaGetTextureReference, called with the const reference.
cudaGetTextureReference(&constTexRefPtr, textureSymbolName);
checkCUDAError("get texture reference");

// remove const from the reference.
texRefPtr = const_cast<textureReference*>( constTexRefPtr );
channelDesc = cudaCreateChannelDesc<float4>();

// use non-const reference to load and configure texture parameters.
cudaMallocArray( &cu_array, &texRefPtr->channelDesc, width, height );
checkCUDAError("malloc device image");

...

// Copy frame data to the array which will be bound to the texture.
cudaMemcpyToArray( cu_array, 0, 0, h_Image, sizeof(float4) * width * height, cudaMemcpyHostToDevice);
checkCUDAError("FilterImage: memcpy");

// Bind the texture to the cu_array.
cudaBindTextureToArray( texRefPtr, cu_array, &texRefPtr->channelDesc );

您可以深入 SingleCudaTexFilter.cu 和 SingleCudaTexFilter.h 文件，了解 CUDA 中纹理使用的详细信息。

在 GPU 上使用纹理进行颜色反转滤镜处理

使用纹理，我们可以通过 CUDA 函数 tex2D 获取任何像素的图像数据。内核启动器的调用在 SingleCudaTexFilter 中完成。

CudaTexInvertFilter 类

class CudaTexInvertFilter : public SingleCudaTexFilter
{

public:

    /**
        \ref deviceTexInvertLaunch metod isaretçisi ve "texInvert1" sembol adi parametresi ile SingleCudaTexFilter yaraticisini çagirir.
    */
    CudaTexInvertFilter()
        : SingleCudaTexFilter(deviceTexInvertLaunch, "texInvert1")
    {
    }

};

deviceTexInvertLaunch 内核

__global__
void gpuTexInvert(
    float* image,
    int width,
    int height
    )
{
    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;

    int cIdx = ( row * width + col ) * 3; // 3 ile çarpim RGB için, linearIndex.


    // threade ait kordinatin texture uzayindaki kordinati bulunur.
    float tu = (float)col / width;
    float tv = (float)row / height;

    // Texture üzerinden görüntü verisi okunur.
    float4 texVal = tex2D( texInvert1, tu, tv );

    // Texture degerleri 1'den çikartilarak global bellege yazilir.
    *( image + cIdx )     = 1 - texVal.x;
    *( image + cIdx + 1 ) = 1 - texVal.y;
    *( image + cIdx + 2 ) = 1 - texVal.z;
}

在滤镜中使用 C++ AMP

在同一个解决方案中处理 CUDA 和 C++ AMP 的技巧

GitHub 上的最新代码使用 CUDA 6.5，以下段落故意保留。

CUDA 4.2 要求使用 VS2008 (v90) 工具集，但，一个很大的但，C++ AMP 要求 VS2012 (v110) 工具集。因此，当您尝试编译包含 CUDA 和 C++ AMP 代码的单个项目时，CUDA 或 C++ AMP 代码将无法编译。我们可以为 C++ AMP 创建一个新的项目文件，并将包含 CUDA 代码的项目转换为 DLL 项目。这样，我们编译好的 CUDA 滤镜就可以被 AMP 项目使用了。

主项目现在是 AmpFilters 项目。CudaFilters 由 FilterFactory.cpp 中的工厂方法实例化。

C++ AMP 上的颜色反转滤镜

在此滤镜中，颜色反转操作实现为 parallel_for_each 内核。滤镜库的唯一缺点是使用 char 作为像素颜色数据类型。char 数据类型不支持在 restrict(amp) 块中使用。我们必须将数据转换为 int 或其他 AMP 支持的类型。

/** Görüntünün RGB kanallarinin tersini alir. */
void AmpInvertFilter::FilterImage(char* imageData)
{
    unsigned int* iImageData = (unsigned int*)malloc(3*width*height * sizeof(int));

    // AMP'de char kullanilamiyor, veriyi int'e donustur.
    for(int i=0; i<3*width*height; i++)
    {
        *( iImageData + i ) = ( unsigned int ) *( imageData + i );
    }

    const int size = 3*width*height;
    
    // Veri üzerinde dogrudan çalisabiliriz. (in-place).
    array_view<unsigned> img(size, iImageData);    

    parallel_for_each(         
        img.extent,         
        [=](index<1> idx) restrict(amp)
        {
            // Her kanalin negatifi alinir.
            img[idx] = 255 - img[idx];
        }
    );

    img.synchronize();

    // AMP'de char kullanilamiyor, veriyi char'e donustur.
    for(int i=0; i<3*width*height; i++)
    {
        *( imageData + i ) = ( char ) *( iImageData + i );
    }
}

滤镜链

如果不能将所有这些滤镜一个接一个地运行，那将毫无乐趣。通过实现 SingleImageFilter 类，SingleImageFilterChain 会按顺序调用其滤镜。

创建实例后，使用 AppendFilter 方法将滤镜添加到队列中。当调用 FilterImage 方法时，所有滤镜都有机会依次处理数据。

您可以在 SingleImageFilterChain 中混合使用 CPU、CUDA、CUDA 纹理滤镜和 C++ AMP 滤镜。

ISingleImageFilter* myFilter1 = new SingleCudaTexFilter(deviceTexAbsDiffLaunch, "texAbsDiff1");
ISingleImageFilter* myFilter2 = new CpuInvertFilter();

SingleImageFilterChain* myFilter = new SingleImageFilterChain();
myFilter->AppendFilter( myFilter1 );
myFilter->AppendFilter( myFilter2 );

已实现类的 UML 图

示例输出

有关全尺寸图像，请访问 http://dissipatedheat.com/2011/05/29/cuda-ile-opencv-kullanarak-webcam-goruntu-isleme/

关注点

通过 OpenCV 学习网络摄像头捕获以及使用带纹理的 CUDA，这是一个很好的体验。在纹理处理例程中遇到了一些棘手的问题，但现在都已解决。

历史

• 第三版发布于 2014 年 8 月 30 日 - 现在可以使用 VS 2013 和 CUDA 6.5 通过 C++ AMP 进行过滤
• 第二版发布于 2012 年 8 月 7 日 - 现在可以使用 VS 2012 RC 上的 C++ AMP 进行过滤
• 第一版发布于 2011 年 3 月 6 日 - 本文也发布在我的博客上：Dissipated Heat