CUDA学习笔记：CUDA中影响内核性能的几个因素

一.GPU资源描述

        众所周知，GPU是CPU的协处理器，是一个硬件加速器。在CUDA里，其容量通过CUDA核心数量和内存大小来描述；相应的，其性能由峰值计算性能由峰值计算性能和内存带宽来描述。

        以PASCAL架构的GTX1080Ti为例：

        其CUDA核心数量3584个(多处理器数量(28SMs)* 每个多处理器上的核心数量(128))，

        显存11GB；

        峰值浮点数计算能力为11.5TFlops(核心频率(1632.5MHz) * (CUDA核心数量(3584)* 浮点单元数量(256))* 2 OPS / 周期)，

        显存带宽484GB/s(GPU芯片数量(2)* 显存等效频率(5505MHz)* 显存位宽(352Bit)/ 8000 = 显存带宽(484.44GB/s))。

二.CUDA编程模型

        内存管理上，CUDA的内存分配函数与C语言几乎一样，其中与内核性能相关的函数为cudaError_tcudaMemcpy(void* dist, const void* src, size_t count, cudaMemcpyKindkind)，其中king有cudaMemcpyHostToHost，cudaMemcpyHostToDevice，cudaMemcpyDeviceToHost，cudaMemcpyDeviceToDevice这几种，这个函数是同步执行的，会阻塞主机应用程序。

        线程管理上，CUDA使用kernel<<grid,block>>的方式。其中，grid可以是1D，2D，3D，表示其上blocks的数量以及排列方式，block也可以是1D，2D，3D，表示其上threads的数量和排列方式，线程是计算的最小单元。CUDA使用这种方式将任务拆分成若干并行子任务，并分配到每个线程上。

三.CUDA执行模型

        GPU架构是围绕一个流式多处理器(SM)的可扩展阵列搭建的，SM的关键组建是：CUDA核心，共享内存/一级缓存，寄存器文件，加载/存储单元，特殊功能单元，线程束调度器。

        当某个block被调度到一个SM上时，其中的线程就只会在这个指定的SM上执行，多个blocks可能被调度到同一块SM上。CUDA采用单指令多线程的架构来管理和执行线程，在SM上，每32个threads为一组，被称为warp。一个warp中的所有线程都执行相同的指令。每个线程都有自己的指令地址计数器和寄存器状态，利用自身的数据执行当前的指令。

四.影响内核性能的因素

1.线程分化

        如上所述一个warp都执行相同的指令，而如果我们使用例如if...then...else、for和while这样的控制流结构，那么因为GPU没有复杂的分支预测机制，如果满足执行条件的threads数量小于一个warp的大小，warp中其他不满足的条件的threads就都会停止执行，这样就会降低降低内核的执行效率。

2.延迟隐藏

        Little'sLaw告诉我们：

        所需warp数量=延迟*吞吐量

        这是由于当某个warp阻塞时，SM的线程束调度器会选取其他满足条件的warp执行，而GPU的线程是轻量级的，在线程中的切换代价几乎可以忽略不计。

        于此，在SM中，一个线程中有很多独立的指令(条件容易满足)，很多并发地符合条件的线程(线程数量多)，都可以提高线程执行并行提升内核性能。

        同样的，更多的独立内存操作，创建更多并发的活跃线程/线程束，可以提高内存加载并行提升内核性能。

3.占用率

        占用率=活跃warp数量/最大warp数量

        为了提高占用率，需要调整block的配置或重新调整资源的使用情况，以允许更多的warp同时处于活跃状态，具体而言：

                (1).保持每个block中的threads数量是warp大小的倍数

                (2).避免block太小，至少要有128或256各线程，(在CUDA6.0以上似乎以256个的倍数为最佳)

                (3).根据内核资源的需求调整块大小

                (4).block的数量要远远多于SM的数量，从而提升并行

4.同步

        系统上，等待主机和设备完成所有的工作；块上，在设备执行过程中等待一个block中的所有线程到达同一点。

        在CUDA中，分别通过cudaError_tcudaDeviceSynchronize(); __device__ void __syncthreads();来完成。

        另外，前文提到的cudaMemcpy函数也会强制同步。

        同步会强制warp处于空闲状态，从而对内核性能产生负面影响。