1. CUDA 基础
1.1. CUDA 简介
GPU 加速计算正在逐步取代 CPU 计算，近年来加速计算带来了越来越多的突破性进展，各类应用程序对加速计算日益增长地需求、便捷地编写加速计算的程序的需求以及不断改进的支持加速计算的硬件设施，所有这一切都在推动着计算方式从 CPU 计算过渡到 GPU 加速计算。
无论是从出色的性能还是易用性来看，CUDA 计算平台均是加速计算的重要实现方式。CUDA 提供了一种可扩展于 C、C++、Python 和 Fortran 等语言的编码接口，并行化后的代码能够在 NVIDIA GPU 上运行，以大幅加速应用程序。它包含有 DNN、BLAS、图形分析 和 FFT 等等库，并且还附带功能强大的命令行和可视化分析器。
CUDA 支持许多领域的超性能计算应用程序：计算流体动力学、分子动力学、量子化学、物理学 和高性能计算 (HPC)等等。
学习 CUDA 将能帮你加速自己的应用程序。应用程序加速后的执行速度会远远超过原本在 CPU 上的执行速度，使那些在 CPU 上性能受限的计算得以进行下去。在本教程中, 你将学习使用 CUDA 的 C/C++ 接口作为加速应用程序编程的入门知识，这些入门知识足以让你加速自己的 CPU 应用程序，以获得性能上的巨大提升并帮你迈入全新的计算领域。
1.2. 学习前的准备工作
如要充分利用本教程学习CUDA，那么你应该要先有如下知识储备：
在 C++/C 中声明变量、编写循环并使用 if/else 语句。
在 C++/C 中定义和调用函数。
在 C++/C 中分配数组。
说白了就是要有C或C++语言的基础，此外不需要事先知道任何关于 CUDA 的知识，当你在本教程完成学习后，你就可以做到：
编写、编译及运行既可调用 CPU 函数也可启动 GPU 核函数的 C/C++ 程序。
通过配置参数控制并行线程的层次结构。
重构串行循环以在 GPU 上并行执行其迭代。
分配和释放可用于 CPU 和 GPU 的内存。
处理 CUDA 代码产生的错误。
加速 CPU 应用程序。
1.3. 加速系统的硬件设施
带有GPU的计算机系统称为加速系统（又称异构系统，即指包含CPU和GPU的系统）。在一个包含 NVIDIA GPU 的加速系统的实验环境上，可以使用 nvidia-smi 命令查询有关此 GPU 的信息。例如：
nvidia-smi
按回车之后，将输出该机器上的GPU信息
需要注意的是，加速系统在运行程序时首先会运行 CPU 程序，在运行到需要GPU进行大规模并行计算的函数时，再将对应函数载入GPU执行。
也就是说，由GPU加速的依然还是纯CPU的应用程序，只是某些模块在运行时调入了GPU中，该模块在同步完毕后将会重新回到CPU中执行主程序的后续代码：
2. 编写在GPU运行的代码
CUDA 为许多编程语言提供了扩展接口，而在本教程用CUDA为 C/C++ 提供的接口来展示。对编程语言的扩展可以让开发人员在 GPU 上更加方便的运行 CUDA 库的函数。
以下是一个 .cu 文件（.cu 是 CUDA 加速程序的文件扩展名，实际上.cu文件只是含有CUDA代码的.cpp文件，没有别的特殊之处）。其中包含两个函数，第一个函数 CPUFunction() 将在 CPU 上运行，第二个函数 GPUFunction() 将在 GPU 上运行：
// 在CPU上运行的函数 void CPUFunction() { printf("This function is defined to run on the CPU.\n"); } // 在GPU上运行的函数 __global__ void GPUFunction() { printf("This function is defined to run on the GPU.\n"); } int main() { CPUFunction(); // 调用CPU函数 GPUFunction<<<1, 1>>>(); // 调用GPU函数 cudaDeviceSynchronize(); // 同步 }
根据上面的代码，我们来讲解一些需要特别注意的重要代码行，以及加速计算中使用的一些其他常用术语：
__global__ void GPUFunction()：
__global__ 关键字表明该函数将在 GPU 上运行并可全局调用（ 既可以由CPU ，也可以由 GPU 调用）；
通常，我们将在 CPU 上执行的代码称为 Host （主机）代码，而将在 GPU 上运行的代码称为 Device （设备）代码；
注意返回类型为 void。使用 __global__ 关键字定义的函数返回值需为 void 类型。
GPUFunction<<<1, 1>>>()：
通常，我们把要运行在 GPU 上的函数称为 kernel （核）函数;
启动核(kernel)函数时，我们必须事先配置GPU参数，使用 <<< ... >>> 语法向核函数传递两个必要的参数;
在 <<< ... >>> 中传递的参数用于为核函数设定线程的层次结构，第一个参数定义线程块(Block)的数量，第二个参数定义Block中含有的线程(Thread)数量。例如本例中的核函数 GPUFunction() 将在包含 1 个线程（第二个配置参数）的 1 个线程块（第一个执行配置参数）上运行。
cudaDeviceSynchronize()：
与其他并行化的代码类似，核函数启动方式为异步，即 CPU 代码将继续执行而不会等待核函数执行完成；
调用 CUDA 提供的函数 cudaDeviceSynchronize 可以让Host 代码(CPU) 等待 Device 代码(GPU) 执行完毕，再在CPU上继续执行。
2.1. 编写运行一个 Hello GPU 核函数
#include <stdio.h> void helloCPU() { printf("Hello from the CPU.\n"); } // __global__ 表明这是一个全局GPU核函数. __global__ void helloGPU() { printf("Hello from the GPU.\n"); } int main() { helloCPU(); // 调用CPU函数 /* 使用 <<<...>>> 配置核函数的GPU参数， * 第一个1表示1个线程块，第二个1表示每个线程块1个线程。*/ helloGPU<<<1, 1>>>(); // 调用GPU函数 cudaDeviceSynchronize(); // `cudaDeviceSynchronize` 同步CPU和GPU }
现在来编译并运行加速后的CUDA代码。将上述文件命名为hello-gpu.cu，执行命令：
nvcc hello-gpu.cu -o hello-gpu ./hello-gpu
得到结果：
3. CUDA 线程的层次结构
从上面的图中可以看出，CUDA线程的层次结构分为三层：Thread（线程）、Block（块）、Grid（网格），网格由块组成，块由线程组成。
3.1. 运行核函数
我们可以通过配置参数指定核函数如何在 GPU 的多个线程中并行运行。具体来说，就可以配置 Block 的数量以及每个 Block 中所包含 Thread 的数量。配置参数的语法如下：
<<< Block 数, 每个Block中的 Thread 数>>>
启动核函数时，核函数代码由我们自行配置的 Block 中的每个 Thread 执行。因此，如果假设已定义一个名为 someKernel 的核函数，则GPU线程可以配置为下列情况：
someKernel<<<1, 1>>() 在GPU中为该核函数分配1个具有1个线程的线程块，核函数中的代码将只运行1次；
someKernel<<<1, 10>>() 在GPU中为该核函数分配1个具有10个线程的线程块，核函数中的代码将运行10次；
someKernel<<<10, 1>>() 在GPU中为该核函数分配10个具有1个线程的线程块，核函数中的代码将运行10次；
someKernel<<<10, 10>>() 在GPU中为该核函数分配10个具有10个线程的线程块，核函数中的代码将运行100次；
启动并行运行的核函数示例：
#include <stdio.h> __global__ void firstParallel() { printf("This is running in parallel.\n"); } int main() { firstParallel<<<5, 5>>>(); // 在GPU中为核函数分配5个具有5个线程的线程块，将运行25次； cudaDeviceSynchronize(); // 同步 }
将上述代码命名为basic-parallel.cu，然后编译运行：
nvcc basic-parallel.cu -o basic-parallel ./basic-parallel
结果如下，数了一下，确实是25次：
3.2. 线程和块的索引
如图所示，每个线程在其线程块的内部都会被分配一个索引，从 0 开始。此外，每个线程块也会被分配一个索引，也是从 0 开始。正如线程组成线程块，线程块又会组成网格(Grid)，而网格是 CUDA 线程层次结构中级别最高的实体，它没有索引。
简言之，CUDA 核函数在由一个或多个线程块组成的网格中执行，且每个线程块中均包含相同数量的一个或多个线程（每个线程块中的线程数量相同）。
在核函数中，可以通过两个变量来获取到索引： threadIdx.x （线程索引）和 blockIdx.x（线程块索引）。
现在让我们来使用索引控制特定的线程和块：
#include <stdio.h> // 核函数 __global__ void printSuccessForCorrectExecutionConfiguration() { // 当执行到第255个线程块的第1023个线程时，才输出 if(threadIdx.x == 1023 && blockIdx.x == 255) { printf("Success!\n"); // 输出 Success！ printf("threadIdx.x: %d\n", threadIdx.x); // 输出线程ID printf("blockIdx.x: %d\n", blockIdx.x); // 输出线程块ID } } int main() { // 配置该核函数由256个含有1024个线程的线程块中执行 printSuccessForCorrectExecutionConfiguration<<<256, 1024>>>(); cudaDeviceSynchronize(); // 同步 }
将上述代码命名为thread-and-block-idx.cu，然后编译运行：
nvcc thread-and-block-idx.cu -o thread-and-block-idx ./thread-and-block-idx
输出：
3.3. 用 CUDA 加速 For 循环
到此为止，加速 for 循环就是一个可行的操作了。在加速计算中，for 循环不再顺序执行每次迭代，而是让每次迭代都在不同的线程中并行执行。
例如，现在有以下在 CPU 中执行的 for 循环：
int N = 10; for (int i = 0; i < N; ++i) { printf("%d\n", i); }
如要并行此循环，必须执行以下 2 个步骤：
编写用于执行单次迭代工作的核函数。
调用核函数时为它配置执行参数，即并行的线程数，每个线程执行一次迭代。
如下例程序：
#include <stdio.h> // 核函数 __global__ void loop() { // 输出每一个线程的线程号（0~9） printf("This is iteration number %d\n", threadIdx.x); } int main() { loop<<<1, 10>>>(); // 执行核函数 cudaDeviceSynchronize(); }
将上述代码命名为single-block-loop.cu，然后编译运行：
nvcc single-block-loop.cu -o single-block-loop ./single-block-loop
输出：
3.4. 管理不同块之间的线程
之前提到过，一个线程块可以包含多个线程，那么我们就可以调整线程块的大小以实现更多类型的并行化。线程块包含的线程具有数量限制：确切地说是 1024 个（即每个块中的线程数量 <= 1024）。通常为了增加加速应用程序中的并行量，我们需要利用多个线程块，并在它们之间进行协调。
CUDA 核函数中，记录了每个块中线程数的变量是 blockDim.x（一个线程块中包含的线程数量，每个块中包含的线程数都是一样的）。通过将此变量与 blockIdx.x 和 threadIdx.x 变量结合使用，并借助表达式 threadIdx.x + blockIdx.x * blockDim.x 计算线程ID。该表达式可以用C++中访问二维数组的索引计算来类比看待，以增强理解。
以下是详细示例：
配置参数 <<<10, 10>>> 将启动共计拥有 100 个线程的网格，该网格又分为由 10 个线程组成的 10 个线程块（即一个线程块中含有10个线程，blockDim.x=10）。这时候，就可以利用表达式 threadIdx.x + blockIdx.x * blockDim.x 来计算某个线程的唯一索引（0 至 99 之间）了。
如果线程块 blockIdx.x 索引为 0，则 blockIdx.x * blockDim.x 为 0。以 0 为起始索引加上可能的 threadIdx.x 值（0 至 9），便可在网格中找到索引为 0 至 9 的线程。
如果线程块 blockIdx.x 索引为 1，则 blockIdx.x * blockDim.x 为 10。以 10 为起始索引加上可能的 threadIdx.x 值（0 至 9），便可在网格中找到索引为 10 至 19 的线程。
如果线程块 blockIdx.x 索引为 5，则 blockIdx.x * blockDim.x 为 50。以 50 为起始索引加上可能的 threadIdx.x 值（0 至 9），便可在网格中找到索引为 50 至 59 的线程。
如果线程块 blockIdx.x 索引为 9，则 blockIdx.x * blockDim.x 为 90。以 90 为起始索引加上可能的 threadIdx.x 值（0 至 9），便可在网格中找到索引为 90 至 99 的线程。
现在我们来加速具有多个线程块的For循环：
#include <stdio.h> __global__ void loop() { // 在Grid中遍历所有thread int i = blockIdx.x * blockDim.x + threadIdx.x; printf("%d\n", i); } int main() { /* * 配置参数还可以试试其他的，例如： * <<<5, 2>>> * <<<10, 1>>> */ loop<<<2, 5>>>(); cudaDeviceSynchronize(); }
将上述代码命名为multi-block-loop.cu，然后编译运行：
nvcc multi-block-loop.cu -o multi-block-loop ./multi-block-loop
输出：
