GPU编程自学9 —— 原子操作

深度学习的兴起，使得多线程以及GPU编程逐渐成为算法工程师无法规避的问题。这里主要记录自己的GPU自学历程。

目录

• 《GPU编程自学1 —— 引言》
• 《GPU编程自学2 —— CUDA环境配置》
• 《GPU编程自学3 —— CUDA程序初探》
• 《GPU编程自学4 —— CUDA核函数运行参数》
• 《GPU编程自学5 —— 线程协作》
• 《GPU编程自学6 —— 函数与变量类型限定符》
• 《GPU编程自学7 —— 常量内存与事件》
• 《GPU编程自学8 —— 纹理内存》
• 《GPU编程自学9 —— 原子操作》
• 《GPU编程自学10 —— 流并行》

九、 原子操作

原子操作 是指对全局和共享内存中的32位或者64位数据进行 “读取-修改-覆写”这一操作。

原子操作可以看作是一种最小单位的执行过程。 在其执行过程中，不允许其他并行线程对该变量进行读取和写入的操作。 如果发生竞争，则其他线程必须等待。

下面先给出原子操作函数的列表，后续会给出一个应用例子。

9.1 原子操作函数列表

9.1.1 atomicAdd()

int atomicAdd(int* address, int val);unsigned int atomicAdd(unsigned int* address, unsigned int val);unsigned long long int atomicAdd(unsigned long long int* address, unsigned long long int val);float atomicAdd(float* address, float val);double atomicAdd(double* address, double val);

读取位于全局或共享存储器中地址address处的32位或64位字old，计算(old + val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

注意：

• 32位浮点数的操作只适用于计算能力大于2.0的GPU
• 64位浮点数的操作只适用于计算能力大于6.0的GPU

但可以通过以下操作在计算能力不足的GPU上实现浮点数原子操作：

#if __CUDA_ARCH__ < 600 __device__ double atomicAdd(double* address, double val) {     unsigned long long int* address_as_ull = (unsigned long long int*)address; unsigned long long int old = *address_as_ull, assumed;     do {         assumed = old;         old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));         // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)         }         while (assumed != old);         return __longlong_as_double(old); } #endif

9.1.2 atomicSub()

int atomicSub(int* address, int val); unsigned int atomicSub(unsigned int* address, unsigned int val);

读取位于全局或共享存储器中地址address处的32位字old，计算(old - val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

9.1.3 atomicExch()

int atomicExch(int* address, int val); unsigned int atomicExch(unsigned int* address, unsigned int val); unsigned long long int atomicExch(unsigned long long int* address, unsigned long long int val); float atomicExch(float* address, float val);

读取位于全局或共享存储器中地址address处的32位或64位字old，并将val 存储在存储器的同一地址中。这两项操作在一次原子事务中执行。该函数将返回old。

9.1.4 atomicMin()

int atomicMin(int* address, int val); unsigned int atomicMin(unsigned int* address, unsigned int val); unsigned long long int atomicMin(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位字或64位字old，计算old 和val 的最小值，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old>

注意：

• 64位的操作只适用于计算能力大于3.5的GPU

9.1.5 atomicMax()

同atomicMin()。

9.1.6 atomicInc()

unsigned int atomicInc(unsigned int* address, unsigned int val);

读取位于全局或共享存储器中地址address处的32位字old，计算 ((old >= val) ? 0 : (old+1))，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

9.1.7 atomicDec()

unsigned int atomicDec(unsigned int* address, unsigned int val);

读取位于全局或共享存储器中地址address处的32位字old，计算 (((old == 0) | (old > val)) ? val : (old-1))，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

9.1.8 atomicCAS()

int atomicCAS(int* address, int compare, int val); unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val); unsigned long long int atomicCAS(unsigned long long int* address, unsigned long long int compare, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位或64位字old，计算 (old == compare ? val : old)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old（比较并交换）。

9.1.9 atomicAnd()

int atomicAnd(int* address, int val); unsigned int atomicAnd(unsigned int* address, unsigned int val); unsigned long long int atomicAnd(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位字或64位字old，计算 (old & val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

注意：

• 64位的操作只适用于计算能力大于3.5的GPU

9.1.10 atomicOr()

int atomicOr(int* address, int val); unsigned int atomicOr(unsigned int* address, unsigned int val); unsigned long long int atomicOr(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位字或64位字old，计算 (old | val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

注意：

• 64位的操作只适用于计算能力大于3.5的GPU

9.1.11 atomicXor()

int atomicXor(int* address, int val); unsigned int atomicXor(unsigned int* address, unsigned int val); unsigned long long int atomicXor(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享存储器中地址address处的32位字或64位字old，计算 (old ^ val)，并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。

注意：

• 64位的操作只适用于计算能力大于3.5的GPU

9.2 直方图统计

在一开始我们就提到过，原子操作是为了保证每次只能有一个线程对变量进行读写，而其它线程必须等待。 这样可以有效地避免多个线程访问和修改一个变量带来的不确定问题。

下面我们的例子就是对一堆范围在[0 255]的数进行直方图统计。由于每遇到一个数我们就要在对应统计值处加1，因此多线程操作往同一位置加1的时候很容易出现问题。因此我们采用原子操作。

关于下面的简单程序还有几点说明:

• 为了验证GPU结果的准确性，我首先采用CPU进行直方图统计，然后在GPU中对数据进行 减计数，如果最后统计结果均为0，则说明GPU和CPU统计结果一致。
• 程序中每一个block包含256个线程，然后我在每一个block中开辟了一块共享内存temp，并将一个block中的统计结果存储到temp上，最后在对所有block的结果进行整合。这些操作主要是为了避免在全局内存上进行原子操作，否则速度会非常慢。
• 对于GPU的计算时间我没有考虑数据在主机和设备之间的通信所花费的时间。 事实上，这里的通信是一件相对很慢的工作。读者可以自己测量下整个包含通信所花费的时间。

#include <iostream>#include "cuda_runtime.h"#include "time.h"using namespace std;#define num (256 * 1024 * 1024)// 核函数// 注意，为了方便验证GPU的统计结果，这里采用了"逆直方图"，// 即每发现一个数字，就从CPU的统计结果中减1__global__ void hist(unsigned char* inputdata, int* outputhist, long size){    // 开辟共享内存，否则在全局内存采用原子操作会非常慢(因为冲突太多)    __shared__ int temp[256];    temp[threadIdx.x] = 0;    __syncthreads();    // 计算线程索引及线程偏移量    int ids = blockIdx.x * blockDim.x + threadIdx.x;    int offset = blockDim.x * gridDim.x;    while (ids < size)    {        //采用原子操作对一个block中的数据进行直方图统计        atomicAdd(&temp[inputdata[ids]],1);        ids += offset;    }    // 等待统计完成，减去统计结果    __syncthreads();    atomicSub(&outputhist[threadIdx.x], temp[threadIdx.x]);}int main(){    // 生成随机数据 [0 255]    unsigned char* cpudata = new unsigned char[num];    for (size_t i = 0; i < num; i++)        cpudata[i] = static_cast<unsigned char>(rand() % 256);    // 声明数组用于记录统计结果    int cpuhist[256];    memset(cpuhist, 0, 256 * sizeof(int));    /*******************************   CPU测试代码   *********************************/    clock_t cpu_start, cpu_stop;    cpu_start = clock();    for (size_t i = 0; i < num; i++)        cpuhist[cpudata[i]] ++;    cpu_stop = clock();    cout << "CPU time: " << (cpu_stop - cpu_start) << "ms" << endl;    /*******************************   GPU测试代码   *********************************/    //定义事件用于计时    cudaEvent_t start, stop;    cudaEventCreate(&start);    cudaEventCreate(&stop);    //开辟显存并将数据copy进显存    unsigned char* gpudata;    cudaMalloc((void**)&gpudata,num*sizeof(unsigned char));    cudaMemcpy(gpudata, cpudata, num*sizeof(unsigned char),cudaMemcpyHostToDevice);    // 开辟显存用于存储输出数据,并将CPU的计算结果copy进去    int* gpuhist;    cudaMalloc((void**)&gpuhist,256*sizeof(int));    cudaMemcpy(gpuhist, cpuhist, 256*sizeof(int), cudaMemcpyHostToDevice);    // 执行核函数并计时    cudaEventRecord(start, 0);    hist << <1024, 256 >> >(gpudata,gpuhist,num);    cudaEventRecord(stop, 0);    // 将结果copy回主机    int histcpu[256];    cudaMemcpy(cpuhist,gpuhist,256*sizeof(int),cudaMemcpyDeviceToHost);    // 销毁开辟的内存    cudaFree(gpudata);    cudaFree(gpuhist);    delete cpudata;    // 计算GPU花费时间并销毁计时事件    cudaEventSynchronize(stop);    float gputime;    cudaEventElapsedTime(&gputime, start, stop);    cudaEventDestroy(start);    cudaEventDestroy(stop);    cout << "GPU time: " << gputime << "ms" << endl;    // 验证结果    long result = 0;    for (size_t i = 0; i < 256; i++)        result += cpuhist[i];    if (result == 0)        cout << "GPU has the same result with CPU." << endl;    else        cout << "Error: GPU has a different result with CPU." << endl;    system("pause");    return 0;}

上面的执行结果为：

CPU time: 188msGPU time: 26.7367msGPU has the same result with CPU.

参考资料

• 《CUDA by Example: An Introduction to General-Purpose GPU Programming》 中文名《GPU高性能编程CUDA实战》
• “CUDA Toolkit Documentation 原子操作 ”http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions