OpenCV CUDA模块设备层-----在 GPU 上执行线程安全的 原子取最小值操作函数atomicMin()

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

算法描述

这是 OpenCV 的 CUDA 模块（cv::cudev） 中封装的原子操作函数之一，用于在 GPU 上执行线程安全的 “原子取最小值” 操作。
线程安全： 在多个线程并发访问时保证数据一致性。

函数原型

__device__ __forceinline__ uint cv::cudev::atomicMin 	( 	uint *  	address,
		uint  	val 
	) 		

返回值

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

参数

• address 一个指向 设备内存（global memory） 中整型变量的指针。这是你要进行原子最小值比较的目标地址。
• val 要和目标地址值比较的整数值。如果 val < *address，则更新为 val。

代码示例

#include <opencv2/cudev/functional/functional.hpp>
#include <cuda_runtime.h>
#include <cstdio>
#include <opencv2/cudev/util/atomic.hpp>

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

    // 线程ID作为输入值尝试更新最小值
    int old = cv::cudev::atomicMin(min_value, tid);

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

int main() {
    int h_min = 9999;  // 初始一个较大的值
    int* d_min;

    cudaMalloc(&d_min, sizeof(int));
    cudaMemcpy(d_min, &h_min, sizeof(int), cudaMemcpyHostToDevice);

    // 启动 10 个线程，threadIdx.x 范围为 0~9
    kernel<<<1, 10>>>(d_min);

    cudaDeviceSynchronize();  // 等待核函数完成

    cudaMemcpy(&h_min, d_min, sizeof(int), cudaMemcpyDeviceToHost);
    printf("Final min value: %d\n", h_min);  // 应为 0

    cudaFree(d_min);
    return 0;
}

运行结果

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