CUDA 程序中的同步

  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 }