详解cuda by example中第九章原子性代码

这两天学到第九章了，发现第九章内容是关于并行线程原子性的问题，这章代码相对简单，主要在于理解共享内存的使用，于是就把我对第九章的代码理解写一下，仅供学习，如果大家有发现不对的地方欢迎指正。（文章代码顺序按照官方示例代码顺序给出，可按顺序食用）

原子性

这个名词我想如果学过计算机操作系统的同学应该不会陌生（至少我考研408这个词见🤮了），简单来说原子性就是指某种操作在执行时不能被打断，必须在执行完这个操作后才能进行其他操作。

代码讲解

#define SIZE    (100*1024*1024)

定义了SIZE的大小也就是我们要计算的直方图中数据的个数

核函数

__global__ void histo_kernel(unsigned char* buffer,
    long size,
    unsigned int* histo) {

    // clear out the accumulation buffer called temp
    // since we are launched with 256 threads, it is easy
    // to clear that memory with one write per thread
    __shared__  unsigned int temp[256];
    temp[threadIdx.x] = 0;
    __syncthreads();

    // calculate the starting index and the offset to the next
    // block that each thread will be processing
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    while (i < size) {
        atomicAdd(&temp[buffer[i]], 1);
        i += stride;
    }
    // sync the data from the above writes to shared memory
    // then add the shared memory values to the values from
    // the other thread blocks using global memory
    // atomic adds
    // same as before, since we have 256 threads, updating the
    // global histogram is just one write per thread!
    __syncthreads();
    atomicAdd(&(histo[threadIdx.x]), temp[threadIdx.x]);
}

这个函数需要三个参数分别为（传入的数据，总数据的大小这里等于我们定义的SIZE，以及直方图的大小这里我们指取值的范围为256）

    __shared__  unsigned int temp[256];
    temp[threadIdx.x] = 0;
    __syncthreads();

这段代码定义了一个大小为256的共享内存区（这里定义为256是因为每个线程块中有256个线程），并将内存区的数值初始化为0，__syncthreads()表示同步同一个线程块（block）内的所有线程。

    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    while (i < size) {
        atomicAdd(&temp[buffer[i]], 1);
        i += stride;
    }

这里面的i计算每个线程的序号，stride的大小为一个线程核中所有线程的总数。while循环表示只要线程序号不超过总数据的大小，那么就让此线程取出buffer[i]中的数据，并在存储他数量的数组的相对位置中加1（整个while循环相当于让每个线程遍历整个数据中他所对应的部分并计算每个数据出现的次数）

    __syncthreads();
    atomicAdd(&(histo[threadIdx.x]), temp[threadIdx.x]);

同步后，将这个线程此次计算的数据与之前计算的数据累加

main函数

 unsigned char* buffer =(unsigned char*)big_random_block(SIZE);

初始化数据

    cudaEvent_t     start, stop;
    HANDLE_ERROR(cudaEventCreate(&start));
    HANDLE_ERROR(cudaEventCreate(&stop));
    HANDLE_ERROR(cudaEventRecord(start, 0));

初始化计算时间的函数

    unsigned char* dev_buffer;
    unsigned int* dev_histo;
    HANDLE_ERROR(cudaMalloc((void**)&dev_buffer, SIZE));
    HANDLE_ERROR(cudaMemcpy(dev_buffer, buffer, SIZE,
        cudaMemcpyHostToDevice));

代码初始化一个大小为SIZE，并且内容为buffer的在GPU上的数组dev_buffer

    HANDLE_ERROR(cudaMalloc((void**)&dev_histo,
        256 * sizeof(int)));
    HANDLE_ERROR(cudaMemset(dev_histo, 0,
        256 * sizeof(int)));

这段代码在GPU上初始化histo数组（用来计算每个数据出现的次数）大小为256（因为数据有256个取值）

    cudaDeviceProp  prop;
    HANDLE_ERROR(cudaGetDeviceProperties(&prop, 0));
    int blocks = prop.multiProcessorCount;
    histo_kernel << <blocks * 2, 256 >> > (dev_buffer,
        SIZE, dev_histo);

    unsigned int    histo[256];
    HANDLE_ERROR(cudaMemcpy(histo, dev_histo,
        256 * sizeof(int),
        cudaMemcpyDeviceToHost));

这段代码获取电脑上第一个gpu设备的属性，获取其能它返回 GPU 上多处理器（Multiprocessor, 简称为 SM）的数量
执行后将计算的结果复制到cpu上

结尾

之后的代码是用来验证gpu计算的结果是否准确以及释放资源这里就不再阐述。这篇文章大体也差不多了，如果有错误还请大家指正。