• 操作系统：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	int*	一个指向设备内存中整型变量的指针。这是你要进行原子最大值比较的目标地址。
val	int	要和目标地址值比较的整数值。如果 val > *address，则更新为 val。

返回值

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

代码示例

#include <cuda_runtime.h>
#include <cstdio>__global__ void kernel(int* max_value) {int tid = threadIdx.x + blockIdx.x * blockDim.x;// 使用 CUDA 内建的 atomicMaxint 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