CUDA学习笔记二_牛客博客

本篇博文记录一些在学习CUDA过程中的一知识点，我的CUDA练手代码github地址：https://github.com/BBuf/CUDA_LEARN_SAMPLES
学习文章：https://blog.csdn.net/sunmc1204953974/article/details/51002061
CUDA程序运行时间计算 首先利用CUDA的clock()函数可以获得CUDA代码的GPU执行单元的频率，也就是GPU的时钟周期（timestamp，然后为了获得时间需要除以GPU的运行频率才能得到以秒为单位的时间。
获取GPU的详细信息 使用cudaGetDeviceProperties这个函数来获取GPU设备的详细信息，这个函数是以*prop形式返回设备dev的属性，一共有两种返回值即cudaSuccess、cudaErrorInvalidDevice，如果是异步启动，可能会返回错误码，整个CudaDeviceProp的结构定义如下：

struct cudaDeviceProp {
char name [256];
size_t totalGlobalMem;
size_t sharedMemPerBlock;
int regsPerBlock;
int warpSize;
size_t memPitch;
int maxThreadsPerBlock;
int maxThreadsDim [3];
int maxGridSize [3];
size_t totalConstMem;
int major;
int minor;
int clockRate;
size_t textureAlignment;
int deviceOverlap;
int multiProcessorCount;
}

其中name代表用于标识设备的ASCII字符串。totalGlobalMem代表设备上可用的全局存储器的总量,以字节为单位。sharedMemPerBlock代表线程块可以使用的共享存储器的最大值,以字节为单位，多处理器上的所有线程块可以同时共享这些存储器。regsPerBlock代表线程块可以使用的32位寄存器的最大值，多处理器上的所有线程块可以同时共享这些寄存器。warpSize是按线程计算的warp块大小。memPitch是允许通过cudaMallocPitch()为包含存储器区域的存储器复制函数分配的最大间距(pitch),以字节为单位。maxThreadsPerBlock是每个块中的最大线程，maxThreadsDim[3]是块各个维度的最大值。maxGridSize[3]是网格各个维度的最大值。totalConstMem是设备上可用的不变存储器总量,以字节为单位。major,minor定义设备计算能力的主要修订号和次要修订号。clockRate是以千赫为单位的时钟频率。textureAlignment对齐要求，与textureAlignment字节对齐的纹理基址无需对纹理取样应用偏移。deviceOverlap如果设备可在主机和设备之间并发复制存储器,同时又能执行内核,则此值为 1，否则此值为 0。multiProcessorCount是设备上多处理器的数量，所以打印出时钟频率clockRate就可以计算时间了。
内存带宽 在GPU上每一秒使用的内存量
并行化程序 在CUDA的函数计算中，第一步是先把主机的数据复制到GPU中，这部分内存叫做global memory，而global memory 并没有 cache，所以使用单线程的话就会有巨大的latency, 也就是说例如求立方和的时候每次都要等到实际读取到数据、累加到 sum 之后，才能进行下一步，这就造成了内存带宽非常低，效率十分地下，所以需要多线程去降低latency，提高运行速度。
threadIdx threadIdx 是 CUDA 的一个内建的变量，表示目前的 thread 是第几个 thread（由 0 开始计算）。
并行优化瓶颈当我把线程数从256变成1024时发现速度已经没有提升了，这个时候计算内存带宽也才几百Mb/s而已，离GPU的理论内存带宽(100G/s)差的太远，那么如何提高内存带宽就要使用到Grid和Block，这都是用来提升内存带宽提速的。
内存存取模式显卡的内存一般是DRAM，所以最有效率的方式是连续存取，而开辟了多线程之后在每个线程里是连续存取内存的，但在实际执行时可能就不连续了，所以在代码里修改存取方式可以加速内存读取。所以在CUDA程序开发中应该更多的去关注如何连续的去操作内存，减少内存存取方面的时间浪费。
thread-block-grid结构 在 CUDA 架构下，显示芯片执行时的最小单位是thread。数个 thread 可以组成一个block。一个 block 中的 thread 能存取同一块共享的内存，而且可以快速进行同步的动作。每一个 block 所能包含的 thread 数目是有限的。不过，执行相同程序的 block，可以组成grid。不同 block 中的 thread 无法存取同一个共享的内存，因此无法直接互通或进行同步。因此，不同 block 中的 thread 能合作的程度是比较低的。不过，利用这个模式，可以让程序不用担心显示芯片实际上能同时执行的 thread 数目限制。例如，一个具有很少量执行单元的显示芯片，可能会把各个 block 中的 thread 顺序执行，而非同时执行。不同的 grid 则可以执行不同的程序（即 kernel）。每个 thread 都有自己的一份 register 和 local memory 的空间。同一个 block 中的每个thread 则有共享的一份 share memory。此外，所有的 thread（包括不同 block 的 thread）都共享一份 global memory、constant memory、和 texture memory。不同的 grid 则有各自的 global memory、constant memory 和 texture memory。
线程越多越好？ 从硬件角度分析，支持CUDA的NVIDIA 显卡，都是由多个multiprocessors 组成每个multiprocessor里包含了8个stream processors，其组成是四个四个一组，也就是两组4D的处理器。每个 multiprocessor 还具有很多个（比如8192个）寄存器，一定的（比如16KB） share memory，以及 texture cache 和 constant cache，结构图如下所示：

在 CUDA 中，大部份基本的运算动作，都可以由 stream processor 进行。每个 stream processor 都包含一个 FMA（fused-multiply-add）单元，可以进行一个乘法和一个加法。比较复杂的运算则会需要比较长的时间。在执行 CUDA 程序的时候，每个 stream processor 就是对应一个 thread。每个 multiprocessor 则对应一个 block。但是我们一个block往往有很大量的线程，之前我们用到了256个和1024个，远超一个multiprocessor所有的8个stream processor。实际上，虽然一个 multiprocessor只有八个stream processor，但是由于 stream processor 进行各种运算都有 latency，更不用提内存存取的 latency，因此 CUDA 在执行程序的时候，是以warp 为单位。比如一个 warp 里面有 32 个 threads，分成两组 16 threads 的 half-warp。由于 stream processor 的运算至少有 4 cycles 的 latency，因此对一个 4D 的stream processors 来说，一次至少执行 16 个 threads（即 half-warp）才能有效隐藏各种运算的 latency。也因此，线程数达到隐藏各种latency的程度后，之后数量的提升就没有太大的作用了。还有一个重要的原因是，由于 multiprocessor 中并没有太多别的内存，因此每个 thread 的状态都是直接保存在multiprocessor 的寄存器中。所以，如果一个 multiprocessor 同时有愈多的 thread 要执行，就会需要愈多的寄存器空间。例如，假设一个 block 里面有 256 个 threads，每个 thread 用到20 个寄存器，那么总共就需要 256x20 = 5,120 个寄存器才能保存每个 thread 的状态。而一般每个 multiprocessor 只有 8,192 个寄存器，因此，如果每个 thread 使用到16 个寄存器，那就表示一个 multiprocessor 的寄存器同时最多只能维持 512 个 thread 的执行。如果同时进行的 thread 数目超过这个数字，那么就会需要把一部份的数据储存在显卡内存中，就会降低执行的效率了。比如一个 warp 里面有 32 个 threads，分成两组 16 threads 的 half-warp。由于 stream processor 的运算至少有 4 cycles 的 latency，因此对一个 4D 的stream processors 来说，一次至少执行 16 个 threads（即 half-warp）才能有效隐藏各种运算的 latency。也因此，线程数达到隐藏各种latency的程度后，之后数量的提升就没有太大的作用了
共享内存 在计算序列每个元素的立方和的和时，通过开辟多线程让计算速度提高了150倍，这是一个很不错的表现了，但是这里面存在一个很大的问题就是当线程数开辟过多，我们在CPU端的求和压力变得很大，所以如何减少CPU端求和压力可以更好的提高我们的代码速度呢？一个 block 内的 thread 可以有共享的内存，也可以进行同步。我们正是可以利用这一点，让每个 block 内的所有 thread 把自己计算的结果加总起来。
Byte字节
Bit 位,二进制数系统中，每个0或1就是一个位(bit)，位是数据存储的最小单位。其中8bit就称为一个字节（Byte）。计算机中的CPU位数指的是CPU一次能处理的最大位数。例如32位计算机的CPU一次最多能处理32位数据。
coalesced 是表示除了连续之外，而且它开始的地址，必须是每个 thread 所存取的大小的 16 倍。例如，如果每个thread 都读取 32 bits 的数据，那么第一个 thread 读取的地址，必须是 16*4 = 64 bytes 的倍数。如果每个 thread 一次存取的数据并不是 32 bits、64 bits、或 128 bits，那就无法符合 coalesced 的条件.不过我们可以通过__align(n)__声明函数来解决这个问题。例子如下：struct __align__(16) vec3d{float x, y, z;}; 这个例子解决的是这样一个核函数：

__global__ void func(struct vec3d* data, float* output)
{
    output[tid] = data[tid].x * data[tid].x + data[tid].y * data[tid].y + data[tid].z * data[tid].z;
}

第二种方法是，把数据结构转换成3个连续的float数组，例如：

__global__ void func(float* x, float* y, float* z, float* output)
{
    output[tid] = x[tid] * x[tid] + y[tid] * y[tid] + z[tid] * z[tid];
}

第3种方法是在第一种和第二种失效的情况下，可以使用shared memrory来调整结构：

__global__ void func(struct vec3d* data, float* output)
{

    __shared__ float temp[THREAD_NUM * 3];
    const float* fdata = (float*) data;
    temp[tid] = fdata[tid];
    temp[tid + THREAD_NUM] = fdata[tid + THREAD_NUM];
    temp[tid + THREAD_NUM*2] = fdata[tid + THREAD_NUM*2];
    //同步
    __syncthreads();
    output[tid] = temp[tid*3] * temp[tid*3] + temp[tid*3+1] * temp[tid*3+1] + temp[tid*3+2] * temp[tid*3+2];
}

们先用连续的方式，把数据从 global memory 读到 shared memory。由于shared memory 不需要担心存取顺序（但要注意 bank conflict 问题，后面马上会讲到），所以可以避开 non-coalesced 读取的问题，这里说一下为什么要用coalesced方式读取内存，因为使用coalesced方式做内存读取的效率比non-coalesced高。
ShareMemory 们要使用的shared memory，是一个 block 中每个 thread 都共享的内存。它会使用在 GPU 上的内存，所以存取的速度相当快，不需要担心 latency 的问题。我们可以直接利用__shared__来声明一个shared memeory变量，例如：__shared__ int sharedata[128];但是从硬件角度分析 Shared memory 有时候会出现一种叫bank conflict的问题。
ShareMemory的bank conflict问题 目前 CUDA 装置中，每个 multiprocessor 有 16KB 的 shared memory。Shared memory 分成16 个 bank。如果同时每个 thread 是存取不同的 bank，就不会产生任何问题，存取 shared memory 的速度和存取寄存器相同。不过，如果同时有两个（或更多个） threads 存取同一个bank 的数据，就会发生 bank conflict，这些 threads 就必须照顺序去存取，而无法同时存取shared memory 了。Shared memory 是以 4 bytes 为单位分成 banks。因此，假设以下的数据：__shared_int data[128];那么dark[0]是bank0，data[1]是bank1，data[2]是bank2、…、data[15]是bank[15]，而data[16]又回到bank0。由于warp 在执行时是以 half-warp 的方式执行，因此分属于不同的 half warp 的 threads，不会造成 bank conflict。在程序存取shared memory时，使用以下方式：int num = data[base+tid];不会有bank conflict, 可以达到最高效率，但换一下：int num = data[base+4*tid];那么thread0和thread4就会存取到同一个bank，thread1和thread5也是一样，这样就会造成bank conflict。在这个例子中，一个 half warp 的 16 个 threads 会有四个 threads 存取同一个 bank，因此存取 share memory 的速度会变成原来的 1/4。一个重要的例外是，当多个 thread 存取到同一个 shared memory 的地址时，shared memory 可以将这个地址的 32 bits 数据「广播」到所有读取的 threads，因此不会造成 bank conflict。例如int num = data[3];但是这里有个问题？广播不需要时间吗？而且通过何种方式广播？除了这些方式，还有很多时候 shared memory 的 bank conflict 可以透过修改数据存放的方式来解决，例如：

data[tid] = global_data[tid];
... 
int number = data[16 * tid];

会有严重的bank conflict，为了避免这个问题，可以把数据的排列方式加以修改，改为：

int row = tid / 16;
int column = tid % 16;
data[row * 17 + column] = global_data[tid];
...
int number = data[17 * tid];

就不会bank conflict。
Thread同步 在CUDA中，想要完成block中的同步还是十分简单的，就是使用一个CUDA 的内部函数：__syncthreads() 它表示block 中所有的 thread 都要同步到这个点才能继续执行。

树状加法 在立方和计算的算法中，blcok的加和是在一个线程thread0上进行的，为了让这一部分也可以并行起来就有了树状加法，这个和ACM的树状数组差不多嘛，看图：

示意图中第一排每一个格子就是一个线程的结果，保存在shared[]，暂且把shared[0]简写为 sh0，我们可以清楚的看到计算的过程，伪代码实现如下：

sh0=sh0+sh1, sh2=sh2+sh3, sh4=sh4+sh5...
同步
sh0=sh0+sh2;sh4=sh4+sh6...
同步
...
最后结果在sh0里

详细代码实现如下：

int offset = 1, mask = 1;
while(offset < THREAD_NUM)
{ 
    if((tid & mask) == 0)  
    {  
         shared[tid] += shared[tid + offset];  
    }  
    offset += offset; 
    mask = offset + mask; 
    __syncthreads(); 
}