OpenCV CUDA模块设备层-----GPU上执行线程安全的 “原子取最大值” 操作函数

• 操作系统：ubuntu22.04
• OpenCV版本：OpenCV4.9
• IDE:Visual Studio Code
• 编程语言：C++11

算法描述

这是一个 OpenCV 的 CUDA 模块（cv::cudev） 中封装的原子操作函数，用于在 GPU 上执行线程安全的 “原子取最大值” 操作。
将 *address 和 val 进行比较，如果 val > *address，就将 *address = val。

函数原型

__device__ __forceinline__ int cv::cudev::atomicMax(int* address, int val)

参数

返回值

返回修改前 *address 的原始值。

代码示例

#include <cuda_runtime.h>
#include <cstdio>

__global__ void kernel(int* max_value) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    // 使用 CUDA 内建的 atomicMax
    int old = atomicMax(max_value, tid);

    printf("Thread %d: old = %d, new = %s\n",
           tid, old, (tid > old ? "updated" : "not updated"));
}

int main() {
    int h_max = 0;
    int* d_max;

    // 分配设备内存
    cudaMalloc(&d_max, sizeof(int));
    if (!d_max) {
        printf("Failed to allocate device memory!\n");
        return -1;
    }

    // 初始化设备内存
    cudaMemcpy(d_max, &h_max, sizeof(int), cudaMemcpyHostToDevice);

    // 启动核函数
    kernel<<<1, 10>>>(d_max);

    // 显式同步设备
    cudaDeviceSynchronize();  // 等待核函数执行完成

    // 拷贝结果回主机
    cudaMemcpy(&h_max, d_max, sizeof(int), cudaMemcpyDeviceToHost);

    printf("Final max value: %d\n", h_max);  // 应该输出 9

    // 清理资源
    cudaFree(d_max);
    return 0;
}

运行结果

Thread 0: old = 0, new = not updated
Thread 1: old = 0, new = updated
Thread 2: old = 1, new = updated
Thread 3: old = 2, new = updated
Thread 4: old = 3, new = updated
Thread 5: old = 4, new = updated
Thread 6: old = 5, new = updated
Thread 7: old = 6, new = updated
Thread 8: old = 7, new = updated
Thread 9: old = 8, new = updated
Final max value: 9