CUDA架构与应用杂谈

RecursiveFrog · 发表于 2022-6-12 15:02

CUDA架构与应用杂谈
CUDA（Compute Unified Device Architecture），是显卡厂商NVIDIA推出的运算平台。 CUDA是一种由NVIDIA推出的通用并行计算架构，该架构使GPU能够解决复杂的计算问题。它包含了CUDA指令集架构（ISA）以及GPU内部的并行计算引擎。开发人员可以使用C语言来为CUDA架构编写程序，所编写出的程序可以在支持CUDA的处理器上以超高性能运行。CUDA3.0已经开始支持C++和FORTRAN。
CUDA 是 NVIDIA 发明的一种并行计算平台和编程模型。它通过利用图形处理器 (GPU) 的处理能力，可大幅提升计算性能。
目前为止基于 CUDA 的 GPU 销量已达数以百万计，软件开发商、科学家以及研究人员正在各个领域中运用 CUDA，其中包括图像与视频处理、计算生物学和化学、流体力学模拟、CT 图像再现、地震分析以及光线追踪等等。
计算行业正在从只使用CPU的“中央处理”向CPU与GPU并用的“协同处理”发展。为打造这一全新的计算典范，NVIDIA（英伟达）发明了CUDA（Compute Unified Device Architecture，统一计算设备架构）这一编程模型，是想在应用程序中充分利用CPU和GPU各自的优点。该架构已应用于GeForce（精视）、ION（翼扬）、Quadro以及Tesla GPU（图形处理器）上，对应用程序开发人员来说，这是一个巨大的市场。
参考文献链接
https://baike.baidu.com/item/CUDA/1186262?fr=aladdin
https://mp.weixin.qq.com/s/kxYSw_fR4QMZ2-O5fvOR8g
https://www.zhihu.com/question/461354739/answer/1964488472
https://blog.csdn.net/liusaoyue/article/details/5100656
1.深蓝学院课程讲解：https://www.shenlanxueyuan.com/course/410
2. D. Kirk and W. Hwu, “Programming Massively Parallel Processors –A Hands-on Approach, Second Edition”
3. CUDA by example, Sanders and Kandrot
4. Nvidia CUDA C Programming Guide：https://docs.nvidia.com/cuda/cuda-c-programming-guide/
5. CS/EE217 GPU Architecture andProgramming

GPU架构
在消费级市场上，几乎每一款重要的消费级视频应用程序都已经使用CUDA加速或很快将会利用CUDA来加速，其中不乏Elemental Technologies公司、MotionDSP公司以及LoiLo公司的产品。
在科研界，CUDA一直受到热捧。例如，CUDA现已能够对AMBER进行加速。AMBER是一款分子动力学模拟程序，全世界在学术界与制药企业中有超过0,000名研究人员使用该程序来加速新药的探索工作。
在金融市场，Numerix以及CompatibL针对一款全新的对手风险应用程序发布了CUDA支持并取得了18倍速度提升。Numerix为近400家金融机构所广泛使用。
CUDA的广泛应用造就了GPU计算专用Tesla GPU的崛起。全球财富五百强企业已经安装了700多个GPU集群，这些企业涉及各个领域，例如能源领域的斯伦贝谢与雪佛龙以及银行业的法国巴黎银行。
随着微软Windows 7与苹果Snow Leopard操作系统的问世，GPU计算必将成为主流。在这些全新的操作系统中，GPU将不仅仅是图形处理器，它还将成为所有应用程序均可使用的通用并行处理器。
CUDA C 编程及 GPU 基本知识

1 CPU 和 GPU 的基础知识

提到处理器结构，有2个指标是经常要考虑的：延迟和吞吐量。所谓延迟，是指从发出指令到最终返回结果中间经历的时间间隔。而所谓吞吐量，就是单位之间内处理的指令的条数。
下图1是 CPU 的示意图。从图中可以看出 CPU 的几个特点：
1. CPU 中包含了多级高速的缓存结构。 因为我们知道处理运算的速度远高于访问存储的速度，那么奔着空间换时间的思想，设计了多级高速的缓存结构，将经常访问的内容放到低级缓存中，将不经常访问的内容放到高级缓存中，从而提升了指令访问存储的速度。
2. CPU 中包含了很多控制单元。 具体有2种，一个是分支预测机制，另一个是流水线前传机制。
3. CPU 的运算单元 (Core) 强大，整型浮点型复杂运算速度快。

图1：CPU 的示意图
所以综合以上三点，CPU 在设计时的导向就是减少指令的时延，我们称之为延迟导向设计，如下图3所示。
下图2是 GPU 的示意图，它与之前 CPU 的示意图相比有着非常大的不同。从图中可以看出 GPU 的几个特点 (注意紫色和黄色的区域分别是缓存单元和控制单元)：
1. GPU 中虽有缓存结构但是数量少。 因为要减少指令访问缓存的次数。
2. GPU 中控制单元非常简单。 控制单元中也没有分支预测机制和数据转发机制。对于复杂的指令运算就会比较慢。
3. GPU 的运算单元 (Core) 非常多，采用长延时流水线以实现高吞吐量。 每一行的运算单元的控制器只有一个，意味着每一行的运算单元使用的指令是相同的，不同的是它们的数据内容。那么这种整齐划一的运算方式使得 GPU 对于那些控制简单但运算高效的指令的效率显著增加。

图2：GPU 的示意图
所以，GPU 在设计过程中以一个原则为核心：增加简单指令的吞吐。因此，我们称 GPU 为吞吐导向设计，，如下图3所示。

图3：CPU 是延迟导向设计，GPU 是吞吐导向设计
那么究竟在什么情况下使用 CPU，什么情况下使用 GPU 呢？
CPU 在连续计算部分，延迟优先，CPU 比 GPU ，单条复杂指令延迟快10倍以上。
GPU 在并行计算部分，吞吐优先，GPU 比 CPU ，单位时间内执行指令数量10倍以上。
适合 GPU 的问题：
1. 计算密集：数值计算的比例要远大于内存操作，因此内存访问的延时可以被计算掩盖。
2. 数据并行：大任务可以拆解为执行相同指令的小任务，因此对复杂流程控制的需求较低。
2 CUDA 编程的重要概念
CUDA (Compute Unified Device Architecture)，由英伟达公司2007年开始推出，初衷是为 GPU 增加一个易用的编程接口，让开发者无需学习复杂的着色语言或者图形处理原语。
OpenCL (Open Computing Languge) 是2008年发布的异构平台并行编程的开放标准，也是一个编程框架。OpenCL 相比 CUDA，支持的平台更多，除了 GPU 还支持 CPU、DSP、FPGA 等设备。
下面我们将以 CUDA 为例，介绍 GPU 编程的基本思想和基本操作。
首先主机端 (host) 和设备端 (device)，主机端一般指我们的 CPU，设备端一般指我们的 GPU。
一个 CUDA 程序，我们可以把它分成3个部分：
第1部分是： 从主机 (host) 端申请 device memory，把要拷贝的内容从 host memory 拷贝到申请的 device memory 里面。
第2部分是： 设备端的核函数对拷贝进来的东西进行计算，来得到和实现运算的结果，图4中的 Kernel 就是指在 GPU 上运行的函数。
第3部分是： 把结果从 device memory 拷贝到申请的 host memory 里面，并且释放设备端的显存和内存。

图4：一个 CUDA 程序可以分成3个部分
CUDA 编程中的内存模型
这里就引出了一个非常重要的概念就是 CUDA 编程中的内存模型。
从硬件的角度来讲：
CUDA 内存模型的最基本的单位就是 SP (线程处理器)。每个线程处理器 (SP) 都用自己的 registers (寄存器) 和 local memory (局部内存)。寄存器和局部内存只能被自己访问，不同的线程处理器之间呢是彼此独立的。
由多个线程处理器 (SP) 和一块共享内存所构成的就是 SM (多核处理器) (灰色部分)。多核处理器里边的多个线程处理器是互相并行的，是不互相影响的。每个多核处理器 (SM) 内都有自己的 shared memory (共享内存)，shared memory 可以被线程块内所有线程访问。
再往上，由这个 SM (多核处理器) 和一块全局内存，就构成了 GPU。一个 GPU 的所有 SM 共有一块 global memory (全局内存)，不同线程块的线程都可使用。
上面这段话可以表述为：每个 thread 都有自己的一份 register 和 local memory 的空间。同一个 block 中的每个 thread 则有共享的一份 share memory。此外，所有的 thread (包括不同 block 的 thread) 都共享一份 global memory。不同的 grid 则有各自的 global memory。

图5：CUDA 内存模型，硬件角度
从软件的角度来讲：
1. 线程处理器 (SP) 对应线程 (thread)。
2. 多核处理器 (SM) 对应线程块 (thread block)。
3. 设备端 (device) 对应线程块组合体 (grid)。

图6：CUDA 内存模型，软件角度
如下图6所示，所谓线程块内存模型在软件侧的一个最基本的执行单位，所以我们从这里开始梳理。线程块就是线程的组合体，它具有如下这些特点：
1. 块内的线程通过共享内存、原子操作和屏障同步进行协作 (shared memory, atomic operations and barrier synchronization)
2. 不同块中的线程不能协作。
如下图7所示的线程块就是由256个线程组成的，它执行的任务就是一个最基本的向量相加的一个操作。在线程块内，这256个线程的计算是彼此互相独立的，并行的。下面的这个，就是如何确定每个线程的索引 (在显存中的位置)。在计算完以后 (图中弯箭头的头部)，会设置一个时钟，将这256个线程的计算结果进行同步。

图7：一个256个线程组成的线程块
以上就是一个256位向量的加的操作的并行处理方法，得到最终的向量加的结果。
所谓网格 (grid)，其实就是线程块的组合体，如下图8所示。
1. 网格 (grid) 内的线程块是彼此互相独立，互不影响的。
2. 全局内存可以由所有的线程块进行访问。
CUDA 核函数由线程网格 (数组) 执行。每个线程都有一个索引，用于计算内存地址和做出控制决策。在计算完以后 (图中所有弯箭头的头部)，会设置一个时钟，将这N个线程块的计算结果进行同步。

图8：网格就是线程块的组合体
线程块 id & 线程 id：定位独立线程的门牌号
核函数需要确定每个线程在显存中的位置，我们之前提到 CUDA 的核函数是要在设备端来进行计算和处理的，在执行核函数时需要访问到每个线程的 registers (寄存器) 和 local memory (局部内存)。在这个过程中需要确定每一个线程在显存上的位置。所以我们需要像图9那样使用线程块的 index 和线程的 index 来确定线程在显存上的位置。

图9：使用线程块的 index 和线程的 index 来确定线程在显存上的位置
如图9所示，图9中的线程块索引是2维的，每个网格都由2×2个线程块组成；线程索引是3维的，每个线程块都由2×4×2个线程组成，所以代码应该是：

图10：线程Id计算
图10中：M=N=2，P,Q,S=2,4,2。
每个线程x的那一维应该是线程块的索引×线程块的x维度大小+线程的索引。(设备端线程x的那一维的索引)。
每个线程y的那一维应该是线程块的索引×线程块的y维度大小+线程的索引。(设备端线程y的那一维的索引)。
线程束 (warp)
前面我们提到，如图11所示的每一行由1个控制单元加上若干计算单元所组成，这些所有的计算单元执行的控制指令是一个。这其实就是个非常典型的 &#34;单指令多数据流机制&#34;。

图11：一个线程束 (warp)：采用单指令多数据流机制
单指令多数据流机制是说：执行的指令是一条，只不过不同的计算单元使用的数据是不一样的。而上面这一行，我们就称之为一个线程束 (warp)。
所以，SM 采用的 SIMT (Single-Instruction, Multiple-Thread，单指令多线程) 架构，warp (线程束) 是最基本的执行单元。一个 warp 包含32个并行 thread，这些 thread 以不同数据资源执行相同的指令。一个 warp 只包含一条指令，所以：warp 本质上是线程在 GPU 上运行的最小单元。
由于warp的大小为32，所以block所含的thread的大小一般要设置为32的倍数。
当一个 kernel 被执行时，grid 中的线程块被分配到 SM (多核处理器) 上，一个线程块的 thread 只能在一个SM 上调度，SM 一般可以调度多个线程块，大量的 thread 可能被分到不同的 SM 上。每个 thread 拥有它自己的程序计数器和状态寄存器，并且用该线程自己的数据执行指令，这就是所谓的 Single Instruction Multiple Thread (SIMT)，如图12所示。

图12：Single Instruction Multiple Thread (SIMT)
3 并行计算向量相加

下面我们就用一个实际的例子来看看 CUDA 编程具体是如何操作的。例子就是两个长度为N的张量相加，如下图13所示。

图13：两个张量相加
在 CPU 中完成相加的操作很简单：
// Compute vector sum C = A+B要在 GPU 中完成这一操作，首先我们想一下它是否适合使用 GPU，我们当时总结了四个特点：
1. 访问内存次数少，满足。
2. 控制指令简单，无复杂分枝预测，跳转指令，满足。
3. 计算指令简单，满足，是简单的加法操作。
4. 并行度高，满足，不同的之间不互相影响。
所以，向量相家的任务适合在 GPU 上编程。
再回顾下 GPU 运算步骤，如图4所示：
一个 CUDA 程序，我们可以把它分成3个部分：
第1部分是： 从主机 (host) 端申请 device memory，把要拷贝的内容从 host memory 拷贝到申请的 device memory 里面。
第2部分是： 设备端的核函数对拷贝进来的东西进行计算，来得到和实现运算的结果，图4中的 Kernel 就是指在 GPU 上运行的函数。
第3部分是： 把结果从 device memory 拷贝到申请的 host memory 里面，并且释放设备端的显存和内存。
如下：
#include <cuda.h>下面我们把这些内容细化到函数。
设备端代码:
1. 读写线程寄存器
2. 读写 Grid 中全局内存
3. 读写 block 中共享内存
主机端代码：
1. 申请显存，内存
2. Grid 中全局内存拷贝转移 (显存，内存互相拷贝)
3. 内存，显存释放
内存是插在主板上的内存插槽上的内存条，而显存是独立显卡上焊在显卡上的内存芯片。
申请显存的函数 cudaMalloc()：
在主机端完成显存的申请，得到相应的指针。

图14：申请显存的函数 cudaMalloc()
释放显存的函数 cudaFree( )：
将指向显存的指针释放掉。

图15：释放显存的函数 cudaFree( )
内存和显存之间互相拷贝的函数 cudaMemcpy( )：
参数含义是：终点的指针，起点的指针，拷贝的大小，模式 (主机端到设备端，设备端到主机端，设备端之间的拷贝)

图16：内存和显存之间互相拷贝的函数 cudaMemcpy( )
以上三个函数是 CUDA 帮我们写好的，如果调用的话需要先：
# include cuda.h下面就是具体的 C++ 代码实现：
申请内存的大小是 n *sizeof(float)，定义3个指针 A_d，B_d，C_d。
cudaMalloc 函数需要传入 1. 指针的指针 (指向申请得到的显存的指针)。2. 申请显存的大小。 所以分别传入 &A_d 和 size。同理后面依次传入 &B_d 和 size，&C_d 和 size。
cudaMemcpy 函数需要传入 1. 终点的指针。2. 起点的指针。3. 拷贝的大小。4. 模式。 所以分别传入 A_d, A, size, cudaMemcpyHostToDevice。同理后面依次传入 B_d, B, size, cudaMemcpyHostToDevice 和 C, C_d, size, cudaMemcpyHostToDevice。
最后把设备端申请的显存都释放掉。cudaFree 函数需要传入设备端申请显存的指针，即 A_d，B_d，C_d。
void下面我们进入最重要的部分，即：如何自己书写一个 kernel 函数。
核函数调用的注意事项
1. 在 GPU 上执行的函数。
2. 一般通过标识符 __global__ 修饰。
3. 调用通过<<<参数1,参数2>>>，用于说明内核函数中的线程数量，以及线程是如何组织的。
4. 以网格 (Grid) 的形式组织，每个线程格由若干个线程块 (block) 组成，而每个线程块又由若干个线程 (thread) 组成。
5. 调用时必须声明内核函数的执行参数。
6. 在编程时，必须先为 kernel 函数中用到的数组或变量分配好足够的空间，再调用 kernel 函数，否则在 GPU 计算时会发生错误。
CUDA 编程的标识符号
不同的表示符号对应着不同的工作地点和被调用地点。核函数使用 __global__ 标识，必须返回 void。__device__ & __host__ 可以一起用。

图17：CUDA 编程的标识符号
下面，按照我们刚才的对核函数的介绍，我们展示了向量相加的代码。
代码讲解：
首先，看到 __global__ 标识，返回的是 void，就意味着 vecAddKernel 函数是一个在 host 端调用，在 device 端执行的核函数。它的三个参数就是我们之前申请好的指向三段显存的指针。
通过 int i= threadIdx.x+ blockDim.x* blockIdx.x; (线程的索引，线程块的索引，线程块维度的大小) 来计算好要访问的线程的索引的位置。
那么如何在主机端调用呢？我们使用尖括号**<<<网格 grid 维度，线程块 block 维度>>>**来包括：线程块数 ceil(n/256) 和一个线程块的线程数256。

图18：向量相加的代码
第1步主机端 __host__ 修饰：申请显存，内存。显存，内存的互相拷贝。内存，显存释放。比如图19中申请的网格是 ceil(n/256) 维的代表一个网格有 ceil(n/256) 个线程块；线程块是256维的，代表一个线程块有256个线程。
第2步设备端 __global__ 修饰：计算索引绝对位置，并行计算。

图19：主机端和设备端代码
详细地讲，核函数只能在主机端调用，调用时必须申明执行参数。调用形式如下：
Kernel<<<Dg,Db, Ns, S>>>(param <<<>>> 运算符内是核函数的执行参数，告诉编译器运行时如何启动核函数，用于说明内核函数中的线程数量，以及线程是如何组织的。
<<<>>> 运算符对 kernel 函数完整的执行配置参数形式是 <<<Dg, Db, Ns, S>>>
1. 参数 Dg 用于定义整个 grid 的维度和尺寸，即一个 grid 有多少个 block。为 dim3 类型。Dim3 Dg(Dg.x, Dg.y, 1) 表示grid中每行有 Dg.x 个 block，每列有 Dg.y 个 block，第三维恒为1(目前一个核函数只有一个grid)。整个 grid 中共有 Dg.x*Dg.y 个 block，其中 Dg.x 和 Dg.y 最大值为65535。
2. 参数 Db 用于定义一个 block 的维度和尺寸，即一个 block 有多少个 thread。为 dim3 类型。Dim3 Db(Db.x, Db.y, Db.z) 表示整个 block 中每行有 Db.x 个 thread，每列有 Db.y 个 thread，高度为 Db.z。Db.x 和 Db.y 最大值为512，Db.z 最大值为62。一个 block 中共有 Db.x*Db.y*Db.z 个 thread。计算能力为1.0,1.1的硬件该乘积的最大值为768，计算能力为1.2,1.3的硬件支持的最大值为1024。
3. 参数 Ns 是一个可选参数，用于设置每个 block 除了静态分配的 shared Memory 以外，最多能动态分配的shared memory 大小，单位为 byte。不需要动态分配时该值为0或省略不写。
4. 参数 S 是一个 cudaStream_t 类型的可选参数，初始值为零，表示该核函数处在哪个流之中。
最后我们简单介绍下 CUDA 编程如何执行编译的过程。因为我们之前在 CPU 上编程，使用 g++ 或 gcc 进行编译，再通过 link 生成可执行程序。那么在 GPU 端，编译器就是 NVCC (NVIDIA Cuda compiler driver)。
通常我们会把和 GPU 相关的头文件放在 .h 文件里，把设备端执行的程序 (__global__ 定义的函数) 放在 .cu 文件里，这些程序我们用 NVCC 来进行编译。主机端的程序放在 .h 和 .cpp 里面，这些程序我们可以继续用 g++ 或 gcc 来进行编译。
通常我们有这几种编译的方法：
1. 逐个文件编译 (GPU 和 CPU 的程序都编译成 .o 文件。最后把它们汇总在一起，并 link 为一个可执行文件 .exe)，但是这只适用于文件数较少的情况，当文件数较多时，这种办法就显得比较复杂。
2. 使用 cmake 方式编译，写一个 cmake.txt，下文有介绍。

图20：CUDA 编程如何执行编译的过程
CUDA 中 threadIdx，blockIdx，blockDim，gridDim 的使用
1. threadIdx是一个uint3类型，表示一个线程的索引。
2. blockIdx是一个uint3类型，表示一个线程块的索引，一个线程块中通常有多个线程。
3. blockDim是一个dim3类型，表示线程块的大小。
4. gridDim是一个dim3类型，表示网格的大小，一个网格中通常有多个线程块。
下面这张图21比较清晰的表示的几个概念的关系：

图21：几个变量的关系
cuda 通过<<< >>>符号来分配索引线程的方式，我知道的一共有15种索引方式。
4 实践

4.1 向量相加 CUDA 代码

这一节我们通过一个实例直观感受下 CUDA 并经计算究竟能使这些计算简单，并行度高的操作加速多少。
我们先看一下 CPU 执行向量相加的代码：
#include <iostream>注释：
float*a =(float*)malloc(size); 分配一段内存，使用指针 a 指向它。
for 循环产生一些随机数，并放在分配的内存里面。
vecAdd(float* A,float* B,float* C,int n) 要输入指向3段内存的指针名，也就是 a, b, c。
gettimeofday 函数来得到精确时间。它的精度可以达到微妙，是C标准库的函数。
最后的 free 函数把申请的3段内存释放掉。
编译：
g++ -O3 main_cpu.cpp -o VectorSumCPU我们再看一下 CUDA 执行向量相加的代码：
#include <iostream>注释：
首先要用 __global__ 来修饰。
vecAdd(float* A,float* B,float* C,int n) 要输入指向3段显存的指针名，也就是 d_a, d_b, d_c。
float*da =NULL; 定义空指针。
cudaMalloc((void**)&da, size); 申请显存，da 指向申请的显存，注意 cudaMalloc 函数传入指针的指针 (指向申请得到的显存的指针)。
cudaMemcpy(da,a,size,cudaMemcpyHostToDevice) 把内存的东西拷贝到显存，也就是把 a, b, c 里面的东西拷贝到 d_a, d_b, d_c 中。
int threadPerBlock =256; int blockPerGrid =(n + threadPerBlock -1)/threadPerBlock; 计算线程块和网格的数量。
vecAddKernel <<< blockPerGrid, threadPerBlock >>> (da, db, dc, n); 调用核函数。
gettimeofday 函数来得到精确时间。它的精度可以达到微妙，是C标准库的函数。
最后的 free 函数把申请的3段内存释放掉。
编译：
/usr/local/cuda/bin/nvcc main_gpu.cu -o VectorSumGPU4.2 实践向量相加

编译之后得到可执行文件 VectorSumCPU 和 VectorSumGPU 之后，我们可以执行一下比较下运行时间 (注意要在 linux 下运行)：
在 CPU 下，执行1000000000次加需要4.18秒。
./VectorSumCPU 1000000000在 GPU 下，执行1000000000次加只需要1.6e-05秒，哇。
(base) wjh19@iccv:~/mage/CUDA/db$ ./VectorSumGPU 1000000000GPU 对于计算简单，并行度高的计算果然可以大幅提速！！！
在 CPU 下，执行1000次加需要1e-06秒。
(base) wjh19@iccv:~/mage/CUDA/db$ ./VectorSumCPU 1000在 GPU 下，执行1000次加需要1.3e-05秒。
(base) wjh19@iccv:~/mage/CUDA/db$ ./VectorSumGPU 1000GPU 对于少量计算效率反倒不如 CPU。
为何AI芯片公司不支持CUDA？
实现真正意义上的完全兼容是极不可能的。最多是常用API和功能类似，减少用户移植成本。CUDA本身涵盖的功能非常广泛，硬件功能上就几乎体现了NV自家GPGPU的所有可能性，再加上驱动和软件上层封装（各种库，比如cuBLAS，cuFFT，cuDNN之类），以及完备的开发工具套件（编译器、调试器、profiler等等）。这些东西就算是全部开源，让各家移植支持自己的硬件，多数公司恐怕也是有心无力。更别说CUDA很多功能与硬件深度耦合，硬件设计不一致，靠软件封装来保持一致性，工作量真不是一般公司消化得了的。更何况其中很多东西并不公开，各家无从下手，功能就更难做得一致了。这还只是功能的一致性问题，性能上就更难保证了。李逵和李鬼，靠长的像是不行的，抡起板斧来就露相了……
其实就算是NV，各代硬件之间发生较大的功能改动时，更新工作量也不小。NV在硬件微架构和指令集上迭代很快，几乎每隔一两代就会有较大的功能变化。这样底层很多东西都要跟着调整。NV的PTX是一个很好的隔离机制，底层指令集之类的改动多数可以在PTX这层兼容，这样上层就不用动了。不过，底层仍然有PTX覆盖不到的改动（主要是运行逻辑的改变，比如Independent Thread Scheduling这种，同样的PTX代码在不同架构上行为不一致）。或者是有些程序没有内嵌PTX文件，那也没法兼容。而驱动或上层编程接口之类的改动跟PTX就没啥关系，当然也没法靠PTX覆盖。
老黄曾说NV是软件公司，也不是随便说说。至少CUDA的各种功能，并不都是跟着硬件版本走，很多都是软件层的封装。新硬件出来，CUDA一般会发个大版本用以提供相应支持。但CUDA自身软件层封装的功能也会不断新增和改进，也会有相应的版本（比如11.1，11.2之类）。只不过软件层的向后兼容可以做的比较好，用户通常不太关注而已。
这些硬件或软件上不同版本的差异，多多少少都会影响到用户的使用。所以很多软件包括AI框架甚至都要安装对应的CUDA版本，否则就可能出错。NV自家都不能完全兼容，外人还想兼容，那就太难了。这些对NV也是沉重的负担，感觉NV应该是在有意识的压缩产品支持周期，比如SM50（Maxwell架构）是14年首发，15、16年甚至再往后都有很多卡还在卖的。可是去年的CUDA 11.2版本已经把deprecated了，这周期也就四五年而已……
其次，API这个东西，应该是没有版权的。之前Google与Oracle在Java的API版权上打过旷日持久的官司，去年最终裁决结果认为API本身是没有版权的。当然这应该说的是API的命名和总体结构设计没有版权，具体API的内部实现可能还是有的。实际上各家API的“借鉴”其实非常普遍，普遍到大家都习以为常，见怪不怪了。
首先CUDA的编程模型很大程度复用了C语言，所以有cudaMalloc、cudaFree、cudaMemset这种定制版的“C API”。CUDA里的数学函数也多数沿用了cmath里的形式，比如exp()是double函数，expf()是float版的exp。cuBLAS就不用说了，多数API都源自LAPACK（虽然BLAS的这套API现在都快成标准接口了）。cuFFT的API与FFTW虽然有不小的差别，但两者之间的传承关系也很明显。这些API基本都来自曾经非常流行的开源库，应该说还算是常规操作。
但这还没完。熟悉Intel开发工具的人应该知道Intel有MKL（Math Kernel Library）和IPP（Intel Integrated Performance Primitives）。MKL相当于Intel的BLAS+数学函数库，IPP主要是做图像和信号处理的，有ippi（image processing）与ipps（signal processing）。NVIDIA有个库叫NPP，也有nppi与npps。Intel还有个很有名的并行库叫TBB，可以基于模板做并行transform、reduce、scan之类的泛型操作，还能做并行任务拆分和调度。CUDA有个库叫Thrust，也是基于模板的泛型编程，也可以做并行transform、reduce、scan…… 当然，TBB和Thrust在接口上差别还是很大的，而且TBB从功能和可编程性上讲比Thrust要完整得多。但你要说Thrust没借鉴过TBB……
这里也不是针对NV。其实绝大多数功能相似的库或软件产品，具有类似的API是再正常不过的事情。比如各种计算机代数系统（以matlab为代表），各种AI框架等等。编译器还会造接口去接受其他编译器的参数输入格式呢！只要不是直接抄代码，接口类似，但内部实现有差别，其实也不是什么见不得人的事情。
最后，API只是个入口，里面也会有非常多的坑。把自家API跟用户常用的主流API做得像，可以大大减少用户的学习成本和移植工作量，甚至一个文本替换脚本就能搞定大部分。但前面也说了，功能的差异是不可避免的。99%相同，1%不同，看起来好像还行。但实际用起来，99%相同的部分占用时间1%，那1%不同的部分埋下的坑可能会耗费你99%的时间…… 根本不知道到底是哪1%不一样……。
所以我觉得API的复用，还是要格外慎重。有很大把握做到一致的可以复用，有差别的还是尽量区分开，否则真的是遗患无穷。有些公司意识不太好，老想着开始先蒙混过关，假装一样，将来有机会再慢慢修补。其实这是非常短视的做法，用户谁没事老关注你各个版本什么变化，一不小心踩坑真是心累又心碎。这种交付就非常不靠谱，严重不推荐。
其实用户最常用的CUDA核心功能也没有太多，能把这些支持好就不错。至少可以覆盖多数用户的多数需求。当前大部分的AI硬件公司应该都没啥2C业务，客户支持的压力会小很多。毕竟如果是2B的话，用户水准一般都还可以，做一些定制和差异化，能有自己的一些优势，也能保证先活下去。而且现在上层框架为了兼容性，接口一般不会对CUDA做那么深的定制。有这些基础功能和通用接口支撑，多数功能移植起来应该也没那么难。这些问题很多都不是技术难度问题，更多的是工作量的问题。CUDA里有些复杂功能，实在不支持就放弃算了。有些市场，吃不下就不要硬啃，真的划不来……
生态也不都是靠用的人多堆出来的，技术先进性和技术发展方向的话语权也很重要。花这么多功夫去做兼容和移植，还不如多研究研究怎么把自己的技术优势发挥出来。
这当然是一个好想法，而且事实上，AMD也在往这个方向上去想。
AMD最近支持了一个叫做HIP的编程模式，这东西跟CUDA没什么区别，甚至于可以直接搜索替换CUDA程序完成。见下表，出处：
比较CUDA、OpenCL和HiP的语法
VINCENT HINDRIKSEN于2016年4月5日发布，附1条评论
CUDA和OpenCL都是众所周知的GPGPU语言。不幸的是，这两种语言之间存在一些细微的差异，如下所示。
可能听说过HiP，AMD为支持现代AMD斐济GPU和CUDA设备而开发的语言。CUDA可以（大部分自动）转换为HiP，从那时起，代码还支持AMD高端设备。
为了概述HiP与其他API的比较，Ben Sanders进行了概述。下面是CUDA、OpenCL和HiP的表格，稍作修改，使其更加完整。语言HC和C++AMP可以在原文中找到。
Both CUDA and OpenCL are well-known GPGPU-languages. Unfortunately there are some slight differences between the languages, which are shown below.
You might have heard of HiP, the language that AMD made to support both modern AMD Fiji GPUs and CUDA-devices. CUDA can be (mostly automatically) translated to HiP and from that moment your code also supports AMD high-end devices.
To give an overview how HiP compares to other APIs, Ben Sanders made an overview. Below you’ll find the table for CUDA, OpenCL and HiP, slightly altered to be more complete. The languages HC and C++AMP can be found in the original.

当然，如果所有类似的API都不应该使用相同的措辞，那么讨论也是理所当然的。最好的做法是最好的混合，因为CUDA的“共享”比OpenCL的“本地”更清晰。另一方面，OpenCL在位置和维度上的功能（get\u global\u id（0）等）往往比CUDA提供的功能更受欢迎。CUDA的“<<>>>”打破了所有C/C++编译器，因此很难制作IDE插件的前端。
更好地理解CUDA和OpenCL之间的差异，同时也有助于了解HiP是如何进入画面的。
编译过程
CUDA的核心部分是专门开发的C编译器。C语言对大多数开发人员都十分熟悉的，可以使编程人员专注于开发并行程序而不是处理负责的图形API。为了简化开发，CUDA的C编译器允许程序员将CPU 和 GPU的代码混合记录到一个程序文件中。一些简单代码被增加到的C程序中，通知CUDA编译器哪些函数由CPU处理，哪些为GPU编译。然后程序被CUDA编译器编译，而CPU处理的代码则由开发者的标准C编译器编。

整个编译过程需要几个阶段。首先，所有的代码都要让CPU来处理，这些都会从文件中提取，并且都会通过标准的编译器。用于GPU处理的代码，首先要转换成中间媒介性语言——PTX。中间语言更像是一种汇编程序，并且能够中和潜在的无效代码。在最后的阶段，中间语言会转换成指令。这些指令会被GPU所认同，并且会以二进制的形式被执行。
NVIDIA CUDA技术基于一种全新的用于开拓GPU运算性能的软件架构，CUDA程序执行时，GPU作为主CPU的协处理器工作，GPU可以处理大量的并行信息，同时CPU组织、解释、与后台交流要处理的信息。应用程序的计算密集型部分要被执行很多次，每次由主应用程序提交的不同数据，经过编译后由GPU并行执行。

CUDA可以用来生产资源，比如生成几何图形，在程序中进行材质贴图等等，同时这些也可以传递到传统的图形API来生成。3D图形API也可以将渲染后的结果发送到CUDA进行后续处理。CUDA本身就是基于图形芯片，而这种图形芯片也具备通用计算的能力。这里有许多交互性的例子，在GPU的显存中存储数据将更具优势，系统可以绕过速度相对较慢的PCI-Express总线，直接调用显存中的数据。
另一方面需要指出的是，针对这种在显存内的资源共享来说，图形数据并不总是短小精悍的，并且也会给程序员带来一些头痛的问题。例如，转换分辨率或者颜色深度时，图形数据就有优先权。因此，如果在缓冲中的资源需要增加的时候，驱动程序会毫不犹豫的将应用程序分配给CUDA来执行。这样CUDA计算和图形处理就不会产生冲突。对于数据的分配和管理，CUDA还有待于更进一步完善。尤其是当系统中有几个GPU的时候，首先就无法使用SLI模式了，只能用一颗GPU来完成显示工作。不过这也是避免系统混乱的最好办法。
CUDA API其本质上来讲是由各种操作显存的函数组成的。cudaMalloc用来分配内存，cudaFree用来释放内存，cudaMemcpy用来互相拷贝内存和显存之间的数据。
名词解析
thread线程在CUDA里定义thread线程的概念。因为这里所指的线程，与传统的“CPU线程”是有所区别的，同时也不是在GPU文章里所指的“线程”。在GPU中，线程是最基本的元素，贯穿于数据处理的始终。与CPU中的线程不同，CUDA的线程是非常轻巧微小的，这就意味着，单独的线程处理起来会非常的简单快速。
warp 不要试图从字面理解warp的概念，因为仅仅是一种象征性的比喻，一个由NVIDIA自创的术语罢了。NVIDIA的意思是CUDA的整个处理工作，就像是一架织布机，织物在织布机内快速的来回穿过。
在CUDA中的一个warp，是由32个线程组成的。这也是SIMD处理中，数据的最小封包单位。CUDA采用的是多处理并行架构，主旨就是尽量能并行处理更多的数据。
grid 栅格，将许多个block块封装起来。这种数据机制的优势就在于可以同时在GPU中处理多个block块。这种方式将GPU所有硬件资源都紧密的联系在一起。
从CUDA原理中得到的优化PC启示:
PhysX物理加速也是建立在CUDA技术之上的,CUDA运行时不但要占用CPU资源,还要在显卡的GPU和显存中划分出一定的资源来用做GUP计算如:物理加速,通用计算等.通过对CUDA的分析就不难理解为什么9500GT级别的显卡在开启物理加速以后为什么性能不升反降的原因。
随着,大量游戏对物理加速的支持和许多软件开始对GPU加速的支持,显卡将不单是图形处理,GPU的性能,流处理器的数量,和显存的大小将直接影响着使用者游戏和软件的运行速度.在新应用下选择显卡应该着重考虑以下几方面:
1.显卡的GPU性能,性能强大的GPU才能够更好的运行物理加速和CUDA通用计算,特别是在运行3D游戏时,GPU要同时负责图形加速和物理加速,对GPU性能有一定要求。
2.流处器的数量和频率,流处理器数目越多频率越高,并行计算能力越强。
3.显存的大小和速度,大容量的高速显存在CUDA计算中能够在更短的时间内交换更多的数据,在3D游戏中也不会因为显存太小而影响性能。
NVIDIA推荐的CUDA和物理加速显示256M的9600GT以上的显卡,但是在目前来看,物理加速和CUDA要能够流畅运行的话,一块512M DDR3的9600GT是基本的要求,512M显存才有足够的显存空间给CUDA作为GPU计算内存使用.而如果显存只有256M,在CUDA计算量大的时候将直接影响性能,如果是3D游戏,图形处理也将受到影响。
参考文献链接
https://baike.baidu.com/item/CUDA/1186262?fr=aladdin
https://mp.weixin.qq.com/s/kxYSw_fR4QMZ2-O5fvOR8g
https://www.zhihu.com/question/461354739/answer/1964488472
https://blog.csdn.net/liusaoyue/article/details/5100656
1.深蓝学院课程讲解：https://www.shenlanxueyuan.com/course/410
2. D. Kirk and W. Hwu, “Programming Massively Parallel Processors –A Hands-on Approach, Second Edition”
3. CUDA by example, Sanders and Kandrot
4. Nvidia CUDA C Programming Guide：https://docs.nvidia.com/cuda/cuda-c-programming-guide/
5. CS/EE217 GPU Architecture andProgramming

		自动登录	找回密码
密码			立即注册

CUDA架构与应用杂谈

本帖子中包含更多资源