3

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

 1 year ago
source link: https://alex-mcavoy.github.io/nvidia/cuda-spring-bootcamp/7acd409c.html
Go to the source link to view the article. You can view the picture content, updated content and better typesetting reading experience. If the link is broken, please click the button below to view the snapshot at that time.
neoserver,ios ssh client

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

2023-02-07NVIDIA ， CUDA春训营

3 2k 2 分钟

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;
}

report

Recommend

About Joyk


Aggregate valuable and interesting links.
Joyk means Joy of geeK