前言

       本文将介绍 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 
     
        7 #include 
      
         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
       
        =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
        
         >> (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
        
       
      
     

运行测试

      

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

小结

       1. 掌握本节知识的关键除了要掌握各个API，还要深刻理解内核函数中的块及线程变量的控制，或者说施展 :) 

       2. 一定要明确传递进 API 的是参数本身，还是参数的地址，这很关键。