本文共 7101 字，大约阅读时间需要 23 分钟。
CUDA编程（二）
CUDA初始化与核函数
CUDA初始化
在上一次中已经说过了，CUDA安装成功之后，新建一个工程还是十分简单的，直接在新建项目的时候选择NVIDIA CUDA项目就可以了，我们先新建一个MyCudaTest 工程，删掉自带的示例kernel.cu，然后新建项，新建一个CUDA C/C++ File ,我们首先看一下如何初始化CUDA，因此我命名为InitCuda.cu
首先我们要使用CUDA的RunTime API 所以 我们需要include cuda_runtime.h
#include//CUDA RunTime API#include
接下来这个函数会调用 runtime API 中 有关初始化CUDA的内容
//CUDA 初始化bool InitCUDA(){ int count; //取得支持Cuda的装置的数目 cudaGetDeviceCount(&count); //没有符合的硬件 if (count == 0) { fprintf(stderr, "There is no device.\n"); return false; } int i; for (i = 0; i < count; i++) { cudaDeviceProp prop; if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) { if (prop.major >= 1) { break; } } } if (i == count) { fprintf(stderr, "There is no device supporting CUDA 1.x.\n"); return false; } cudaSetDevice(i); return true;} 这段程序首先会调用cudaGetDeviceCount 函数，获得支持 CUDA 的GPU的数量，如果计算机上没有支持 CUDA 的装置，则会传回 1，而这个1是device 0 ，device0 只是一个仿真装置，但是CUDA的很多功能都不支持（不支持CUDA1.0以上版本），因此我们要真正确定系统上是否有支持CUDA的装置，需要对每个device调用cudaGetDeviceProperties，来获得它们的具体参数，以及所支持的CUDA版本（prop.major 和 prop.minor 分别代表装置支持的版本号码，例如 6.5 则 prop.major 为 6 而prop.minor 为 5）
cudaGetDeviceProperties除了可以获得装置支持的 CUDA 版本之外，还有装置的名称、内存的大小、最大的 thread 数目、执行单元的频率等等。详情可参考NVIDIA 的 CUDA Programming Guide。
在找到支持 CUDA 1.0 以上的装置之后，就可以呼叫 cudaSetDevice 函式，把它设为目前要使用的显卡。
下面我们在Main函数中调用InitCUDA函数，由于我们使用VS，所以直接ctrl+F5编译执行就可以了，执行时如果系统上有支持 CUDA 的装置，应该会显示 CUDA initialized。
int main() { if (!InitCUDA()) { return 0; } printf("CUDA initialized.\n"); return 0;} 
完整程序：
#include//CUDA RunTime API#include //CUDA 初始化bool InitCUDA(){ int count; //取得支持Cuda的装置的数目 cudaGetDeviceCount(&count); //没有符合的硬件 if (count == 0) { fprintf(stderr, "There is no device.\n"); return false; } int i; for (i = 0; i < count; i++) { cudaDeviceProp prop; if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) { if (prop.major >= 1) { break; } } } if (i == count) { fprintf(stderr, "There is no device supporting CUDA 1.x.\n"); return false; } cudaSetDevice(i); return true;}int main() { if (!InitCUDA()) { return 0; } printf("CUDA initialized.\n"); return 0;}
CUDA核函数
完成了CUDA的初始化检查操作，下面我们就可以使用CUDA完成一些简单计算了，这里我们打算计算一系列数字的立方和。
所以我们先写了一个随机函数：
#define DATA_SIZE 1048576int data[DATA_SIZE];//产生大量0-9之间的随机数void GenerateNumbers(int *number, int size){ for (int i = 0; i < size; i++) { number[i] = rand() % 10; }}//生成随机数(main中调用)//GenerateNumbers(data, DATA_SIZE); 该函数会产生一大堆 0 ~ 9 之间的随机数，然后我们要对他们进行立方和操作。
那么我们如何让这个工作在显卡上完成呢？首先第一件事很显而易见，这些数字不能放在内存里了，而是要复制到GPU的显存上。下面我们就来看一下数据复制的部分。
Host&Device架构：
上一次已经讲过关于CUDA架构的一些基础了，这里再稍微复习一下，在 CUDA 的架构下，一个程序分为两个部份：host 端和 device 端。Host 端是指在 CPU 上执行的部份，而 device 端则是在显示芯片上执行的部份。Device 端的程序又称为 “kernel”。通常 host 端程序会将数据准备好后，复制到显卡的内存中，再由显示芯片执行 device 端程序，完成后再由 host 端程序将结果从显卡的内存中取回。
我们需要把产生的数据复制到Device端的RAM，才能在显卡上完成计算，因此我们首先开辟一块合适的显存，然后把随机数从内存复制进去。
//生成随机数 GenerateNumbers(data, DATA_SIZE); /*把数据复制到显卡内存中*/ int* gpudata, *result; //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果 ) cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE); cudaMalloc((void**)&result, sizeof(int)); //cudaMemcpy 将产生的随机数复制到显卡内存中 //cudaMemcpyHostToDevice - 从内存复制到显卡内存 //cudaMemcpyDeviceToHost - 从显卡内存复制到内存 cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE,cudaMemcpyHostToDevice);
注释已经写得比较明白了。cudaMalloc 和 cudaMemcpy 的用法和一般的 malloc 及 memcpy 类似，不过 cudaMemcpy 则多出一个参数，指示复制内存的方向。在这里因为是从主内存复制到显卡内存，所以使用 cudaMemcpyHostToDevice。如果是从显卡内存到主内存，则使用cudaMemcpyDeviceToHost。
完成了从内存到显存的数据拷贝之后，我们接下来就要在显卡上完成计算了，如何让程序跑在显卡上？答案是核函数。
CUDA核函数：
要写在显示芯片上执行的程序。在 CUDA 中，在函数前面加上__global__ 表示这个函式是要在显示芯片上执行的,所以我们只要在正常函数之前加上一个__global__就行了：
// __global__ 函数 (GPU上执行) 计算立方和__global__ static void sumOfSquares(int *num, int* result){ int sum = 0; int i; for (i = 0; i < DATA_SIZE; i++) { sum += num[i] * num[i] * num[i]; } *result = sum;} 在显示芯片上执行的程序有一些限制，首先最明显的一个限制——不能有传回值，还有一些其他的限制，后面会慢慢提到。
执行核函数：
写好核函数之后需要让CUDA执行这个函数。
在 CUDA 中，要执行一个核函数，使用以下的语法：
函数名称<<>>(参数...);
这里我们先不去并行，只是单纯地完成GPU计算，所以我们让block = 1，thread = 1，share memory = 0
sumOfSquares<<<1, 1, 0>>>(gpudata, result);
计算完了，千万别忘了还要把结果从显示芯片复制回主内存上，然后释放掉内存~
int sum; //cudaMemcpy 将结果从显存中复制回内存 cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost); //Free cudaFree(gpudata); cudaFree(result);
最后我们把结果打印出来就大功告成了：
printf("GPUsum: %d \n", sum); 之后我们再用CPU计算一下来验证一下上面的过程是否有错，这一步还是十分必要的：
sum = 0; for (int i = 0; i < DATA_SIZE; i++) { sum += data[i] * data[i] * data[i]; } printf("CPUsum: %d \n", sum); 完整程序：
程序代码：
#include#include //CUDA RunTime API#include #define DATA_SIZE 1048576int data[DATA_SIZE];//产生大量0-9之间的随机数void GenerateNumbers(int *number, int size){ for (int i = 0; i < size; i++) { number[i] = rand() % 10; }}//CUDA 初始化bool InitCUDA(){ int count; //取得支持Cuda的装置的数目 cudaGetDeviceCount(&count); if (count == 0) { fprintf(stderr, "There is no device.\n"); return false; } int i; for (i = 0; i < count; i++) { cudaDeviceProp prop; if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) { if (prop.major >= 1) { break; } } } if (i == count) { fprintf(stderr, "There is no device supporting CUDA 1.x.\n"); return false; } cudaSetDevice(i); return true;}// __global__ 函数 (GPU上执行) 计算立方和__global__ static void sumOfSquares(int *num, int* result){ int sum = 0; int i; for (i = 0; i < DATA_SIZE; i++) { sum += num[i] * num[i] * num[i]; } *result = sum;}int main(){ //CUDA 初始化 if (!InitCUDA()) { return 0; } //生成随机数 GenerateNumbers(data, DATA_SIZE); /*把数据复制到显卡内存中*/ int* gpudata, *result; //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果 ) cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE); cudaMalloc((void**)&result, sizeof(int)); //cudaMemcpy 将产生的随机数复制到显卡内存中 //cudaMemcpyHostToDevice - 从内存复制到显卡内存 //cudaMemcpyDeviceToHost - 从显卡内存复制到内存 cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice); // 在CUDA 中执行函数 语法：函数名称<< >>(参数...); sumOfSquares << <1, 1, 0 >> >(gpudata, result); /*把结果从显示芯片复制回主内存*/ int sum; //cudaMemcpy 将结果从显存中复制回内存 cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost); //Free cudaFree(gpudata); cudaFree(result); printf("GPUsum: %d \n", sum); sum = 0; for (int i = 0; i < DATA_SIZE; i++) { sum += data[i] * data[i] * data[i]; } printf("CPUsum: %d \n", sum); return 0;}
运行结果：
总结：
这次给大家介绍了CUDA的初始化和如何在显卡上运行程序，即先将数据从内存复制到显存，再写好运算的核函数，之后用CUDA调用核函数，完成GPU上的计算，之后当然不要忘记将结果复制回内存，释放掉显存。
总的来说一个CUDA程序的骨架已经搭建起来了，而GPU计算的重中之重即并行加速还没有进行介绍，不过在加速之前我们还有一件非常重要的事情需要考虑，那就是我们的程序到底有没有加速，也就是我们要输出程序的运行时间，这个时间我们需要使用CUDA提供的一个Clock函数，可以取得GPU执行单元的频率，所以下一篇博客我将主要讲解这个函数~希望能给大家的学习带来帮助~
参考资料：《深入浅出谈CUDA》