CUDA 编程之对原子操作与并行性的理解

相关概念

CUDA 的原子操作可以理解为对一个变量进行“读取-修改-写入”这三个操作的一个最小单位的执行过程，这个执行过程不能够再分解为更小的部分，在它执行过程中，不允许其他并行线程对该变量进行读取和写入的操作。基于这个机制，原子操作实现了对在多个线程间共享的变量的互斥保护，确保任何一次对变量的操作的结果的正确性。

原子操作确保了在多个并行线程间共享的内存的读写保护，每次只能有一个线程对该变量进行读写操作，一个线程对该变量操作的时候，其他线程如果也要操作该变量，只能等待前一线程执行完成。原子操作确保了安全，代价是牺牲了性能。

原子操作分为很多，此处以 atomicAdd() 为例，做一些并行的加法。

代码编写与测试

以下代码实现这样的功能：

• 定义一百万个线程，并且线程块的大小为一千。从而有一千个线程块。
• 在全局 memory 中开辟一段大小为ARRAY_BYTES的空间用来放置数据（初始化为0）。
• 令一百万个线程去全局 memory 中按线程 ID 取数据（线程 ID 要对数组大小取余，从而保证有些线程 ID 过大的线程也会取到数据），并且进行累加（+1）。
• 对运行过程进行计时。

#include <stdio.h>
#include "gputimer.h"

#define NUM_THREADS 1000000
#define ARRAY_SIZE  100

#define BLOCK_WIDTH 1000

void print_array(int *array, int size)
{
   
   
    printf("{ ");
    for (int i = 0; i < size; i++)  {
   
    printf("%d ", array[i]); }
    printf("}\n");
}

__global__ void increment_naive(