OpenCV-CUDA 图像处理

CUDA编程

GPU并行计算

架构本质与定位

1. GPU的从属角色：GPU并非独立计算平台，而是CPU的协处理器，需与CPU协同工作，因此GPU并行计算的本质是基于CPU+GPU的异构计算架构。
2. 核心组成与连接：CPU（主机端/host）与GPU（设备端/device）通过PCIe总线连接，共同构成异构计算体系。

CPU与GPU的核心差异

二者在核心数量、适用任务、线程特性上存在显著区别，具体对比如下：

异构架构的优势互补逻辑

CPU与GPU通过分工协作发挥最大功效，具体分工为：

• CPU：负责处理逻辑复杂的串行程序，承担任务调度、逻辑判断等核心控制工作。
• GPU：重点处理数据密集型的并行计算程序，利用多核心优势高效完成大规模重复简单运算。

CUDA编程的基础知识和核心概念

1. 内核 (Kernel) 单指令多线程

内核是在GPU上执行的函数，是CUDA并行计算的核心。它由__global__关键字声明，只能从主机(CPU)调用，但在设备(GPU)上执行。

特点：

• 内核启动时会产生大量并行线程，所有线程执行相同的内核函数（单指令多线程SIMT模型）
• 线程通过索引区分自己需要处理的数据
• 内核启动语法：kernel_name<<<gridDim, blockDim, sharedMem, stream>>>(args)
  • gridDim：网格维度（线程块数量）
  • blockDim：线程块维度（每个线程块中的线程数量）
  • sharedMem：每个线程块分配的共享内存大小（可选）
  • stream：指定流（可选，用于异步执行）

示例：

// 声明一个内核函数
__global__ void myKernel(int *data) {
    // 内核函数体
}

// 启动内核，1024个线程块，每个线程块256个线程
myKernel<<<1024, 256>>>(d_data);

2. 线程 (Thread)

线程是GPU上的基本执行单位，是并行计算的最小单元。每个线程独立执行内核函数，拥有自己的私有内存（寄存器、局部内存）。

特点：

• 每个线程有唯一的标识符，用于确定自己需要处理的数据
• 线程执行时通过内置变量获取自身索引：
  • threadIdx.x：线程在块内的x维度索引
  • threadIdx.y：线程在块内的y维度索引（用于2D线程块）
  • threadIdx.z：线程在块内的z维度索引（用于3D线程块）
• 线程的执行顺序不确定，不能假设执行顺序

线程索引计算示例：

__global__ void kernel(int *array, int n) {
    // 计算1D全局索引
    int globalIndex = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 确保不越界访问
    if (globalIndex < n) {
        array[globalIndex] = ...; // 处理数据
    }
}

3. 线程块 (Block)

线程块是一组线程的集合（最多1024个线程），块内线程可以：

• 通过共享内存(__shared__)交换数据
• 通过同步函数(__syncthreads())协调执行

特点：

• 线程块具有三维索引：blockIdx.x, blockIdx.y, blockIdx.z
• 块内线程共享__shared__内存（比全局内存快得多）
• 线程块大小通常选择2的幂（如128、256、512、1024）以优化性能
• 同一线程块的线程必须在同一流式多处理器(SM)上执行
• __syncthreads()用于块内线程同步，确保所有线程都完成某阶段后再继续

线程块示例：

__global__ void sharedMemoryKernel(float *input, float *output, int n) {
    // 声明共享内存（每个线程块有一份）
    __shared__ float s_data[256];
    
    // 计算索引
    int tid = threadIdx.x;
    int globalIdx = blockIdx.x * blockDim.x + tid;
    
    // 加载数据到共享内存
    s_data[tid] = input[globalIdx];
    
    // 同步：确保所有线程都加载完数据
    __syncthreads();
    
    // 处理共享内存中的数据（比全局内存快）
    output[globalIdx] = s_data[tid] * 2.0f;
}

4. 网格 (Grid)

网格是一组线程块的集合，一个内核启动对应一个网格。网格中的线程块可以在GPU的多个流式多处理器(SM)上并行执行。

特点：

• 网格可以是1D、2D或3D结构，便于处理多维数据（如图像、矩阵）
• 线程块之间不能直接通信或同步，必须通过全局内存间接通信
• 网格中的线程块数量理论上没有上限（受GPU内存限制）
• 全局内存是网格中所有线程共享的内存空间

2D网格和2D线程块示例（适合图像处理）：

__global__ void imageProcessingKernel(unsigned char *input, unsigned char *output, int width, int height) {
    // 计算2D索引
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    // 计算线性索引
    int idx = y * width + x;
    
    // 处理像素（确保在图像范围内）
    if (x < width && y < height) {
        output[idx] = 255 - input[idx]; // 简单的图像反转
    }
}

// 启动内核：2D网格和2D线程块
dim3 blockSize(16, 16); // 16x16 = 256线程 per block
dim3 gridSize((width + blockSize.x - 1) / blockSize.x, 
              (height + blockSize.y - 1) / blockSize.y);
imageProcessingKernel<<<gridSize, blockSize>>>(d_input, d_output, width, height);

层级关系与执行模型

这些概念形成了一个层级结构：

网格(Grid) → 包含多个线程块(Block)
线程块(Block) → 包含多个线程(Thread)
线程(Thread) → 执行内核函数的基本单元

GPU执行时：

1. 主机启动内核，指定网格和线程块配置
2. GPU将网格中的线程块分配到不同的流式多处理器(SM)
3. 每个SM将线程块划分为32线程的线程束(Warp)进行调度
4. 线程束内的线程执行相同的指令，但可以有不同的数据路径

关键总结

• 线程：最小执行单位，有私有内存，通过索引区分工作
• 线程块：线程的集合，支持共享内存和同步，有硬件数量限制
• 网格：线程块的集合，覆盖整个问题空间，线程块间无直接同步
• 内核：GPU上执行的函数，由整个网格的所有线程共同执行

CUDA内存模型与GPU硬件架构核心要点

1、CUDA内存模型

1. 内存层次结构：

   • 本地内存（Local Memory）：每个线程私有，仅线程自身可访问
   • 共享内存（Shared Memory）：线程块内所有线程共享，生命周期与线程块一致
   • 全局内存（Global Memory）：所有线程可访问，是最常用的全局存储空间
   • 只读内存：包括常量内存（Constant Memory）和纹理内存（Texture Memory），适用于读取频繁、写入极少的数据

2. 内存与优化：内存结构设计对程序性能优化至关重要，不同类型内存的访问效率差异显著

2、GPU硬件核心架构

1. 核心组件：

   • SM（流式多处理器）：GPU的核心执行单元，包含CUDA核心、共享内存和寄存器等
   • CUDA核心：并行计算的基本单元，SM通过调度线程实现并行计算

2. 线程执行机制：

   • 逻辑与物理并行：Kernel启动的大量线程是逻辑并行，物理上依赖SM的CUDA核心数量
   • 线程块分配：线程块被分配到SM上执行，一个线程块只能在一个SM上调度，一个SM可调度多个线程块
   • SIMT架构：SM采用单指令多线程架构，基本执行单元是线程束（warps）
   • 线程束特性：每个线程束包含32个线程，同时执行相同指令，但拥有独立的寄存器状态和执行路径

3、性能关键要点

1. 线程束分化：当线程束内线程遇到分支结构时，部分线程可能闲置等待，导致性能下降
2. 资源限制：SM同时并发的线程束数量受限于共享内存和寄存器等资源
3. 配置影响：Kernel的grid和block配置直接影响性能，需合理设计
4. 块大小建议：线程块大小通常应设置为32的倍数，以匹配线程束的32线程结构

OpenCV-CUDA图像处理

利用OpenCV和CUDA实现基础的图像处理算法。

Cmake

cmake_minimum_required(VERSION 3.10)
project(CUDAImageProcessing)

# 设置C++标准
set(CMAKE_CXX_STANDARD 11)

# 启用CUDA支持
enable_language(CUDA)

# 针对RTX 3060设置CUDA架构 (计算能力8.6)
set(CMAKE_CUDA_ARCHITECTURES 86)

# 查找OpenCV库
find_package(OpenCV REQUIRED)
message(STATUS "OpenCV VERSION: ${OpenCV_VERSION}")
message(STATUS "OpenCV库: ${OpenCV_LIBS}")
message(STATUS "OpenCV头文件: ${OpenCV_INCLUDE_DIRS}")

# 添加可执行文件
add_executable(test
    src/main.cu  # 你的CUDA源文件
)

# 设置CUDA编译选项
set_target_properties(test PROPERTIES
    CUDA_SEPARABLE_COMPILATION ON
)

# 包含目录
target_include_directories(test PRIVATE
    ${OpenCV_INCLUDE_DIRS}
)

# 链接库
target_link_libraries(test PRIVATE
    ${OpenCV_LIBS}
)

Code

main.cu

#include <iostream>
#include <opencv2/opencv.hpp>
#include <chrono>
#include <string>

// CUDA核函数：彩色转灰度
__global__ void rgbToGray(const uchar3* const input, uchar* const output, int width, int height) {
    // 计算全局线程索引
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    // 检查坐标是否在图像范围内
    if (x < width && y < height) {
        // 获取当前像素位置
        int idx = y * width + x;
        
        // 获取RGB值
        uchar3 rgb = input[idx];
        
        // 计算灰度值 (标准转换公式)
        output[idx] = static_cast<uchar>(0.299f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z);
    }
}

// 验证CUDA和OpenCV结果是否一致
bool verifyResults(const cv::Mat& cvResult, const cv::Mat& cudaResult) {
    if (cvResult.size() != cudaResult.size() || cvResult.type() != cudaResult.type()) {
        return false;
    }
    
    for (int i = 0; i < cvResult.rows; ++i) {
        for (int j = 0; j < cvResult.cols; ++j) {
            if (cvResult.at<uchar>(i, j) != cudaResult.at<uchar>(i, j)) {
                std::cerr << "验证失败: 像素差异在 (" << i << ", " << j << ")" << std::endl;
                return false;
            }
        }
    }
    
    return true;
}

int main(int argc, char** argv) {
    // 读取图像
    std::string imagePath = "../images/input.jpg";
    cv::Mat image = cv::imread(imagePath, cv::IMREAD_COLOR);
    if (image.empty()) {
        std::cout << "fail to load: " << imagePath << std::endl;
        return -1;
    }
    
    std::cout << "image size: " << image.cols << "x" << image.rows << std::endl;
    
    // 获取图像尺寸
    int width = image.cols;
    int height = image.rows;
    
    // --------------- OpenCV直接灰度转换 ---------------
    cv::Mat grayOpenCV;
    auto startCV = std::chrono::high_resolution_clock::now();
    
    cv::cvtColor(image, grayOpenCV, cv::COLOR_BGR2GRAY);
    
    auto endCV = std::chrono::high_resolution_clock::now();
    auto durationCV = std::chrono::duration_cast<std::chrono::milliseconds>(endCV - startCV).count();
    std::cout << "OpenCV cost time: " << durationCV << " ms" << std::endl;
    
    // --------------- CUDA灰度转换 ---------------
    // 分配GPU内存
    uchar3* d_input;
    uchar* d_output;
    cudaMalloc(&d_input, width * height * sizeof(uchar3));
    cudaMalloc(&d_output, width * height * sizeof(uchar));
    
    // 将图像数据传输到GPU
    cudaMemcpy(d_input, image.ptr<uchar3>(), width * height * sizeof(uchar3), cudaMemcpyHostToDevice);
    
    // 定义线程块和网格尺寸
    dim3 blockSize(16, 16);
    dim3 gridSize((width + blockSize.x - 1) / blockSize.x, 
                  (height + blockSize.y - 1) / blockSize.y);
    
    // 预热运行 (排除编译和初始化开销)
    rgbToGray<<<gridSize, blockSize>>>(d_input, d_output, width, height);
    cudaDeviceSynchronize();  // 确保核函数执行完成
    
    // 性能计时
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    
    cudaEventRecord(start);
    
    // 执行核函数
    rgbToGray<<<gridSize, blockSize>>>(d_input, d_output, width, height);
    
    // 同步并测量时间
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    std::cout << "CUDA cost time: " << milliseconds << " ms" << std::endl;
    
    // 创建输出图像
    cv::Mat grayCUDA(height, width, CV_8UC1);
    
    // 将结果从GPU传回CPU
    cudaMemcpy(grayCUDA.ptr(), d_output, width * height * sizeof(uchar), cudaMemcpyDeviceToHost);
    
    // 释放GPU内存
    cudaFree(d_input);
    cudaFree(d_output);
    
    // 保存结果
    cv::imwrite("output_opencv.jpg", grayOpenCV);
    cv::imwrite("output_cuda.jpg", grayCUDA);
    
    // 验证结果
    bool resultsMatch = verifyResults(grayOpenCV, grayCUDA);
    std::cout << "Val: " << (resultsMatch ? "Success" : "Fail") << std::endl;
    
    // 计算加速比
    if (milliseconds > 0) {
        float speedup = durationCV / milliseconds;
        std::cout << speedup << "x "  << " CUDA faster than Opencv "<< std::endl;
    }
    
    std::cout << "End!" << std::endl;
    
    return 0;
}