CUDA

计算统一设备架构（Computer Unified Device Architecture, CUDA）是英伟达开发的并行计算平台和编程模型，只支持NVIDIA GPU卡。OpenCL可以为其他的GPU（AMD和Inter的设备）编写并行代码，其比CUDA更为复杂。

CUDA可以使用简单的编程API在GPU上创建大规模的并行应用程序。

何为并行？

随着手持设备的功能要求越来越高，我们需要一个可以快速运行的处理器以较高的时钟速度、较小的体积和更小的功率执行各项任务。

随着越来越多的晶体管封装到芯片中，会导致时钟速度提高；但是最近几年，时钟速度基本保持不变，原因是：高功率损耗和高时钟速率。晶体管在小面积内封装的芯片中高速工作，其能耗很高，很难保持处理器的低温。

举个例子，需要你在很短的时间内挖一个大洞，会有几种方法：

  1. 挖的更快 → 更快的时钟速度
  2. 买一把好铲子 → 拥有更多可以在每个时钟周期做更多工作的晶体管。
  3. 雇佣很多台挖掘机，帮助你完成工作 → 并行处理

但是，由于功耗的限制，方法一和二受到限制。方法三就类似拥有很多可以并行执行任务的更小更简单的处理器。

所以GPU不是一个可以执行复杂任务的强大处理器，而是由许多小而简单的且可以并行工作的处理器。

所以提高性能的三种方法：

  1. 更快的时钟速度；
  2. 单处理器每时钟周期做更多的工作；
  3. 许多小型处理器并行工作，GPU使用此模式来提高性能。

    GPU vs. CPU

    CPU具有复杂的控制硬件和较少的数据计算硬件。
  • 负责的控制硬件在性能上提供了CPU的灵活性和简单的编程接口，但是其具有很高的功耗。

GPU具有简单的控制硬件和更多的数据计算硬件。

  • 其具有并行计算的能力，这种结构更加节能，但是不够灵活，有更为严格的编程模型。

任何硬件架构的性能都是根据延迟和吞吐量来度量。

  • 延迟：完成给定任务所花费的时间
  • 吞吐量：在给定时间内完成任务的数量

两个指标并不相互矛盾，改进延迟将会提高吞吐量；

比如顾客在商店排队买东西，其目的是：很短的时间内买完所需东西，这就是改进延迟；而商店老板目的：一天内看到的顾客越来越多，这就是提高吞吐量。

正常的串行CPU被设计为优化延迟，在最短时间内执行所有指令；GPU被设计为优化吞吐量，在给定时间内执行更多指令。而对于计算机视觉任务来说，我们不介意单个像素处理的延迟，而是在给定时间内处理更多的像素。

综上所述，并行计算可以满足我们想要在相同的时钟速度和功率下，提高计算性能。而GPU通过让许多简单的计算单元达到并行工作的目的。

CUDA架构

GPU由许多并行处理单元 - Core 核心。在硬件层面，这些 Core 被分为流处理器和流多处理器。GPU有这些流多处理器的网格 Grid。在软件层面，CUDA程序是作为一系列并行运行的多线程Thread 来执行的，每个线程都在不同的核心上运行。

将GPU看作多个块block的组合，每个块可以执行多个线程。每个块绑定到GPU上的不同流多处理器。来自同一block的线程之间可以相互通信。gpu有一个分层的内存结构，处理多个block内线程之间的通信。

CUDA编程模型

CPU及其内存称为主机Host，GPU及其显存称为设备Device。CUDA代码包含主机和设备的代码。主机代码由普通的C/C++编译器在CPU上编译，设备代码由GPU编译器在GPU上编译。主机代码通过内核调用设备代码，并在设备上并行启动多个线程，线程启动多少由程序员决定。

设备代码和普通的C代码类似，只是设备代码在大量内核上并行执行。要使设备代码工作，需要设备显存上的数据。因此，在启动线程之前，主机将数据从主机内存复制到设备显存，线程处理来自设备显存的数据，并将结果存储到设备显存中；最后将这些数据复制回主机内存等待进一步处理。

综上所述，CUDA C开发步骤：

  1. 为主机和设备显存中的数据分配内存；
  2. 将数据从主机内存复制到设备显存；
  3. 通过执行并行度（指定线程数量）来启动内核；
  4. 所有线程执行完，将数据从设备显存复制回主机内存；
  5. 释放主机和设备上使用的所有内存。

CUDA C程序

主机代码由标准C编译器GCC编译的，设备代码由NVIDIA GPU编译器执行。所有的CUDA代码必须保存为.cu 后缀名。

  1. #include <iostream>
  3. __global__ void myfirstkernel(void)
  4. {
  5. }
  7. int main(void)
  8. {
  9.     myfirstkernel <<<1, 1>>>();
  10.     printf("Hello, CUDA!\n");
  11.     return 0;
  12. }

CUDA代码与标准C代码，两个主要区别：

  1. 一个名为 myfirstkernel 的空函数，前缀为 __global__
  2. 使用 <<<1, 1>>> 调用 myfirstkernel 函数

__global__是 CUDA C 在标准C中添加一个限定符，它告诉编译器在这个限定符后面的函数定义应该在设备上而不是在主机上运行。所以 myfirstkernel 将在设备上而不是和主机上。而NVCC编译器会将 main() 函数提供给C编译器，所以其会在主机上运行。

main 函数中在调用myfirstkernel 函数之后带有一些尖括号和数值，被称为内核调用，这代表：从主机代码调用设备代码，尖括号中的两个数字代表块的数量和将在设备上并行运行的线程数。所以myfirstkernel函数只会运行在一个只有一个块和一个线程的设备上。

CUDA编译器不能在没有设备代码的情况下编译代码。 ×