GPU编程

1、GPU  ：CPU的协处理器（通过PCle总线协同工作）

特点  ： 

1）多ALU核，并行处理计算，

2）线程较轻，数量多，实现大规模的并发

3）少量Cache （利用逻辑控制思维替代了缓存单元，ps CPU中需要很多的逻辑缓存区，存储临时数据）

    作用： 不是为了存储临时数据，而是提高线程访问效率从而进行访问节点的合并，存储其访问的指针（指向dram）

绿色：计算单元
橙色：存储单元（硬盘+缓存）
黄色：控制单元

 

从图中可得：GPU中，每个控制器都配有缓存，有很多控制器

 

特点：CPU秉承着低延时性，CPU适合用于逻辑控制，串行的计算场景

           GPU的思路是高吞吐量，GPU适合用于大规模的并行计算场景

 

 

2、 若假定C语言工作的对象是CPU和内存条（接下来,称为主机内存），那么CUDA C工作的的对象就是GPU及GPU上的内存 。

     CUDA是操纵GPU的工具，即程序（软件体系），本质是个异构模型，协同控制GPU 和CPU 。

      CUDA的并行计算通过控制成千上万个线程实现的。

由图可得 ：  一个kernel 对用一个 grid （网格） ，一个grid 有很多 block (线程块） , 各block是并行执行的，block间无法通信，                      也没有执行顺序。

                     一个GPU ≈ 16个 SM ( streaming processor )  ≈ 上千个 SM (streanming multiprocessor)  。

、

 

由图可得，CUDA分为三层（函数库，运行API,驱动API）

 

3、核函数： 在GPU 上执行的函数称之为核函数

1）  调用方法： 标识符 __global__ 修饰  :

      

       调用 <<<参数1，参数2>>> , 表示: 内核函数中线程数量和组织方式

         # 一个线程grid 中只有一个线程block， 一个线程block 中只有一个 一个threads

2） 必须预先位kernel函数分配好足够的空间

 

dim3结构类型
        1. dim3是基亍uint3定义的矢量类型，相当亍由3个unsigned int型组成的结构体。uint3类型有三个数据成员unsigned int x; unsigned int y; unsigned int z;
        2. 可使用亍一维、二维或三维的索引来标识线程，构成一维、二维或三维线程块。
        3. dim3结构类型变量用在核函数调用的<<<,>>>中。
        4. 相关的几个内置变量
        4.1. threadIdx，顾名思义获取线程thread的ID索引；如果线程是一维的那么就取threadIdx.x，二维的还可以多取到一个值threadIdx.y，以此类推到三维threadIdx.z。
        4.2. blockIdx，线程块的ID索引；同样有blockIdx.x，blockIdx.y，blockIdx.z。
        4.3. blockDim，线程块的维度，同样有blockDim.x，blockDim.y，blockDim.z。
        4.4. gridDim，线程格的维度，同样有gridDim.x，gridDim.y，gridDim.z。
        5. 对于一维的block，线程的threadID=threadIdx.x。
        6. 对于大小为（blockDim.x, blockDim.y）的 二维 block，线程的threadID=threadIdx.x+threadIdx.y*blockDim.x。
        7. 对于大小为（blockDim.x, blockDim.y, blockDim.z）的 三维 block，线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。
        8. 对于计算线程索引偏移增量为已启动线程的总数。如stride = blockDim.x * gridDim.x; threadId += stride。
函数修饰符
        1. __global__，表明被修饰的函数在设备上执行，但在主机上调用。
        2. __device__，表明被修饰的函数在设备上执行，但只能在其他__device__函数或者__global__函数中调用。
常用的GPU内存函数
cudaMalloc()
        1. 函数原型： cudaError_t cudaMalloc (void **devPtr, size_t size)。
        2. 函数用处：与C语言中的malloc函数一样，只是此函数在GPU的内存你分配内存。
        3. 注意事项：
        3.1. 可以将cudaMalloc()分配的指针传递给在设备上执行的函数；
        3.2. 可以在设备代码中使用cudaMalloc()分配的指针进行设备内存读写操作；
        3.3. 可以将cudaMalloc()分配的指针传递给在主机上执行的函数；
        3.4. 不可以在主机代码中使用cudaMalloc()分配的指针进行主机内存读写操作（即不能进行解引用）。
cudaMemcpy()
        1. 函数原型：cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)。
        2. 函数作用：与c语言中的memcpy函数一样，只是此函数可以在主机内存和GPU内存之间互相拷贝数据。
        3. 函数参数：cudaMemcpyKind kind表示数据拷贝方向，如果kind赋值为cudaMemcpyDeviceToHost表示数据从设备内存拷贝到主机内存。
        4. 与C中的memcpy()一样，以同步方式执行，即当函数返回时，复制操作就已经完成了，并且在输出缓冲区中包含了复制进去的内容。
        5. 相应的有个异步方式执行的函数cudaMemcpyAsync()，这个函数详解请看下面的流一节有关内容。
cudaFree()
        1. 函数原型：cudaError_t cudaFree ( void* devPtr )。
        2. 函数作用：与c语言中的free()函数一样，只是此函数释放的是cudaMalloc()分配的内存。
        下面实例用于解释上面三个函数

 

GPU内存分类: 全局内存，   固定内存  ，  共享内存 ，常量内存，纹理内存。

全局内存
        通俗意义上的设备内存。

共享内存
        1. 位置：设备内存。
        2. 形式：关键字__shared__添加到变量声明中。如__shared__ float cache[10]。
        3. 目的：对于GPU上启动的每个线程块，CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。

常量内存
        1. 位置：设备内存
        2. 形式：关键字__constant__添加到变量声明中。如__constant__ float s[10];。
        3. 目的：为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下，用常量内存替换全局内存能有效地减少内存带宽。
        4. 特点：常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时，静态地分配空间。
        5. 要求：当我们需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol()，而cudaMemcpy()会复制到全局内存。
        6. 性能提升的原因：
        6.1. 对常量内存的单次读操作可以广播到其他的“邻近”线程。这将节约15次读取操作。（为什么是15，因为“邻近”指半个线程束，一个线程束包含32个线程的集合。）
        6.2. 常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。

纹理内存
        1. 位置：设备内存
        2. 目的：能够减少对内存的请求并提供高效的内存带宽。是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序设计，意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。如下图：

        3. 纹理变量（引用）必须声明为文件作用域内的全局变量。
        4. 形式：分为一维纹理内存 和 二维纹理内存。
        4.1. 一维纹理内存
        4.1.1. 用texture<类型>类型声明，如texture<float> texIn。
        4.1.2. 通过cudaBindTexture()绑定到纹理内存中。
        4.1.3. 通过tex1Dfetch()来读取纹理内存中的数据。
        4.1.4. 通过cudaUnbindTexture()取消绑定纹理内存。
        4.2. 二维纹理内存
        4.2.1. 用texture<类型,数字>类型声明，如texture<float，2> texIn。
        4.2.2. 通过cudaBindTexture2D()绑定到纹理内存中。
        4.2.3. 通过tex2D()来读取纹理内存中的数据。
        4.2.4. 通过cudaUnbindTexture()取消绑定纹理内存。

固定内存
        1. 位置：主机内存。
        2. 概念：也称为页锁定内存或者不可分页内存，操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此操作系统能够安全地使某个应用程序访问该内存的物理地址，因为这块内存将不会破坏或者重新定位。
        3. 目的：提高访问速度。由于GPU知道主机内存的物理地址，因此可以通过“直接内存访问DMA（Direct Memory Access)技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU介入。因此DMA复制过程中使用固定内存是非常重要的。
        4. 缺点：使用固定内存，将失去虚拟内存的所有功能；系统将更快的耗尽内存。
        5. 建议：对cudaMemcpy()函数调用中的源内存或者目标内存，才使用固定内存，并且在不再需要使用它们时立即释放。
        6. 形式：通过cudaHostAlloc()函数来分配；通过cudaFreeHost()释放。
        7. 只能以异步方式对固定内存进行复制操作。

 

4、原子性：  即最小的操作单位

函数调用： 如atomicAdd（addr,y)将生成一个原子的操作序列，这个操作序列包括读取地址addr处的值，将y增加到这个值，以                       及将结果保存回地址addr。

常用线程操作函数：

1. 同步方法__syncthreads()，这个函数的调用，将确保线程块中的每个线程都执行完__syscthreads()前面的语句后，才会执行下一条语句。
使用事件来测量性能
        1. 用途：为了测量GPU在某个任务上花费的时间。CUDA中的事件本质上是一个GPU时间戳。由于事件是直接在GPU上实现的。因此不适用于对同时包含设备代码和主机代码的混合代码设计。
        2. 形式：首先创建一个事件，然后记录事件，再计算两个事件之差，最后销毁事件。如：

 

5、流  ： CUDA流表示一个GPU操作队列，每个流是GPU上的一个任务。

并发和并行： 并发指代同一时刻运行不同的任务， 并行指代同一时刻运行同一任务。

硬件前提： 支持设备重叠功能的GPU （即，执行核函数同时，能够支持设备和主机间的复制操作）

声明与创建：声明cudaStream_t stream;，创建cudaSteamCreate(&stream);。
        6. cudaMemcpyAsync()：前面在cudaMemcpy()中提到过，这是一个以异步方式执行的函数。在调用cudaMemcpyAsync()时，只是放置一个请求，表示在流中执行一次内存复制操作，这个流是通过参数stream来指定的。当函数返回时，我们无法确保复制操作是否已经启动，更无法保证它是否已经结束。我们能够得到的保证是，复制操作肯定会当下一个被放入流中的操作之前执行。传递给此函数的主机内存指针必须是通过cudaHostAlloc()分配好的内存。（流中要求固定内存）
        7. 流同步：通过cudaStreamSynchronize()来协调。
        8. 流销毁：在退出应用程序之前，需要销毁对GPU操作进行排队的流，调用cudaStreamDestroy()。
        9. 针对多个流：
        9.1. 记得对流进行同步操作。
        9.2. 将操作放入流的队列时，应采用宽度优先方式，而非深度优先的方式，换句话说，不是首先添加第0个流的所有操作，再依次添加后面的第1，2,…个流。而是交替进行添加，比如将a的复制操作添加到第0个流中，接着把a的复制操作添加到第1个流中，再继续其他的类似交替添加的行为。
        9.3. 要牢牢记住操作放入流中的队列中的顺序影响到CUDA驱动程序调度这些操作和流以及执行的方式。
技巧
        1. 当线程块的数量为GPU中处理数量的2倍时，将达到最优性能。
        2. 核函数执行的第一个计算就是计算输入数据的偏移。每个线程的起始偏移都是0到线程数量减1之间的某个值。然后，对偏移的增量为已启动线程的总数。

 

 

 

 

 

问题： GPU中存在IO瓶颈？ （GPU的存储有限，所以数据需要从外部导入；导入的过程带来IO的问题，使得IO的时间远大于GPU处理的时间）

1、 算法层面： 优化显存使用（ 只copy 必要数据，优化数据结构等）

2、优化CUDA，即从控制器上优化内存的申请，共享等

3、优化硬件。放弃Pci-e，使用ibm的 power 架构，直接使用 nv-link进行内存传输。