CUDA 标准编程模式

前言

  本文将介绍 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 的是参数本身,还是参数的地址,这很关键。

posted on 2014-05-27 20:56  空山悟  阅读(890)  评论(0编辑  收藏

导航