CUDA学习笔记（四）

        今天学习一下GPU的全局内存，当然，除了了解通用的GPU外，还会介个jetson nano进行了解和学习。

        主要内容包括几个方面：cuda内存模型、内存管理、内存访问模式、转置矩阵问题。

1.cuda内存模型

        cuda编程中最主要的两个层次模型，一个是之前了解到的线程层次结构，也就是将线程的组织分为了由内到外的thread、block、grid；另一个就是内存层次结构。

        GPU的内存层次结构划分为了寄存器、缓存/共享内存、本地内存、全局内存、常量内存、纹理内存。这些内存的读写速度一次降低，但容量一次增大。

        寄存器：寄存器是距离计算核心最近的内存，读写速度是最快的。在GPU中，每一个线程都会私有一个寄存器用来计算核函数。一个核函数使用了较多的寄存器会影响到性能。

        共享内存：共享内存也是片上内存，所以与全局内存相比拥有较高的带宽和较低的延迟。共享内存是一个block中的所有线程共享的。因为共享内存的大小是较小的，所以当一个核函数使用共享内存时，共享内存的使用大小会直接影响一个block中的warp的活跃数量。当核函数使用的共享内存较大时，一个block中活跃的warp数量就会很少，这会导致sm中的计算资源没有充分使用。共享内存的定义是使用修饰符__shared__。

        本地内存：这块内存一般是当寄存器不够使用时，也就是寄存器溢出才会使用到。

        全局内存：全局内存是内存空间最大，延迟最高，最常使用的内存。在主机端使用cudamalloc函数申请内存时，就是在这块内存上进行开辟空间。所有的线程都可以访问全局内存。

        常量内存：常量内存是用来存储只读不可写的数据，其定义方式是使用修饰符__constant__。线程束中的所有线程都读取同一个地址时，常量内存表现最好，因为内从一个常量内存中读取一个数据，就会在线程束中广播传递给线程束内的所有线程。

        纹理内存：纹理内存也是只可读的聂村，但它为各种数据布局提供了不同的寻址模式和滤波模式。

        缓存：上面所有的内存都是可编程的。但缓存是不可编程的。GPU的缓存主要有四种：一级缓存、二级缓存、只读常量缓存、只读纹理缓存。

        一级缓存和二级缓存都是用来存储本地内存和全局内存的数据，以便保证空间局部性和时间局部性。而只读常量缓存是针对常量内存数据的存储，只读纹理缓存是针对纹理内存数据的存储。

        ps：GPU中的一级缓存和共享内存使用的是一块内存，普遍大小为64kb，但是可以通过手动设置，使缓存和共享内存的划分来配合我们的需要。

        cudaFuncSetCacheConfig（const void *func,enum cudaFuncCache cacheConfig）;  

        其中cacheConfig有几种配置：        

                cudaFuncCachePreferNone：没有参考值（默认）；

                cudaFuncCachePreferShared：48kb共享内存，16kb的一级缓存；

                cudaFuncCachePreferL1：16kb的共享内存，48kb的一级缓存；

                cudaFuncCachePreferEqual：两个都是32kb。

2.内存管理

        内存管理其实就是对gpu上的内存进行分配、释放、以及数据传输等。

        具体api主要有：

                cudaMalloc((void **)d_ptr,size_t count);  全局内存申请；

                cudaMemset(void *d_ptr,int value,size_t count);  对dd_ptr指向内存的指定字节赋值value；

                cudaFree(void *d_ptr);  释放掉申请的全局内存空间。

                cudamemcpy(void *dst,const void *src,size_t count,enum cudaMemcpyKind kind); 将src指向的内存数据拷贝到dst指向内存，其中kind参数是表明数据拷贝的方向，具体有：

                cudaMemcpyHostToDevice:主机端拷贝到设备端；

                cudaMemcpyDeviceToHost：设别断拷贝到主机端；

                cudaMemcpyDeviceToDevice：设别拷贝到设备；

                cudaMemcpyHostToHost:主机拷贝到主机；

        另外说一下固定内存。在主机端，内存是默认是可分页的，所以直接从主机端的内存向设备端进行数据拷贝是肯呢个存在不安全的，原因就是主机端的页面是可能会进行替换。为了利用上主机端内存的大吞吐量的数据传输优点，cuda提供了接口在主机端申请一个固定的内存，也就是说这块申请的固定内存不会随着系统运行而出现替换的情况。

        cudaMallocHost(void **ptr,size_t count);

        使用上述接口可以在主机端分配一块固定内存，使用这块内存向设备端进行数据传输有更高的数据吞吐量。但是，使用固定内存会导致主机端能够用来存储虚拟内存的可分页内存减少，导致主机端性能降低。在使用完固定内存后，使用以下接口释放掉内存：

        cudaFreeHost(void *ptr);

        使用固定内存会导致主机端的性能降低，但其实在数据传输的时候使用cudamemcpy时，主机和设备端时同步的，也就是主机端必须等待设别断数据传输完毕才能向下运行，所以尽管主机端性能会降低，但在大数据传输的时候，固定内存还是比较适用的。

        另外还有一个就是零拷贝内存。通常情况下，主机是不能访问设别内存，设备也不能访问主机内存，但1通过零拷贝内存可以实现让主机和设备都可以直接访问这块内存。

        cudaHostAlloc(void **phost,size_t count,usigned int flags);

        其中flags有：  cudaHostAllocDefat：使该函数行为和使用固定内存一致；

                                cudaHostAllocPortable：

                                cudaHostAllocWriteCombined

                                cudaHostAllocMapped：返回一个可以映射的主机内存，方便映射到设备端，使设备端能够直接访问这块内存。

        所以，使用零拷贝内存的流程是：

                现在主机端使用cudaHostAlloc()申请一块零拷贝地址，flag使用cudaHostAllocMapped，然后使用接口cudaHostGetDevicePointer()将主机内存引用到一个设备端指针，然后设备端就可以直接使用映射的设备端指针使用零拷贝内存了。

        ps：使用零拷贝内存时，由于是主机端和设备端都能够访问该内存，所以当主机端和设备端在对该内存进行读写的时候应该做好同步，不然会导致不可预估的问题。

3.内存访问模式

        cuda执行模式的显著特征之一就是执行必须以线程束为单位进行发布和执行。一个线程束中的32个线程，每个线程都会提出一个包含请求地址的单一内存访问请求。根据线程数中内存地址的分布，内存访问可以被分为不同的模式。

        核函数的内存请求通常是在DRAM设备和片上内存间以128字节或32字节内存事务来实现的。所有对全局内存的访问都会通过二级缓存，也有许多访问会通过一级缓存。如果这两级缓存都被用到，那么内存访问是由一个128字节的内存事务实现的。如果只使用了二级缓存，那么这个内存访问是由一个32字节的内存事务实现的。

        因此在优化应用程序时，需要注意设备内存访问的两个特性：

                对齐内存访问：当设备内存事务的第一个地址是用于事务服务的缓存粒度的偶数倍时（32个字节的二级缓存或128字节的一级缓存），就会出现对齐内存访问。

                合并内存访问：当一个线程束中的32个线程访问一块连续的内存块时，就会出现合并内存访问。

        所以在使用结构体数组和数组结构体的时候，有限使用数组结构体。因为在结构体数组中，一个元素内部是一个结构体，存放着不同的数据，那么整个结构体数组在内存中的存放就会是以下排列方式abc abc abc abc，要访问该结构体数组中的全部a元素时，就无法满足合并内存访问，造成带宽浪费。而数组结构体中，内存存放是aaaabbbbcccc，非常方便合并内存访问。

        ps：cpu上的一级缓存优化了时间局部性和空间局部性。gpu一级缓存是专为空间局部性而不是时间局部性设计的。频繁访问一级缓存上的数据不会增加数据在一级缓存上的保留概率。