CUDA 程序中的同步

前言

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

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

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

块内同步函数：__syncthreads ()

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

代码示例

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

// 相关 CUDA 库
#include "cuda_runtime.h"
#include "cuda.h"
#include "device_launch_parameters.h"

// 此头文件包含 __syncthreads ()函数
#include "device_functions.h"

#include <iostream>
#include <cstdlib>

using namespace std;

const int N = 100;

// 块数
const int BLOCK_data = 3; 
// 各块中的线程数
const int THREAD_data = 10; 

// CUDA初始化函数
bool InitCUDA()
{
    int deviceCount; 

    // 获取显示设备数
    cudaGetDeviceCount (&deviceCount);

    if (deviceCount == 0) 
    {
        cout << "找不到设备" << endl;
        return EXIT_FAILURE;
    }

    int i;
    for (i=0; i<deviceCount; i++)
    {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop,i)==cudaSuccess) // 获取设备属性
        {
            if (prop.major>=1) //cuda计算能力
            {
                break;
            }
        }
    }

    if (i==deviceCount)
    {
        cout << "找不到支持 CUDA 计算的设备" << endl;
        return EXIT_FAILURE;
    }

    cudaSetDevice(i); // 选定使用的显示设备

    return EXIT_SUCCESS;
}

// 此函数在主机端调用，设备端执行。
__global__ 
static void Sum (int *data,int *result)
{
    // 声明共享内存 (数组)
    extern __shared__ int shared[]; 
    // 取得线程号
    const int tid = threadIdx.x; 
    // 获得块号
    const int bid = blockIdx.x; 

    shared[tid] = 0;
    // 有点像网格计算的思路
    for (int i=bid*THREAD_data+tid; i<N; i+=BLOCK_data*THREAD_data)
    {
        shared[tid] += data[i];
    }
   
    // 块内线程同步函数
    __syncthreads ();

    // 每个块内索引为 0 的线程对其组内所有线程的求和结果再次求和
    if (tid == 0) {
        for(int i = 1; i < THREAD_data; i++) { 
            shared[0] += shared[i];
        }
        // result 数组存放各个块的计算结果
        result[bid] = shared[0];
    }
}

int main ()
{
    // 初始化 CUDA 编译环境
    if (InitCUDA()) {
        return EXIT_FAILURE;
    }
    cout << "成功建立 CUDA 计算环境" << endl << endl;

    // 建立，初始化，打印测试数组
    int *data = new int [N];
    cout << "测试矩阵: " << endl;
    for (int i=0; i<N; i++)
    {
        data[i] = rand()%10;
        cout << data[i] << " ";
        if ((i+1)%10 == 0) cout << endl;
    }
    cout << endl;

    int *gpudata, *result; 
    
    // 在显存中为计算对象开辟空间
    cudaMalloc ((void**)&gpudata, sizeof(int)*N); 
    // 在显存中为结果对象开辟空间
    cudaMalloc ((void**)&result, sizeof(int)*BLOCK_data);
    
    // 将数组数据传输进显存
    cudaMemcpy (gpudata, data, sizeof(int)*N, cudaMemcpyHostToDevice); 
    // 调用 kernel 函数 - 此函数可以根据显存地址以及自身的块号，线程号处理数据。
    Sum<<<BLOCK_data,THREAD_data,THREAD_data*sizeof (int)>>> (gpudata,result);
    
    // 在内存中为计算对象开辟空间
    int *sumArray = new int[BLOCK_data];
    // 从显存获取处理的结果
    cudaMemcpy (sumArray, result, sizeof(int)*BLOCK_data, cudaMemcpyDeviceToHost);
    
    // 释放显存
    cudaFree (gpudata); 
    cudaFree (result);

    // 计算 GPU 每个块计算出来和的总和
    int final_sum=0;
    for (int i=0; i<BLOCK_data; i++)
    {
        final_sum += sumArray[i];
    }

    cout << "GPU 求和结果为: " << final_sum << endl;

    // 使用 CPU 对矩阵进行求和并将结果对照
    final_sum = 0;
    for (int i=0; i<N; i++)
    {
        final_sum += data[i];
    }
    cout << "CPU 求和结果为: " << final_sum << endl;

    getchar();

    return 0;
}

运行结果

　　

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

小结

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