NVIDIA CUDA2023春训营（九）CUDA 原子操作

原子操作

CUDA 编程的基本思想利用 GPU 来尽可能地并行执行相同的核函数，对于大多数并行任务，线程间不需要合作或使用其他线程的资源，只需要保证自己能够正常执行即可

但对于某些需要同步执行的操作，例如多个核函数需要对同一个变量进行读取-修改-写入，由于核函数之间是异步的，当试图同时执行时，就会导致出现问题

CUDA 的原子操作是针对全局内存或共享内存中的变量，其是对全局变量或共享变量进行读取-修改-写入这三个操作的一个最小单位的执行过程，在这个执行过程中，不允许其他并行线程对该变量进行读取和写入

也就是说，原子操作实现了多个线程间共享的变量的互斥保护，每次只能有一个线程对全局变量或共享变量进行读写操作，能够确保任何一次对变量的操作的结果的正确性

常用函数

原子操作的常用函数如下

// 加法：value = value + num
atomicAdd(&value, num);
// 减法：value = value - num
atomicSub(&value, num);
// 赋值：value = num
atomicExch(&value, num);
// 求最大值：value = max(value, num)
atomicMax(&value, num);
// 求最小值：value = min(value, num)
atomicMin(&value, num);
// 向上计数：value = value <= num ? value + 1 : 0
atomicInc(&value, num);
// 向下计数：value = value > num || value == 0 ? value - 1 : 0
atomicDec(&value, num);
// 比较并交换：value = value == compare ? val : value
atomicCAS(&value, compare, val);
// 与运算：value = value & num
atomicAnd(&value, compare, val);
// 或运算：value = value | num
atomicOr(&value, compare, val);
// 异或运算：value = value ^ num
atomicXor(&value, compare, val);

实例

下面代码给出使用原子操作求和的实例

#include <stdio.h>
#define N 10
#define BLOCKS 32
#define BLOCK_SIZE 256

__global__ void sum_gpu(int *data, int *res) {
    int index = blockDim.x * blockIdx.x + threadIdx.x;
    if (index < N) {
        atomicAdd(res, data[index]);
    }
}

int sum_cpu(int *data) {
    int sum = 0;
    for (int i = 0; i < N; i++) {
        sum += data[i];
    }
    return sum;
}

void check(int res_cpu, int res_gpu) {
    printf("CPU 计算结果：%d\n", res_cpu);
    printf("GPU 计算结果：%d\n", res_gpu);
    printf("%s\n", res_cpu == res_gpu ? "Pass" : "Error");
}

int main() {

    // 申请统一内存
    int *data, *res;
    cudaMallocManaged(&data, sizeof(int) * N);
    cudaMallocManaged(&res, sizeof(int));

    // 初始化
    res[0] = 0;
    for (int i = 0; i < N; i++) {
        data[i] = rand() % 10;
    }

    // cpu计算
    int res_cpu = sum_cpu(data);
    
    // gpu计算
    sum_gpu<<<BLOCKS, BLOCK_SIZE>>>(data, res);
    cudaDeviceSynchronize();
    int res_gpu = res[0];

    // 检查结果
    check(res_cpu, res_gpu);
    
    // 释放统一内存
    cudaFree(data);
    cudaFree(res);

    return 0;
}