前言
本文将介绍 CUDA 编程的基本模式，所有 CUDA 程序都基于此模式编写，即使是调用库，库的底层也是这个模式实现的。
模式描述
1. 定义需要在 device 端执行的核函数。( 函数声明前加 _golbal_ 关键字 )
2. 在显存中为待运算的数据以及需要存放结果的变量开辟显存空间。( cudaMalloc 函数实现 )
3. 将待运算的数据传输进显存。( cudaMemcpy，cublasSetVector 等函数实现 )
4. 调用 device 端函数，同时要将需要为 device 端函数创建的块数线程数等参数传递进 <<<>>>。( 注: <<<>>>下方编译器可能显示语法错误，不用管 )
5. 从显存中获取结果变量。( cudaMemcpy，cublasGetVector 等函数实现 )
6. 释放申请的显存空间。( cudaFree 实现 )
PS：每个 device 端函数在被调用时都能获取到调用它的具体块号，线程号，从而实现并行( 获取方法请参考下面的编程规范说明以及代码示例 )。
编程规范说明
在 CUDA 标准编程模式中，增加了一些编程规范，在这里简要说明：
函数声明关键字：
1. __device__
表明此函数只能在 GPU 中被调用，在 GPU 中执行。这类函数只能被 __global__ 类型函数或 __device__ 类型函数调用。
2. __global__
表明此函数在 CPU 上调用，在 GPU 中执行。这也是以后会常提到的 "内核函数"，有时为了便于理解也称 "device" 端函数。
3. __host__
表明此函数在 CPU 上调用和执行，这也是默认情况。
内核函数配置运算符 <<<>>> - 这个运算符在调用内核函数的时候使用，一般情况下传递进三个参数：
1. 块数
2. 线程数
3. 共享内存大小 (此参数默认为0 )
内核函数中的几个系统变量 - 这几个变量可以在内核函数中使用，从而控制块与线程的工作：
1. gridDim：块数
2. blockDim：块中线程数
3. blockIdx：块编号 (0 - gridDim-1)
4. threadIdx：线程编号 (0 - blockDim-1)
知道这些已经足够编写 CUDA 程序了，更多的编程说明将在以后的文章中介绍。
代码示例
该程序采用 CUDA 并行化思想来对数组进行求和 (代码下方如果出现红色波浪线无视之)：
1 // 相关 CUDA 库
2 #include "cuda_runtime.h"
3 #include "cuda.h"
4 #include "device_launch_parameters.h"
5
6 #include <iostream>
7 #include <cstdlib>
8
9 using namespace std;
10
11 const int N = 100;
12
13 // 块数
14 const int BLOCK_data = 3;
15 // 各块中的线程数
16 const int THREAD_data = 10;
17
18 // CUDA初始化函数
19 bool InitCUDA()
20 {
21 int deviceCount;
22
23 // 获取显示设备数
24 cudaGetDeviceCount (&deviceCount);
25
26 if (deviceCount == 0)
27 {
28 cout << "找不到设备" << endl;
29 return EXIT_FAILURE;
30 }
31
32 int i;
33 for (i=0; i<deviceCount; i++)
34 {
35 cudaDeviceProp prop;
36 if (cudaGetDeviceProperties(&prop,i)==cudaSuccess) // 获取设备属性
37 {
38 if (prop.major>=1) //cuda计算能力
39 {
40 break;
41 }
42 }
43 }
44
45 if (i==deviceCount)
46 {
47 cout << "找不到支持 CUDA 计算的设备" << endl;
48 return EXIT_FAILURE;
49 }
50
51 cudaSetDevice(i); // 选定使用的显示设备
52
53 return EXIT_SUCCESS;
54 }
55
56 // 此函数在主机端调用，设备端执行。
57 __global__
58 static void Sum (int *data,int *result)
59 {
60 // 取得线程号
61 const int tid = threadIdx.x;
62 // 获得块号
63 const int bid = blockIdx.x;
64
65 int sum = 0;
66
67 // 有点像网格计算的思路
68 for (int i=bid*THREAD_data+tid; i<N; i+=BLOCK_data*THREAD_data)
69 {
70 sum += data[i];
71 }
72
73 // result 数组存放各个线程的计算结果
74 result[bid*THREAD_data+tid] = sum;
75 }
76
77 int main ()
78 {
79 // 初始化 CUDA 编译环境
80 if (InitCUDA()) {
81 return EXIT_FAILURE;
82 }
83 cout << "成功建立 CUDA 计算环境" << endl << endl;
84
85 // 建立，初始化，打印测试数组
86 int *data = new int [N];
87 cout << "测试矩阵: " << endl;
88 for (int i=0; i<N; i++)
89 {
90 data[i] = rand()%10;
91 cout << data[i] << " ";
92 if ((i+1)%10 == 0) cout << endl;
93 }
94 cout << endl;
95
96 int *gpudata, *result;
97
98 // 在显存中为计算对象开辟空间
99 cudaMalloc ((void**)&gpudata, sizeof(int)*N);
100 // 在显存中为结果对象开辟空间
101 cudaMalloc ((void**)&result, sizeof(int)*BLOCK_data*THREAD_data);
102
103 // 将数组数据传输进显存
104 cudaMemcpy (gpudata, data, sizeof(int)*N, cudaMemcpyHostToDevice);
105 // 调用 kernel 函数 - 此函数可以根据显存地址以及自身的块号，线程号处理数据。
106 Sum<<<BLOCK_data,THREAD_data,0>>> (gpudata,result);
107
108 // 在内存中为计算对象开辟空间
109 int *sumArray = new int[THREAD_data*BLOCK_data];
110 // 从显存获取处理的结果
111 cudaMemcpy (sumArray, result, sizeof(int)*THREAD_data*BLOCK_data, cudaMemcpyDeviceToHost);
112
113 // 释放显存
114 cudaFree (gpudata);
115 cudaFree (result);
116
117 // 计算 GPU 每个线程计算出来和的总和
118 int final_sum=0;
119 for (int i=0; i<THREAD_data*BLOCK_data; i++)
120 {
121 final_sum += sumArray[i];
122 }
123
124 cout << "GPU 求和结果为: " << final_sum << endl;
125
126 // 使用 CPU 对矩阵进行求和并将结果对照
127 final_sum = 0;
128 for (int i=0; i<N; i++)
129 {
130 final_sum += data[i];
131 }
132 cout << "CPU 求和结果为: " << final_sum << endl;
133
134 getchar();
135
136 return 0;
137 }运行测试
PS：矩阵元素是随机生成的
小结
1. 掌握本节知识的关键除了要掌握各个API，还要深刻理解内核函数中的块及线程变量的控制，或者说施展 :)
2. 一定要明确传递进 API 的是参数本身，还是参数的地址，这很关键。
