CUDA图像处理加速demo

环境搭建和安装就不提了。
主要步骤是申请显存，将内存复制到显存，执行核函数，将显存复制回内存。
核函数是可以认为是线程的worker函数。

直接上代码：

这段是初始化

   cv::Mat inputImage = cv::imread("");
    cv::cvtColor(inputImage, inputImage, cv::COLOR_BGR2GRAY);
    int w = 43;
    int h = 43;
    cv::Mat inputImage2;

    int numGPUs;
    cudaGetDeviceCount(&numGPUs);
    if (numGPUs <= 0) {
        std::cerr << "No CUDA-capable devices found." << std::endl;
        return 1;
    }
    int deviceId = 0; 
    cudaDeviceProp deviceProp;

    // 查询设备属性
    cudaGetDeviceProperties(&deviceProp, deviceId);

下面这段显存初始化和拷贝显存

cudaMalloc 和 cudaMemcpy，注意cudaMemcpyHostToDevice这是从Host到Device，device就是显卡

    int inRows = inputImage.rows;
    int inCols = inputImage.cols;
    uchar* d_inputImage, * d_outputImage, * d_tempImage;
    size_t inputSize = inRows * inCols * sizeof(uchar);
    size_t outputSize = (inRows) * (inCols) * sizeof(uchar);

    cudaMalloc(&d_inputImage, inputSize);
    cudaMalloc(&d_outputImage, outputSize);
    cudaMalloc(&d_tempImage, outputSize);
    cudaMemcpy(d_inputImage, inputImage.data, inputSize, cudaMemcpyHostToDevice);

下面这段是分块，执行核函数，还有计时功能

一般的卡是几百个核心。block是每个核心的线程数，grid是核心数。
我们让每个线程只处理一个像素，因此grid和block有图像大小确定关系。
erodeKernel2 是定义的核函数。

    dim3 block(16, 16);
    dim3 grid((inCols + block.x - 1) / block.x, (inRows + block.y - 1) / block.y);

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


    erodeKernel2 << <grid, block >> > (d_inputImage, d_tempImage, inRows, inCols, w, h);
    erodeKernel2_1 << <grid, block >> > (d_tempImage, d_outputImage, inRows, inCols, w, h);
    // 记录结束事件
    cudaEventRecord(stop);
    // 等待事件完成
    cudaEventSynchronize(stop);
    {
        // 计算运行时间
        float milliseconds = 0;
        cudaEventElapsedTime(&milliseconds, start, stop);
        std::cout << "GPU runtime: " << milliseconds << " ms" << std::endl;
    }

下面来看核函数，先拿一个最普通的核函数。

核函数要用__global__ 修饰。这是一个腐蚀的核函数，也就是最小值滤波。
熟悉多线程的知道可以通过线程的ID来确定要处理哪个位置的像素，CUDA也不例外。x和y是图像坐标，我们以上面的划分为例，则blockDim.x是0-gridx的一个数，blockIdx.x是16，threadIdx.x是0-16的一个数。在这个核函数中，只处理一个像素。
为了能处理全部的图像我们划分的block*grid往往大于图像的尺寸，因此要判断x，y是否在图像范围内，如果在范围内才处理。

__global__ void erodeKernel_junk(const uchar* input, uchar* output, int rows, int cols, int kernelSize, int kernelSize2) {
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    int y = blockDim.y * blockIdx.y + threadIdx.y;
    if (x < cols && y < rows) {
        const int halfSize = kernelSize / 2;
        const int halfSize2 = kernelSize2 / 2;
        uchar minVal = 255;
        for (int i = -halfSize2; i <= halfSize2; ++i) {
            for (int j = -halfSize; j <= halfSize; ++j) {
                int px = x + j;
                int py = y + i;
                if (px >= 0 && px < cols && py >= 0 && py < rows) {
                    uchar val = input[py * cols + px];
                    if (val < minVal) {
                        minVal = val;
                    }
                }
            }
        }
        output[y * cols + x] = minVal;
    }
}

下面是一个有共享内存的核函数

共享内存可以加快计算效率，减少内存读写消耗。
线程块内可以共享内存，我们定义的线程块为 dim3 block(16, 16)，也就是在每个1616线程可以共享内存。因为处理的像素是和线程相关的，因此也可以说1616大小的图像共享内存。
我们定义共享内存的大小：__ shared __ uchar sharedBlock[33 * 33] 一般来说这个数要比线程块要大，因为在每个线程中我们都要往共享内存中存放数据。

sharedBlock[py * blockDim.x + px] = input[y * cols + x];给共享内存写入值。从这里可以明显看出共享内存中的内容为input[y * cols + x];而在一个线程块中x和y的范围是0-16，因此就16*16的邻域图像放在了共享内存中，这对滤波是极为有利的。

__syncthreads();是等待所有线程块线程都执行到此处，也就说共享内存已经写入完毕，不会出现空值。

然后在处理图像时要从共享内存取值，而不是从input显存中取值。
但是这里有一种特殊情况，加入滤波器宽高很大，超过了共享内存有效大小怎么办（这里的有效大小为16*16），那需要判断一下， 如果超过了范围就从input中取值。
也就是 if (px + j < 0 || px + j >= blockDim.x) { val = input[y * cols + x + j]; } else { val = sharedBlock[py * blockDim.x + px + j]; }

/*行处理 共享内存*/
__global__ void erodeKernel2(const uchar* input, uchar* output, int rows, int cols, int kernelSize, int kernelSize2) {
    __shared__ uchar sharedBlock[33 * 33];  // 假设每个线程块的共享内存大小为32x32
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    int y = blockDim.y * blockIdx.y + threadIdx.y;
    const int halfSize = kernelSize / 2;
    if (x > halfSize && x < cols - halfSize && y < rows - halfSize && y>halfSize) {

        int px = threadIdx.x;
        int py = threadIdx.y;
        sharedBlock[py * blockDim.x + px] = input[y * cols + x];
        __syncthreads();
        uchar minVal = 255;
        for (int j = -halfSize; j <= halfSize; ++j) {
            uchar val = 254;
            if (px + j < 0 || px + j >= blockDim.x) {
                //continue;
                val = input[y * cols + x + j];
                //val = sharedBlock[py * blockDim.x + 0];
            }
            else {
                val = sharedBlock[py * blockDim.x + px + j];
            }
            //val = sharedBlock[py * blockDim.x + px + j];
            //val = input[y * cols + x + j];
            if (val < minVal) {
                minVal = val;
            }
        }

        output[y * cols + x] = minVal;
    }

}