前言

       在并发，多线程环境下，同步是一个很重要的环节。同步即是指进程/线程之间的执行顺序约定。

       本文将介绍如何通过共享内存机制实现块内多线程之间的同步。

       至于块之间的同步，需要使用到 global memory，代价较为高昂，目前使用的情况也不多，就先不介绍了。

块内同步函数：__syncthreads ()

       线程调用此函数后，该线程所属块中的所有线程均运行到这个调用点后才会继续往下运行。

代码示例

       使用同步思想优化之前一篇博文中提到的数组求和程序。在新的程序中，让每个块中的第一个线程将块中所有线程的运算结果都加起来，然后再存入到结果数组中。这样，结果数组的长度与块数相等 (原来是和总线程数相等)，大大降低了 CPU 端程序求和的工作量以及需要传递进/出显存的数据 (代码下方如果出现红色波浪线无视之)：

1 // 相关 CUDA 库  2 #include "cuda_runtime.h"  3 #include "cuda.h"  4 #include "device_launch_parameters.h"  5   6 // 此头文件包含 __syncthreads ()函数  7 #include "device_functions.h"  8   9 #include 
     
       10 #include 
      
        11  12 using namespace std; 13  14 const int N = 100; 15  16 // 块数 17 const int BLOCK_data = 3;  18 // 各块中的线程数 19 const int THREAD_data = 10;  20  21 // CUDA初始化函数 22 bool InitCUDA() 23 { 24     int deviceCount;  25  26     // 获取显示设备数 27     cudaGetDeviceCount (&deviceCount); 28  29     if (deviceCount == 0)  30     { 31         cout << "找不到设备" << endl; 32         return EXIT_FAILURE; 33     } 34  35     int i; 36     for (i=0; i
       
        =1) //cuda计算能力 42             { 43                 break; 44             } 45         } 46     } 47  48     if (i==deviceCount) 49     { 50         cout << "找不到支持 CUDA 计算的设备" << endl; 51         return EXIT_FAILURE; 52     } 53  54     cudaSetDevice(i); // 选定使用的显示设备 55  56     return EXIT_SUCCESS; 57 } 58  59 // 此函数在主机端调用，设备端执行。 60 __global__  61 static void Sum (int *data,int *result) 62 { 63     // 声明共享内存 (数组) 64     extern __shared__ int shared[];  65     // 取得线程号 66     const int tid = threadIdx.x;  67     // 获得块号 68     const int bid = blockIdx.x;  69  70     shared[tid] = 0; 71     // 有点像网格计算的思路 72     for (int i=bid*THREAD_data+tid; i
        
         >> (gpudata,result);120     121     // 在内存中为计算对象开辟空间122     int *sumArray = new int[BLOCK_data];123     // 从显存获取处理的结果124     cudaMemcpy (sumArray, result, sizeof(int)*BLOCK_data, cudaMemcpyDeviceToHost);125     126     // 释放显存127     cudaFree (gpudata); 128     cudaFree (result);129 130     // 计算 GPU 每个块计算出来和的总和131     int final_sum=0;132     for (int i=0; i
        
       
      
     

运行结果

　　

　　PS：矩阵元素是随机生成的

小结

　　共享内存，或者说这个共享数组是 CUDA 中实现同步最常用的方法。