CUDA C实践

未经本人同意，不得私自转载，本文只做学习之用。

PS:我们常写python的，可能对c不太熟悉，但是科班出身的/算法工程师，自己去学。我这不做基础科普。

PPS：感觉越学越多，我分好几章来写吧。争取每章字数一万字左右，减少阅读和学习难度。

简单了解下背景，

GPU的设计初衷是为了加速应用程序中的图形绘制运算，因此开发人员在访问GPU的时候，必须通过OpenGL或者DirectX等API来访问。

这里就存在了一定的难度，我们希望专业负责的人在垂直领域能力越强越好，但是根据现在的需求，开发人员掌握一定的图形编程只是，而且要想办法将计算问题转换为图形计算问题。

其次，GPU与多核CPU在计算架构上有着明显的区别，GPU更注重并行数据计算，但是对并行计算中的互斥性、同步性以及原子性等方面，又存在问题。

OK，找到了问题了，如何解决？

CUDA！！！

CUDA架构专门为GPU计算设计了一种全新的架构，开发人员可以通过CUDA C对GPU编程。而且，CUDA C是对标准C的一种简单扩展，学习和使用起来都非常容易~~

笔者自己是在云服务器，linux环境下开发。

（正经人谁用window写算法？）

大多数的linux发行版本都会带有一个GNU C编译器，也就是我们常说的gcc

这也是我们在跑算法常遇到的一个问题：

gcc和glibc的版本问题。

来来来，先hello先hello

有几个点我先放在这里：

a) 写一段cuda c的代码

b) 了解为主机（Host）编写的代码----为设备（Device）编写的代码， 两者之间的区别和联系。

c) 如何从主机上运行设备代码

d) 如何使用设备内存

e）如何获取或者查询cuda的设备信息。

这里主要是唤醒大家的记忆

在linux下执行，需要依次执行

然后我们就可以看到命令行出现了，hello..

这是一个最基础的C代码。

但是这里展示的是

我们将CPU和系统的内存称为主机，把GPU和内存称为设备。

那么同样的代码，我们如何放在GPU（设备）上来执行？

我们引出了第一个概念：

“在GPU设备上执行的函数通常被称为核函数-Kernel”

先看代码

看到第一行，是不是以为是C++？ 我们把文件的后缀从.c 变成了.cpp

实际上不是，实际上是.cu文件

怎么运行呢？

先说几个问题：

看完3.3你再回来。

我们继续：

cuda c为标准c增加了__global__修饰符，作用是告诉编译器，函数应该编译为设备而不是在主机上运行。

函数kernel（）由编译设备代码的编译器执行（GPU）

main（）函数则交给主机编译器（cpu）

kernel()的调用究竟代表什么含义？为什么必须加上尖括号和两个数值？

cuda c的优势在于它提供了与c在语言级别上的集成，因此这个设备函数调用非常像主机函数调用。

尖括号表示将一些参数传递给运行时的系统，这些参数并不是传递给设备代码的参数，而是告诉运行时如何启动设备的代码，这一点会后面再讲。

简单来说，这个奇怪的函数：实际上表达的是调用设备（GPU）代码

要在linux上运行CUDA程序，首先得先装环境。。

1. 安装CUDA Toolkit：首先需要在Linux系统上安装NVIDIA CUDA Toolkit。通过访问NVIDIA官方网站

按照官方文档中的说明进行安装。

2. 选择正确的编译器：CUDA程序需要使用特定的编译器进行构建。默认情况下，CUDA Toolkit会安装NVIDIA提供的`nvcc`编译器。请确保系统中已安装CUDA Toolkit并配置了正确的环境变量，以便在终端中使用`nvcc`命令。

大概是这样子，敲入“nvcc -V”

3. 编写CUDA程序：CUDA程序保存为以`.cu`为扩展名的文件。

4. 编译CUDA程序：在终端中并使用`nvcc`编译器进行编译。使用以下命令：

注意的是CUDA程序需要具有兼容的GPU设备才能运行

另外，要在CUDA程序中使用GPU并行计算，需要正确设置并行计算的配置，代码中``表示仅使用1个线程块和1个线程。

我们继续深入研究，

核函数核函数，本质是一个函数，作为函数，应该是可以传参数。那我们也讲了，因为是C，因此在内存使用这一块，有更严苛的要求。

我们先来看一段进阶代码：

HelloWorldPlus

在这段代码里面，我定义了一个核函数，实现的功能是c=a+b

那么在这里，我们需要注意两个问题：

1。核函数的传参与C函数传参类似。

2。变量在使用之前，，记得要申请内存，malloc操作，当然，谁申请，谁释放，别忘了使用完之后Free。

在这里我们先看看cudaMalloc（）

第一个参数是一个指针：指向用于保存新分配内存地址的变量。

第二个参数是分配内存的大小。

这个函数的定义和c里面的malloc（）是相同的，并且返回类型为void*

HANDLE_ERROR()是一个宏定义。作用只是判断函数调用是否返回了一个错误值。

我把book.h放在后面。

这段代码抛出了我们在学习CUDA C时候的第一个重点问题：

CUDA C为什么简单，为什么功能强大，是因为它淡化了主机代码和设备代码之间的差异。

但是

但是

但是！

程序猿一定不能在主句代码中对cudaMalloc()返回的至臻进行解引用(Dereference).

这也就是说：主机代码可以将这个指针作为参数传递，对其执行算术运算，甚至可以将其转换为另一种不同的类型。

但是，绝对不可以使用这个指针来读取或者写入内存。

更烦人的是，编译器无法防止这种错误、、、

1。 可以将cudaMalloc（）分配的指针传递给在设备上执行的函数

2。可以在设备代码中使用cudaMalloc（）分配的指针进行内存读/写操作。

3。可以将cudaMalloc（）分配的至臻传递给在主机上执行的函数。

4。不能在主机代码中使用cudaMalloc（）分配的至臻进行内存读/写操作。

简单来讲，不能用标准C的free（）函数来施放cudaMalloc（）分配的内容。

那么如果是这样的话，就会涉及到一个问题：

在主机上不能对设备商的内存做任何修改，我们应该如何使用呢？

答：在设备代码中使用设备指针和调用cudaMemcpy()

我们总结一下：

主机指针只能访问主机代码中的内存，设备只能访问设备代码中的内存。0