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

素颜马尾好姑娘i 2023-07-25 10:00 6阅读 0赞

# 相关概念 #

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(int *g)
    { 
    	// which thread is this?
    	int i = blockIdx.x * blockDim.x + threadIdx.x; 
    
    	// each thread to increment consecutive elements, wrapping at ARRAY_SIZE
    	i = i % ARRAY_SIZE;  
    	g[i] = g[i] + 1;
    }
    
    __global__ void increment_atomic(int *g)
    { 
    	// which thread is this?
    	int i = blockIdx.x * blockDim.x + threadIdx.x; 
    
    	// each thread to increment consecutive elements, wrapping at ARRAY_SIZE
    	i = i % ARRAY_SIZE;  
    	atomicAdd(& g[i], 1);
    }
    
    int main(int argc,char **argv)
    {    
        GpuTimer timer;
        printf("%d total threads in %d blocks writing into %d array elements\n",
               NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, ARRAY_SIZE);
    
        // declare and allocate host memory
        int h_array[ARRAY_SIZE];
        const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int);
     
        // declare, allocate, and zero out GPU memory
        int * d_array;
        cudaMalloc((void **) &d_array, ARRAY_BYTES);
        cudaMemset((void *) d_array, 0, ARRAY_BYTES); 
    
        // launch the kernel - comment out one of these
        timer.Start();
        // increment_naive<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);
        increment_atomic<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);
        timer.Stop();
        
        // copy back the array of sums from GPU and print
        cudaMemcpy(h_array, d_array, ARRAY_BYTES, cudaMemcpyDeviceToHost);
        print_array(h_array, ARRAY_SIZE);
        printf("Time elapsed = %g ms\n", timer.Elapsed());
     
        // free GPU memory allocation and exit
        cudaFree(d_array);
        return 0;
    }

如果我们按照该代码运行会出现以下结果。  
![在这里插入图片描述][2020041416355852.png]  
为什么会出现这样的结果？？？  
其实这就是原子操作在对线程的**锁定操作**——即一个线程在取数据的时候原子操作会锁定该块区域，从而令其他线程无法取该区域的数据，所以数组每个元素都是10000——表示每个“坑位”分别有10000个线程在此累加过，并且是按照顺序的。

如果我们运行\*\*increment\_naive()\*\*函数，则会发现累加结果不是10000，并且每次运行结果都不一样，并且运行时间也会少一些——**因为各个线程在以不可预知的顺序异步的访问内存区域进行累加。**

这个程序很好的体现了 CUDA 程序中线程的**并行性**。

下面附上头文件的代码：

#ifndef __GPU_TIMER_H__
    #define __GPU_TIMER_H__
    
    struct GpuTimer
    { 
          cudaEvent_t start;
          cudaEvent_t stop;
     
          GpuTimer()
          { 
                cudaEventCreate(&start);
                cudaEventCreate(&stop);
          }
     
          ~GpuTimer()
          { 
                cudaEventDestroy(start);
                cudaEventDestroy(stop);
          }
     
          void Start()
          { 
                cudaEventRecord(start, 0);
          }
     
          void Stop()
          { 
                cudaEventRecord(stop, 0);
          }
     
          float Elapsed()
          { 
                float elapsed;
                cudaEventSynchronize(stop);
                cudaEventElapsedTime(&elapsed, start, stop);
                return elapsed;
          }
    };
    
    #endif /* __GPU_TIMER_H__ */

[2020041416355852.png]: /images/20230528/a5094c9566ca4a4db9f3b48f2b561cfb.png