# CUDA入门第一课
## 一、CUDA是什么
CUDA（Compute Unified Device Architecture，统一计算设备架构）是由英伟达（NVIDIA）推出的一种并行计算平台和编程模型。它允许开发者利用NVIDIA GPU（图形处理单元）的强大计算能力来加速计算密集型任务。
### （一）GPU与CPU的区别
- **CPU（中央处理器）**：
- 通常拥有少量的核心（如4核、8核等），每个核心性能强大，擅长处理复杂的逻辑控制和顺序执行的任务。例如，在运行操作系统、执行复杂的算法逻辑（如深度优先搜索等）时，CPU能够很好地发挥作用。它就像是一个“精巧的工匠”，能够精确地处理各种复杂的任务细节。
- CPU的设计注重单线程性能，指令执行是顺序的，虽然现代CPU有超线程技术等，但总体上还是以处理少量线程的复杂任务为主。
- **GPU**：
- 拥有大量的核心（可以达到数千个核心），每个核心相对简单，但能够同时处理大量的简单任务。这就好比一个“庞大的工人团队”，虽然单个工人的能力有限，但当有大量重复性工作（如在工厂中同时组装多个相同的零件）时，这个团队能够发挥巨大的优势。
- GPU最初是用于图形渲染，如在游戏中快速计算像素颜色、纹理映射等。后来人们发现，它在处理并行计算任务（如矩阵运算、图像处理等）时也表现出色，因为这些任务可以被分解成大量独立的小任务，每个核心可以同时处理一个小任务。
### （二）CUDA的作用
- CUDA提供了一套编程接口和工具，让开发者能够编写程序，将计算任务分配给GPU执行。通过CUDA，开发者可以将原本在CPU上执行的计算密集型代码（如科学计算、机器学习中的矩阵运算等）移植到GPU上，从而实现加速。
- 例如，在深度学习中，神经网络的训练需要进行大量的矩阵乘法运算。使用CUDA编程，可以将这些矩阵乘法运算任务分配给GPU的多个核心同时执行，相比于在CPU上单线程执行，能够大大缩短训练时间。
## 二、CUDA的编程模型
### （一）线程层次结构
- **线程（Thread）**：
- 这是最基本的执行单元，类似于CPU中的线程。在GPU中，一个线程执行一段代码（称为核函数，后面会详细介绍），处理数据的一个小部分。例如，在处理一个图像的像素值计算任务时，每个线程可以负责计算一个像素点的新值。
- **线程块（Thread Block）**：
- 一组线程组成一个线程块。线程块内的线程可以方便地进行协作，例如通过共享内存（Shared Memory）进行数据交换。线程块内的线程执行是相对独立的，但它们可以被分配到同一个GPU核心上执行，这样可以更好地利用GPU核心的资源。
- 例如，一个线程块可以包含256个线程，这些线程共同处理图像的一个小区域，通过共享内存快速传递中间计算结果，提高计算效率。
- **网格（Grid）**：
- 一个或多个线程块组成一个网格。网格是GPU上执行任务的最高层次结构。一个核函数的执行对应一个网格，网格中的所有线程块共同完成整个计算任务。
- 例如，对于一个大型矩阵运算任务，可以将矩阵划分为多个小块，每个小块对应一个线程块，所有线程块组成一个网格，整个网格执行完成就代表矩阵运算任务完成。
### （二）内存层次结构
- **寄存器（Register）**：
- 每个线程有自己的寄存器，用于存储线程私有的变量。寄存器访问速度非常快，但数量有限。例如，线程中用于临时存储计算中间结果的变量通常存储在寄存器中。
- **共享内存（Shared Memory）**：
- 线程块内的线程共享一块共享内存。它可以用于线程之间的数据交换和协作。共享内存的访问速度比全局内存快，但容量相对较小。例如，在进行矩阵分块计算时，线程块内的线程可以将矩阵块的数据加载到共享内存中，然后进行高效的计算和数据交互。
- **全局内存（Global Memory）**：
- 所有线程都可以访问全局内存，它是GPU内存的主要部分。全局内存容量较大，但访问速度相对较慢。通常，程序开始时会将数据从CPU内存（主机内存）拷贝到GPU的全局内存中，计算完成后，再将结果从全局内存拷贝回CPU内存。
- **常量内存（Constant Memory）和纹理内存（Texture Memory）**：
- 常量内存用于存储只读的常量数据，访问速度较快，适用于存储一些不会改变的参数，如物理常数等。纹理内存主要用于图形纹理映射等操作，它具有特殊的缓存机制，可以提高对图像纹理数据的访问效率。
## 三、CUDA编程的基本步骤
### （一）安装CUDA工具包
- 首先需要安装NVIDIA的CUDA工具包，它包括编译器（nvcc）、运行时库、开发文档等。安装完成后，就可以开始编写CUDA程序了。
### （二）编写CUDA程序
- **核函数（Kernel Function）**：
- 核函数是运行在GPU上的函数，它是由CPU（主机）调用的。核函数的定义方式与C语言函数类似，但有一些特殊的修饰符和规则。例如：
```c
__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements)
{
int i = threadIdx.x;
if (i < numElements)
{
C[i] = A[i] + B[i];
}
}
```
这是一个简单的核函数，用于计算两个向量的和。`__global__`修饰符表示这是一个核函数。`threadIdx.x`是线程的索引，用于确定每个线程处理数据的位置。
- **主机代码（Host Code）**：
- 主机代码运行在CPU上，负责管理GPU的执行过程，如分配内存、拷贝数据、调用核函数等。例如：
```c
int main()
{
int numElements = 50000;
size_t size = numElements * sizeof(float);
float *h_A = (float *)malloc(size);
float *h_B = (float *)malloc(size);
float *h_C = (float *)malloc(size);
// 初始化数据
for (int i = 0; i < numElements; ++i)
{
h_A[i] = rand()/(float)RAND_MAX;
h_B[i] = rand()/(float)RAND_MAX;
}
float *d_A, *d_B, *d_C;
cudaMalloc((void **)&d_A, size);
cudaMalloc((void **)&d_B, size);
cudaMalloc((void **)&d_C, size);
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
int threadsPerBlock = 256;
int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
free(h_A);
free(h_B);
free(h_C);
return 0;
}
```
在主机代码中，首先分配CPU内存和GPU内存，然后将数据从CPU内存拷贝到GPU内存。接着，通过`<<<blocksPerGrid, threadsPerBlock>>>`语法调用核函数，其中`blocksPerGrid`和`threadsPerBlock`分别表示网格中线程块的数量和每个线程块中线程的数量。最后，将计算结果从GPU内存拷贝回CPU内存，并释放内存资源。
### （三）编译和运行
- 使用CUDA编译器`nvcc`对CUDA程序进行编译。例如，对于上面的程序，可以使用命令`nvcc -o vectorAdd vectorAdd.cu`进行编译，其中`vectorAdd.cu`是源文件名，`vectorAdd`是生成的可执行文件名。编译完成后，就可以在支持CUDA的设备上运行程序了。
## 四、CUDA的优势和应用场景
### （一）优势
- **强大的计算能力**：GPU拥有大量的核心，能够同时处理大量的计算任务，对于计算密集型任务（如科学计算、图像处理、机器学习等）具有显著的加速效果。例如，在深度学习训练中，使用CUDA编程可以在GPU上实现高效的矩阵运算，将训练时间从几天缩短到几小时甚至更短。
- **易于编程**：CUDA提供了一套相对完善的编程接口和工具，使得开发者能够比较容易地将计算任务移植到GPU上。虽然CUDA编程需要一定的学习成本，但相比于其他一些并行计算框架，它的
