《GPU高性能编程CUDA实战》中代码整理

转自：点击打开链接

CUDA架构专门为GPU计算设计了一种全新的模块，目的是减轻早期GPU计算中存在的一些限制，而正是这些限制使得之前的GPU在通用计算中没有得到广泛的应用。

         使用CUDA C来编写代码的前提条件包括：(1)、支持CUDA的图形处理器，即由NVIDIA推出的GPU显卡，要求显存超过256MB；(2)、NVIDIA设备驱动程序，用于实现应用程序与支持CUDA的硬件之间的通信，确保安装最新的驱动程序，注意选择与开发环境相符的图形卡和操作系统；(3)、CUDA开发工具箱即CUDA Toolkit，此工具箱中包括一个编译GPU代码的编译器；(4)、标准C编译器，即CPU编译器。CUDA C应用程序将在两个不同的处理器上执行计算，因此需要两个编译器。其中一个编译器为GPU编译代码，而另一个为CPU编译代码。 

         一般，将CPU以及系统的内存称为主机(Host)，而将GPU及其内存称为设备(Device)。在GPU设备上执行的函数通常称为核函数(Kernel)。

         cudaMalloc函数使用限制总结：(1)、可以将cudaMalloc()分配的指针传递给在设备上执行的函数；(2)、可以在设备代码中使用cudaMalloc()分配的指针进行内存读/写操作；(3)、可以将cudaMalloc()分配的指针传递给在主机上执行的函数；(4)、不能在主机代码中使用cudaMalloc()分配的指针进行内存读/写操作。

         不能使用标准C的free()函数来释放cudaMalloc()分配的内存；要释放cudaMalloc()分配的内存，需要调用cudaFree()。

         设备指针的使用方式与标准C中指针的使用方式完全一样。主机指针只能访问主机代码中的内存，而设备指针也只能访问设备代码中的内存。

         在主机代码中可以通过调用cudaMemcpy()来访问设备上的内存。

         有可能在单块卡上包含了两个或多个GPU。

         在集成的GPU上运行代码，可以与CPU共享内存。

         计算功能集的版本为1.3或者更高的显卡才能支持双精度浮点数的计算。

         尖括号的第一个参数表示设备在执行核函数时使用的并行线程块的数量，

         并行线程块集合也称为一个线程格(Grid)。线程格既可以是一维的线程块集合，也可以是二维的线程块集合。

         GPU有着完善的内存管理机制，它将强行结束所有违反内存访问规则的进程。

         在启动线程块数组时，数组每一维的最大数量都不能超过65535.这是一种硬件限制,如果启动的线程块数量超过了这个限值,那么程序将运行失败。   

         CUDA运行时将线程块(Block)分解为多个线程。当需要启动多个并行线程块时，只需将尖括号中的第一个参数由1改为想要启动的线程块数量。在尖括号中，第二个参数表示CUDA运行时在每个线程块中创建的线程数量。

         硬件将线程块的数量限制为不超过65535.同样，对于启动核函数时每个线程块中的线程数量，硬件也进行了限制。具体来说，最大的线程数量不能超过设备属性结构中maxThreadsPerBlock域的值。这个值并不固定，有的是512，有的是1024.

         内置变量blockDim，对于所有线程块来说，这个变量是一个常数，保存的是线程块中每一维的线程数量。

         内置变量gridDim，对于所有线程块来说，这个变量是一个常数，用来保存线程格每一维的大小，即每个线程格中线程块的数量。

         内置变量blockIdx，变量中包含的值就是当前执行设备代码的线程块的索引。

         内置变量threadIdx，变量中包含的值就是当前执行设备代码的线程索引。

         CUDA运行时允许启动一个二维线程格，并且线程格中的每个线程块都是一个三维的线程数组。

         CUDA C支持共享内存：可以将CUDA  C的关键字__share__添加到变量声明中，这将使这个变量驻留在共享内存中。CUDA C编译器对共享内存中的变量与普通变量将分别采取不同的处理方式。

         CUDA架构将确保，除非线程块中的每个线程都执行了__syncthreads()，否则没有任何线程能执行__syncthreads()之后的指令。

         由于在GPU上包含有数百个数学计算单元，因此性能瓶颈通常并不在于芯片的数学计算吞吐量，而是在于芯片的内存带宽。

         常量内存用于保存在核函数执行期间不会发生变化的数据。NVIDIA硬件提供了64KB的常量内存，并且对常量内存采取了不同于标准全局内存的处理方式。在某些情况下，用常量内存来替换全局内存能有效地减少内存带宽。要使用常量内存，需在变量前面加上__constant__关键字。

         在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。在程序中的每一行，线程束中的每个线程都将在不同的数据上执行相同的指令。

         纹理内存是在CUDA C程序中可以使用的另一种只读内存。与常量内存类似的是，纹理内存同样缓存在芯片上，因此在某些情况中，它能够减少对内存的请求并提供更高效的内存带宽。纹理缓存是专门为那些在内存访问模式中存在大量空间局部性(Spatial Locality)的图形应用程序而设计的。

         NVIDIA将GPU支持的各种功能统称为计算功能集(Compute Capability)。高版本计算功能集是低版本计算功能集的超集。

         只有1.1或者更高版本的GPU计算功能集才能支持全局内存上的原子操作。此外，只有1.2或者更高版本的GPU计算功能集才能支持共享内存上的原子操作。CUDA C支持多种原子操作。

         C库函数malloc函数将分配标准的，可分页的(Pageble)主机内存。而cudaHostAlloc函数将分配页锁定的主机内存。页锁定内存也称为固定内存(Pinned Memory)或者不可分页内存，它有一个重要的属性：操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存上。因此，操作系统能够安全地使某个应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位。

         固定内存是一把双刃剑。当使用固定内存时，你将失去虚拟内存的所有功能。特别是，在应用程序中使用每个页锁定内存时都需要分配物理内存，因为这些内存不能交换到磁盘上。这意味着，与使用标准的malloc函数调用相比，系统将更快地耗尽内存。因此，应用程序在物理内存较少的机器上会运行失败，而且意味着应用程序将影响在系统上运行的其它应用程序的性能。

         建议，仅对cudaMemcpy()调用中的源内存或者目标内存，才使用页锁定内存，并且在不再需要使用它们时立即释放，而不是等到应用程序关闭时才释放。

         CUDA流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。

         通过使用零拷贝内存，可以避免CPU和GPU之间的显式复制操作。

         对于零拷贝内存，独立GPU和集成GPU，带来的性能提升是不同的。对于集成GPU，使用零拷贝内存通常都会带来性能提升，因为内存在物理上与主机是共享的。将缓冲区声明为零拷贝内存的唯一作用就是避免不必要的数据复制。但是，所有类型的固定内存都存在一定的局限性，零拷贝内存同样不例外。每个固定内存都会占用系统的可用物理内存，这最终将降低系统的性能。对于独立GPU，当输入内存和输出内存都只能使用一次时，那么在独立GPU上使用零拷贝内存将带来性能提升。但由于GPU不会缓存零拷贝内存的内容，如果多次读取内存，那么最终将得不偿失，还不如一开始就将数据复制到GPU。

         CUDA工具箱(CUDAToolkit)包含了两个重要的工具库：(1)、CUFFT(Fast FourierTransform，快速傅里叶变换)库；(2)、CUBLAS(Basic Linear Algebra Subprograms,BLAS)是一个线性代数函数库。

         NPP(NVIDIA Performance Primitives)称为NVIDIA性能原语，它是一个函数库，用来执行基于CUDA加速的数据处理操作，它的基本功能集合主要侧重于图像处理和视频处理。

新建一个基于CUDA的测试工程testCUDA，此工程中除了包括common文件外，还添加了另外三个文件，分别为testCUDA.cu、funset.cu、funset.cuh，这三个文件包括了书中绝大部分的测试代码，见原文；