CUDA学习笔记（对图像进行并行加速运算）

在对CUDA的kernel函数有了一定了解之后，就可以对图像进行多线程加速计算，下面是我在CPU和GPU上各自遍历一张图像，对比两者的性能差异。
这里需要用到opencv的函数读取图像以及显示图像，因此opencv的头文件需要包含进去。其他的注意事项主要是图像的内存分配和存储类型问题，比如彩色图像是三通道的，申请的内存为 row * col * sizeof(uchar3)。

首先在CPU上遍历一张彩色图像的所有像素：

    for (int i = 0; i < row; i++)
    {
        for (int j = 0; j < col; j++)
        {
            int index = i * row + j;
            dst.data[3 * index + 0] = (uchar)src.data[3 * index + 0];
            dst.data[3 * index + 1] = (uchar)src.data[3 * index + 1];
            dst.data[3 * index + 2] = (uchar)src.data[3 * index + 2];

        }
    }

在GPU上遍历一张图像的所有像素，其kernel函数为：

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col,int row )
{

    int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
        + threadIdx.z * (blockDim.x * blockDim.y) \
        + threadIdx.y * blockDim.x \
        + threadIdx.x;
    if (tid < col * row* 3) {
        dst_dev[tid].x = _src_dev[tid].x ;
        dst_dev[tid].y = _src_dev[tid].y ;
        dst_dev[tid].z = _src_dev[tid].z ;
    }
}

线程索引计算方法各有不同，很多新手会在这个地方犯迷糊，从其他地方看到了这个万能计算方法（仅限于单个维度的线程索引），用起来相当省事。
下面是完整的代码：

#include "cuda_runtime.h"
#include "iostream"
#include "opencv2/opencv.hpp"
#include

using namespace std;
using namespace cv;

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col,int row )
{

    int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
        + threadIdx.z * (blockDim.x * blockDim.y) \
        + threadIdx.y * blockDim.x \
        + threadIdx.x;
    if (tid < col * row* 3) {
        dst_dev[tid].x = _src_dev[tid].x ;
        dst_dev[tid].y = _src_dev[tid].y ;
        dst_dev[tid].z = _src_dev[tid].z ;
    }
}

int main()
{
    Mat src = cv::imread("E:\\picture\\source\\lena.jpg");
    Mat dst=src;
    int col = src.cols;
    int row=src.rows;

    clock_t start = clock();
    for (int i = 0; i < row; i++)
    {
        for (int j = 0; j < col; j++)
        {
            int index = i * row + j;
            dst.data[3 * index + 0] = (uchar)src.data[3 * index + 0];
            dst.data[3 * index + 1] = (uchar)src.data[3 * index + 1];
            dst.data[3 * index + 2] = (uchar)src.data[3 * index + 2];

        }
    }
    cout << "cpu所耗费的时间：" << (double)(clock() - start)<< "ms" << "\n";
    cv::imshow("dst", dst);

    uchar3* src_dev, * dst_dev;
    Mat _dst = Mat(row, col , CV_8UC3);
    cudaMalloc((void**)&src_dev, row * col * sizeof(uchar3));
    cudaMalloc((void**)&dst_dev, row * col * sizeof(uchar3));
    cudaMemcpy(src_dev, src.data, row * col * sizeof(uchar3), cudaMemcpyHostToDevice);

    dim3 grid(1 + (col * row / (32 * 32 + 1)), 1, 1);
    dim3 block(32, 32, 1);

    cudaEvent_t start1, stop1;
    cudaEventCreate(&start1);
    cudaEventCreate(&stop1);
    cudaEventRecord(start1, 0);

    Traverse << < grid, block >> > (src_dev, dst_dev,col,row);
    cudaEventRecord(stop1, 0);
    cudaEventSynchronize(stop1);
    float time1;
    cudaEventElapsedTime(&time1, start1, stop1);
    cout << "Gpu所耗费的时间：" << time1 << "ms" << "\n";
    cudaMemcpy(_dst.data, dst_dev, row * col * sizeof(uchar3), cudaMemcpyDeviceToHost);
    cv::imshow("_dst", _dst);

    cudaFree(src_dev);
    cudaFree(dst_dev);

    waitKey(0);
    system("pause");
    return 0;
}

最后的计算结果为：

需要注意的是，由于一张图片的计算量有限，除非是特别大的图像或者数据集，否则cuda并行计算的优势并不明显，甚至会比CPU的计算时间还要长，此外这里是对像素进行遍历拷贝复制，需要频繁的进行内存读取操作，这并不能体现出GPU的并行计算优势。
刚才演示的是单个维度的线程索引，现在用两个维度的线程进行测试，核函数改为：

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col, int row)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    dst_dev[offset].x = _src_dev[offset].x;
    dst_dev[offset].y = _src_dev[offset].x;
    dst_dev[offset].z = _src_dev[offset].x;
}

上面的测试用例是512*512的lena图像，使用CUDA的时间优势还是没有体现出来，为了证明GPU的加速性能，下面换用 3000*2000的图片进行测试。

这的例子只是为了介绍一下最基本的CUDA对图片进行并行加速的思路，方便大家举一反三，触类旁通，其实在这个基础上稍加改动就可以尝试用CUDA对图像进行各种图像处理的操作，比如阈值分割，两个图象叠加等。
比如核函数改为：

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col,int row )
{

    int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
        + threadIdx.z * (blockDim.x * blockDim.y) \
        + threadIdx.y * blockDim.x \
        + threadIdx.x;
    if (tid < col * row * 3) {
        dst_dev[tid].x = 255- _src_dev[tid].x;
        dst_dev[tid].y = 255- _src_dev[tid].y;
        dst_dev[tid].z = 255- _src_dev[tid].z;
    }
}

核函数改为：

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col,int row )
{

    int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
        + threadIdx.z * (blockDim.x * blockDim.y) \
        + threadIdx.y * blockDim.x \
        + threadIdx.x;
    if (tid < col * row * 3) {
        dst_dev[tid].x = 0.3 * _src_dev[tid].x;
        dst_dev[tid].y = 0.6 * _src_dev[tid].y;
        dst_dev[tid].z = 0.1 * _src_dev[tid].z;
    }
}

核函数改为：

__global__ void Traverse(uchar3* _src_dev, uchar* dst_dev, int col,int row )
{

    int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
        + blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
        + threadIdx.z * (blockDim.x * blockDim.y) \
        + threadIdx.y * blockDim.x \
        + threadIdx.x;
    if (tid < col * row * 3) {
        dst_dev[tid] = 0.3 * _src_dev[tid].x+ 0.6 * _src_dev[tid].y+ 0.1 * _src_dev[tid].z;
    }

}

Original: https://blog.csdn.net/Asabc12345/article/details/122413671
Author: 滑了丝的螺丝钉
Title: CUDA学习笔记（对图像进行并行加速运算）