cuda线程索引 cuda __syncthreads

转载

技术领航者之声 2024-03-20 20:14:56

文章标签 cuda线程索引 CUDA 数组 #include 文章分类 数据仓库大数据

前言

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

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

　　至于块之间的同步，需要使用到 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 <iostream>
 10 #include <cstdlib>
 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<deviceCount; i++)
 37     {
 38         cudaDeviceProp prop;
 39         if (cudaGetDeviceProperties(&prop,i)==cudaSuccess) // 获取设备属性
 40         {
 41             if (prop.major>=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<N; i+=BLOCK_data*THREAD_data)
 73     {
 74         shared[tid] += data[i];
 75     }
 76    
 77     // 块内线程同步函数
 78     __syncthreads ();
 79 
 80     // 每个块内索引为 0 的线程对其组内所有线程的求和结果再次求和
 81     if (tid == 0) {
 82         for(int i = 1; i < THREAD_data; i++) { 
 83             shared[0] += shared[i];
 84         }
 85         // result 数组存放各个块的计算结果
 86         result[bid] = shared[0];
 87     }
 88 }
 89 
 90 int main ()
 91 {
 92     // 初始化 CUDA 编译环境
 93     if (InitCUDA()) {
 94         return EXIT_FAILURE;
 95     }
 96     cout << "成功建立 CUDA 计算环境" << endl << endl;
 97 
 98     // 建立，初始化，打印测试数组
 99     int *data = new int [N];
100     cout << "测试矩阵: " << endl;
101     for (int i=0; i<N; i++)
102     {
103         data[i] = rand()%10;
104         cout << data[i] << " ";
105         if ((i+1)%10 == 0) cout << endl;
106     }
107     cout << endl;
108 
109     int *gpudata, *result; 
110     
111     // 在显存中为计算对象开辟空间
112     cudaMalloc ((void**)&gpudata, sizeof(int)*N); 
113     // 在显存中为结果对象开辟空间
114     cudaMalloc ((void**)&result, sizeof(int)*BLOCK_data);
115     
116     // 将数组数据传输进显存
117     cudaMemcpy (gpudata, data, sizeof(int)*N, cudaMemcpyHostToDevice); 
118     // 调用 kernel 函数 - 此函数可以根据显存地址以及自身的块号，线程号处理数据。
119     Sum<<<BLOCK_data,THREAD_data,THREAD_data*sizeof (int)>>> (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<BLOCK_data; i++)
133     {
134         final_sum += sumArray[i];
135     }
136 
137     cout << "GPU 求和结果为: " << final_sum << endl;
138 
139     // 使用 CPU 对矩阵进行求和并将结果对照
140     final_sum = 0;
141     for (int i=0; i<N; i++)
142     {
143         final_sum += data[i];
144     }
145     cout << "CPU 求和结果为: " << final_sum << endl;
146 
147     getchar();
148 
149     return 0;
150 }

运行结果

cuda线程索引 cuda __syncthreads_CUDA